diff --git a/CHANGELOG.md b/CHANGELOG.md
index d0d627a8..072d2c53 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -10,6 +10,7 @@
* `cppOverload` demonstrating C++ function overloading (reason: obsolete)
* `simpleSeparateCompilation` demonstrating NVCC compilation to a static library (reason: trivial)
* `simpleTemplates_nvrtc` demonstrating NVRTC usage for `simpleTemplates` sample (reason: redundant)
+ * `simpleVoteIntrinsics_nvrtc` demonstrating NVRTC usage for `simpleVoteIntrinsics` sample (reason: redundant)
### CUDA 12.5
diff --git a/Samples/0_Introduction/CMakeLists.txt b/Samples/0_Introduction/CMakeLists.txt
index a54161b8..e195e4ec 100644
--- a/Samples/0_Introduction/CMakeLists.txt
+++ b/Samples/0_Introduction/CMakeLists.txt
@@ -36,8 +36,7 @@ add_subdirectory(simpleTemplates)
add_subdirectory(simpleTexture)
add_subdirectory(simpleTexture3D)
add_subdirectory(simpleTextureDrv)
-#add_subdirectory(simpleVoteIntrinsics)
-#add_subdirectory(simpleVoteIntrinsics_nvrtc)
+add_subdirectory(simpleVoteIntrinsics)
add_subdirectory(simpleZeroCopy)
add_subdirectory(systemWideAtomics)
#add_subdirectory(template)
diff --git a/Samples/0_Introduction/simpleVoteIntrinsics/CMakeLists.txt b/Samples/0_Introduction/simpleVoteIntrinsics/CMakeLists.txt
new file mode 100644
index 00000000..c98318f6
--- /dev/null
+++ b/Samples/0_Introduction/simpleVoteIntrinsics/CMakeLists.txt
@@ -0,0 +1,11 @@
+# Include directories and libraries
+include_directories(../../../Common)
+
+# Source file
+set(SRC_FILES
+ simpleVoteIntrinsics.cu
+)
+
+# Add target for simpleVoteIntrinsics
+add_executable(simpleVoteIntrinsics ${SRC_FILES})
+set_target_properties(simpleVoteIntrinsics PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
diff --git a/Samples/0_Introduction/simpleVoteIntrinsics/Makefile b/Samples/0_Introduction/simpleVoteIntrinsics/Makefile
deleted file mode 100644
index 4e338815..00000000
--- a/Samples/0_Introduction/simpleVoteIntrinsics/Makefile
+++ /dev/null
@@ -1,357 +0,0 @@
-################################################################################
-# Copyright (c) 2022, NVIDIA CORPORATION. All rights reserved.
-#
-# Redistribution and use in source and binary forms, with or without
-# modification, are permitted provided that the following conditions
-# are met:
-# * Redistributions of source code must retain the above copyright
-# notice, this list of conditions and the following disclaimer.
-# * Redistributions in binary form must reproduce the above copyright
-# notice, this list of conditions and the following disclaimer in the
-# documentation and/or other materials provided with the distribution.
-# * Neither the name of NVIDIA CORPORATION nor the names of its
-# contributors may be used to endorse or promote products derived
-# from this software without specific prior written permission.
-#
-# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
-# EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
-# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
-# PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
-# CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
-# EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
-# PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
-# PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
-# OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
-# (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
-# OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-#
-################################################################################
-#
-# Makefile project only supported on Mac OS X and Linux Platforms)
-#
-################################################################################
-
-# Location of the CUDA Toolkit
-CUDA_PATH ?= /usr/local/cuda
-
-##############################
-# start deprecated interface #
-##############################
-ifeq ($(x86_64),1)
- $(info WARNING - x86_64 variable has been deprecated)
- $(info WARNING - please use TARGET_ARCH=x86_64 instead)
- TARGET_ARCH ?= x86_64
-endif
-ifeq ($(ARMv7),1)
- $(info WARNING - ARMv7 variable has been deprecated)
- $(info WARNING - please use TARGET_ARCH=armv7l instead)
- TARGET_ARCH ?= armv7l
-endif
-ifeq ($(aarch64),1)
- $(info WARNING - aarch64 variable has been deprecated)
- $(info WARNING - please use TARGET_ARCH=aarch64 instead)
- TARGET_ARCH ?= aarch64
-endif
-ifeq ($(ppc64le),1)
- $(info WARNING - ppc64le variable has been deprecated)
- $(info WARNING - please use TARGET_ARCH=ppc64le instead)
- TARGET_ARCH ?= ppc64le
-endif
-ifneq ($(GCC),)
- $(info WARNING - GCC variable has been deprecated)
- $(info WARNING - please use HOST_COMPILER=$(GCC) instead)
- HOST_COMPILER ?= $(GCC)
-endif
-ifneq ($(abi),)
- $(error ERROR - abi variable has been removed)
-endif
-############################
-# end deprecated interface #
-############################
-
-# architecture
-HOST_ARCH := $(shell uname -m)
-TARGET_ARCH ?= $(HOST_ARCH)
-ifneq (,$(filter $(TARGET_ARCH),x86_64 aarch64 sbsa ppc64le armv7l))
- ifneq ($(TARGET_ARCH),$(HOST_ARCH))
- ifneq (,$(filter $(TARGET_ARCH),x86_64 aarch64 sbsa ppc64le))
- TARGET_SIZE := 64
- else ifneq (,$(filter $(TARGET_ARCH),armv7l))
- TARGET_SIZE := 32
- endif
- else
- TARGET_SIZE := $(shell getconf LONG_BIT)
- endif
-else
- $(error ERROR - unsupported value $(TARGET_ARCH) for TARGET_ARCH!)
-endif
-
-# sbsa and aarch64 systems look similar. Need to differentiate them at host level for now.
-ifeq ($(HOST_ARCH),aarch64)
- ifeq ($(CUDA_PATH)/targets/sbsa-linux,$(shell ls -1d $(CUDA_PATH)/targets/sbsa-linux 2>/dev/null))
- HOST_ARCH := sbsa
- TARGET_ARCH := sbsa
- endif
-endif
-
-ifneq ($(TARGET_ARCH),$(HOST_ARCH))
- ifeq (,$(filter $(HOST_ARCH)-$(TARGET_ARCH),aarch64-armv7l x86_64-armv7l x86_64-aarch64 x86_64-sbsa x86_64-ppc64le))
- $(error ERROR - cross compiling from $(HOST_ARCH) to $(TARGET_ARCH) is not supported!)
- endif
-endif
-
-# When on native aarch64 system with userspace of 32-bit, change TARGET_ARCH to armv7l
-ifeq ($(HOST_ARCH)-$(TARGET_ARCH)-$(TARGET_SIZE),aarch64-aarch64-32)
- TARGET_ARCH = armv7l
-endif
-
-# operating system
-HOST_OS := $(shell uname -s 2>/dev/null | tr "[:upper:]" "[:lower:]")
-TARGET_OS ?= $(HOST_OS)
-ifeq (,$(filter $(TARGET_OS),linux darwin qnx android))
- $(error ERROR - unsupported value $(TARGET_OS) for TARGET_OS!)
-endif
-
-# host compiler
-ifdef HOST_COMPILER
- CUSTOM_HOST_COMPILER = 1
-endif
-
-ifeq ($(TARGET_OS),darwin)
- ifeq ($(shell expr `xcodebuild -version | grep -i xcode | awk '{print $$2}' | cut -d'.' -f1` \>= 5),1)
- HOST_COMPILER ?= clang++
- endif
-else ifneq ($(TARGET_ARCH),$(HOST_ARCH))
- ifeq ($(HOST_ARCH)-$(TARGET_ARCH),x86_64-armv7l)
- ifeq ($(TARGET_OS),linux)
- HOST_COMPILER ?= arm-linux-gnueabihf-g++
- else ifeq ($(TARGET_OS),qnx)
- ifeq ($(QNX_HOST),)
- $(error ERROR - QNX_HOST must be passed to the QNX host toolchain)
- endif
- ifeq ($(QNX_TARGET),)
- $(error ERROR - QNX_TARGET must be passed to the QNX target toolchain)
- endif
- export QNX_HOST
- export QNX_TARGET
- HOST_COMPILER ?= $(QNX_HOST)/usr/bin/arm-unknown-nto-qnx6.6.0eabi-g++
- else ifeq ($(TARGET_OS),android)
- HOST_COMPILER ?= arm-linux-androideabi-g++
- endif
- else ifeq ($(TARGET_ARCH),aarch64)
- ifeq ($(TARGET_OS), linux)
- HOST_COMPILER ?= aarch64-linux-gnu-g++
- else ifeq ($(TARGET_OS),qnx)
- ifeq ($(QNX_HOST),)
- $(error ERROR - QNX_HOST must be passed to the QNX host toolchain)
- endif
- ifeq ($(QNX_TARGET),)
- $(error ERROR - QNX_TARGET must be passed to the QNX target toolchain)
- endif
- export QNX_HOST
- export QNX_TARGET
- HOST_COMPILER ?= $(QNX_HOST)/usr/bin/q++
- else ifeq ($(TARGET_OS), android)
- HOST_COMPILER ?= aarch64-linux-android-clang++
- endif
- else ifeq ($(TARGET_ARCH),sbsa)
- HOST_COMPILER ?= aarch64-linux-gnu-g++
- else ifeq ($(TARGET_ARCH),ppc64le)
- HOST_COMPILER ?= powerpc64le-linux-gnu-g++
- endif
-endif
-HOST_COMPILER ?= g++
-NVCC := $(CUDA_PATH)/bin/nvcc -ccbin $(HOST_COMPILER)
-
-# internal flags
-NVCCFLAGS := -m${TARGET_SIZE}
-CCFLAGS :=
-LDFLAGS :=
-
-# build flags
-
-# Link flag for customized HOST_COMPILER with gcc realpath
-GCC_PATH := $(shell which gcc)
-ifeq ($(CUSTOM_HOST_COMPILER),1)
- ifneq ($(filter /%,$(HOST_COMPILER)),)
- ifneq ($(findstring gcc,$(HOST_COMPILER)),)
- ifneq ($(GCC_PATH),$(HOST_COMPILER))
- LDFLAGS += -lstdc++
- endif
- endif
- endif
-endif
-
-ifeq ($(TARGET_OS),darwin)
- LDFLAGS += -rpath $(CUDA_PATH)/lib
- CCFLAGS += -arch $(HOST_ARCH)
-else ifeq ($(HOST_ARCH)-$(TARGET_ARCH)-$(TARGET_OS),x86_64-armv7l-linux)
- LDFLAGS += --dynamic-linker=/lib/ld-linux-armhf.so.3
- CCFLAGS += -mfloat-abi=hard
-else ifeq ($(TARGET_OS),android)
- LDFLAGS += -pie
- CCFLAGS += -fpie -fpic -fexceptions
-endif
-
-ifneq ($(TARGET_ARCH),$(HOST_ARCH))
- ifeq ($(TARGET_ARCH)-$(TARGET_OS),armv7l-linux)
- ifneq ($(TARGET_FS),)
- GCCVERSIONLTEQ46 := $(shell expr `$(HOST_COMPILER) -dumpversion` \<= 4.6)
- ifeq ($(GCCVERSIONLTEQ46),1)
- CCFLAGS += --sysroot=$(TARGET_FS)
- endif
- LDFLAGS += --sysroot=$(TARGET_FS)
- LDFLAGS += -rpath-link=$(TARGET_FS)/lib
- LDFLAGS += -rpath-link=$(TARGET_FS)/usr/lib
- LDFLAGS += -rpath-link=$(TARGET_FS)/usr/lib/arm-linux-gnueabihf
- endif
- endif
- ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-linux)
- ifneq ($(TARGET_FS),)
- GCCVERSIONLTEQ46 := $(shell expr `$(HOST_COMPILER) -dumpversion` \<= 4.6)
- ifeq ($(GCCVERSIONLTEQ46),1)
- CCFLAGS += --sysroot=$(TARGET_FS)
- endif
- LDFLAGS += --sysroot=$(TARGET_FS)
- LDFLAGS += -rpath-link=$(TARGET_FS)/lib -L$(TARGET_FS)/lib
- LDFLAGS += -rpath-link=$(TARGET_FS)/lib/aarch64-linux-gnu -L$(TARGET_FS)/lib/aarch64-linux-gnu
- LDFLAGS += -rpath-link=$(TARGET_FS)/usr/lib -L$(TARGET_FS)/usr/lib
- LDFLAGS += -rpath-link=$(TARGET_FS)/usr/lib/aarch64-linux-gnu -L$(TARGET_FS)/usr/lib/aarch64-linux-gnu
- LDFLAGS += --unresolved-symbols=ignore-in-shared-libs
- CCFLAGS += -isystem=$(TARGET_FS)/usr/include -I$(TARGET_FS)/usr/include -I$(TARGET_FS)/usr/include/libdrm
- CCFLAGS += -isystem=$(TARGET_FS)/usr/include/aarch64-linux-gnu -I$(TARGET_FS)/usr/include/aarch64-linux-gnu
- endif
- endif
- ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-qnx)
- NVCCFLAGS += -D_QNX_SOURCE
- NVCCFLAGS += --qpp-config 8.3.0,gcc_ntoaarch64le
- CCFLAGS += -DWIN_INTERFACE_CUSTOM -I/usr/include/aarch64-qnx-gnu
- LDFLAGS += -lsocket
- LDFLAGS += -L/usr/lib/aarch64-qnx-gnu
- CCFLAGS += "-Wl\,-rpath-link\,/usr/lib/aarch64-qnx-gnu"
- ifdef TARGET_OVERRIDE
- LDFLAGS += -lslog2
- endif
-
- ifneq ($(TARGET_FS),)
- LDFLAGS += -L$(TARGET_FS)/usr/lib
- CCFLAGS += "-Wl\,-rpath-link\,$(TARGET_FS)/usr/lib"
- LDFLAGS += -L$(TARGET_FS)/usr/libnvidia
- CCFLAGS += "-Wl\,-rpath-link\,$(TARGET_FS)/usr/libnvidia"
- CCFLAGS += -I$(TARGET_FS)/../include
- endif
- endif
-endif
-
-ifdef TARGET_OVERRIDE # cuda toolkit targets override
- NVCCFLAGS += -target-dir $(TARGET_OVERRIDE)
-endif
-
-# Install directory of different arch
-CUDA_INSTALL_TARGET_DIR :=
-ifeq ($(TARGET_ARCH)-$(TARGET_OS),armv7l-linux)
- CUDA_INSTALL_TARGET_DIR = targets/armv7-linux-gnueabihf/
-else ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-linux)
- CUDA_INSTALL_TARGET_DIR = targets/aarch64-linux/
-else ifeq ($(TARGET_ARCH)-$(TARGET_OS),sbsa-linux)
- CUDA_INSTALL_TARGET_DIR = targets/sbsa-linux/
-else ifeq ($(TARGET_ARCH)-$(TARGET_OS),armv7l-android)
- CUDA_INSTALL_TARGET_DIR = targets/armv7-linux-androideabi/
-else ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-android)
- CUDA_INSTALL_TARGET_DIR = targets/aarch64-linux-androideabi/
-else ifeq ($(TARGET_ARCH)-$(TARGET_OS),armv7l-qnx)
- CUDA_INSTALL_TARGET_DIR = targets/ARMv7-linux-QNX/
-else ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-qnx)
- CUDA_INSTALL_TARGET_DIR = targets/aarch64-qnx/
-else ifeq ($(TARGET_ARCH),ppc64le)
- CUDA_INSTALL_TARGET_DIR = targets/ppc64le-linux/
-endif
-
-# Debug build flags
-ifeq ($(dbg),1)
- NVCCFLAGS += -g -G
- BUILD_TYPE := debug
-else
- BUILD_TYPE := release
-endif
-
-ALL_CCFLAGS :=
-ALL_CCFLAGS += $(NVCCFLAGS)
-ALL_CCFLAGS += $(EXTRA_NVCCFLAGS)
-ALL_CCFLAGS += $(addprefix -Xcompiler ,$(CCFLAGS))
-ALL_CCFLAGS += $(addprefix -Xcompiler ,$(EXTRA_CCFLAGS))
-
-SAMPLE_ENABLED := 1
-
-ALL_LDFLAGS :=
-ALL_LDFLAGS += $(ALL_CCFLAGS)
-ALL_LDFLAGS += $(addprefix -Xlinker ,$(LDFLAGS))
-ALL_LDFLAGS += $(addprefix -Xlinker ,$(EXTRA_LDFLAGS))
-
-# Common includes and paths for CUDA
-INCLUDES := -I../../../Common
-LIBRARIES :=
-
-################################################################################
-
-# Gencode arguments
-ifeq ($(TARGET_ARCH),$(filter $(TARGET_ARCH),armv7l aarch64 sbsa))
-SMS ?= 53 61 70 72 75 80 86 87 90
-else
-SMS ?= 50 52 60 61 70 75 80 86 89 90
-endif
-
-ifeq ($(SMS),)
-$(info >>> WARNING - no SM architectures have been specified - waiving sample <<<)
-SAMPLE_ENABLED := 0
-endif
-
-ifeq ($(GENCODE_FLAGS),)
-# Generate SASS code for each SM architecture listed in $(SMS)
-$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm)))
-
-# Generate PTX code from the highest SM architecture in $(SMS) to guarantee forward-compatibility
-HIGHEST_SM := $(lastword $(sort $(SMS)))
-ifneq ($(HIGHEST_SM),)
-GENCODE_FLAGS += -gencode arch=compute_$(HIGHEST_SM),code=compute_$(HIGHEST_SM)
-endif
-endif
-
-ALL_CCFLAGS += --threads 0 --std=c++11
-
-ifeq ($(SAMPLE_ENABLED),0)
-EXEC ?= @echo "[@]"
-endif
-
-################################################################################
-
-# Target rules
-all: build
-
-build: simpleVoteIntrinsics
-
-check.deps:
-ifeq ($(SAMPLE_ENABLED),0)
- @echo "Sample will be waived due to the above missing dependencies"
-else
- @echo "Sample is ready - all dependencies have been met"
-endif
-
-simpleVoteIntrinsics.o:simpleVoteIntrinsics.cu
- $(EXEC) $(NVCC) $(INCLUDES) $(ALL_CCFLAGS) $(GENCODE_FLAGS) -o $@ -c $<
-
-simpleVoteIntrinsics: simpleVoteIntrinsics.o
- $(EXEC) $(NVCC) $(ALL_LDFLAGS) $(GENCODE_FLAGS) -o $@ $+ $(LIBRARIES)
- $(EXEC) mkdir -p ../../../bin/$(TARGET_ARCH)/$(TARGET_OS)/$(BUILD_TYPE)
- $(EXEC) cp $@ ../../../bin/$(TARGET_ARCH)/$(TARGET_OS)/$(BUILD_TYPE)
-
-run: build
- $(EXEC) ./simpleVoteIntrinsics
-
-testrun: build
-
-clean:
- rm -f simpleVoteIntrinsics simpleVoteIntrinsics.o
- rm -rf ../../../bin/$(TARGET_ARCH)/$(TARGET_OS)/$(BUILD_TYPE)/simpleVoteIntrinsics
-
-clobber: clean
diff --git a/Samples/0_Introduction/simpleVoteIntrinsics/NsightEclipse.xml b/Samples/0_Introduction/simpleVoteIntrinsics/NsightEclipse.xml
deleted file mode 100644
index 47662a24..00000000
--- a/Samples/0_Introduction/simpleVoteIntrinsics/NsightEclipse.xml
+++ /dev/null
@@ -1,80 +0,0 @@
-
-
-
- simpleVoteIntrinsics
-
- cudaMemcpy
- cudaFree
- cudaDeviceSynchronize
- cudaMalloc
- cudaGetDeviceProperties
-
-
- whole
-
- ./
- ../
- ../../../Common
-
-
- Vote Intrinsics
-
-
- CUDA
- GPGPU
- vote
-
-
-
-
-
- true
- simpleVoteIntrinsics.cu
-
- 1:CUDA Basic Topics
-
- sm50
- sm52
- sm53
- sm60
- sm61
- sm70
- sm72
- sm75
- sm80
- sm86
- sm87
- sm89
- sm90
-
-
- x86_64
- linux
-
-
- windows7
-
-
- x86_64
- macosx
-
-
- arm
-
-
- aarch64
-
-
- sbsa
-
-
- ppc64le
- linux
-
-
-
- all
-
- Simple Vote Intrinsics
- exe
-
diff --git a/Samples/0_Introduction/simpleVoteIntrinsics_nvrtc/.vscode/c_cpp_properties.json b/Samples/0_Introduction/simpleVoteIntrinsics_nvrtc/.vscode/c_cpp_properties.json
deleted file mode 100644
index f0066b0f..00000000
--- a/Samples/0_Introduction/simpleVoteIntrinsics_nvrtc/.vscode/c_cpp_properties.json
+++ /dev/null
@@ -1,18 +0,0 @@
-{
- "configurations": [
- {
- "name": "Linux",
- "includePath": [
- "${workspaceFolder}/**",
- "${workspaceFolder}/../../../Common"
- ],
- "defines": [],
- "compilerPath": "/usr/local/cuda/bin/nvcc",
- "cStandard": "gnu17",
- "cppStandard": "gnu++14",
- "intelliSenseMode": "linux-gcc-x64",
- "configurationProvider": "ms-vscode.makefile-tools"
- }
- ],
- "version": 4
-}
diff --git a/Samples/0_Introduction/simpleVoteIntrinsics_nvrtc/.vscode/extensions.json b/Samples/0_Introduction/simpleVoteIntrinsics_nvrtc/.vscode/extensions.json
deleted file mode 100644
index c7eb54dc..00000000
--- a/Samples/0_Introduction/simpleVoteIntrinsics_nvrtc/.vscode/extensions.json
+++ /dev/null
@@ -1,7 +0,0 @@
-{
- "recommendations": [
- "nvidia.nsight-vscode-edition",
- "ms-vscode.cpptools",
- "ms-vscode.makefile-tools"
- ]
-}
diff --git a/Samples/0_Introduction/simpleVoteIntrinsics_nvrtc/.vscode/launch.json b/Samples/0_Introduction/simpleVoteIntrinsics_nvrtc/.vscode/launch.json
deleted file mode 100644
index fc69f52b..00000000
--- a/Samples/0_Introduction/simpleVoteIntrinsics_nvrtc/.vscode/launch.json
+++ /dev/null
@@ -1,10 +0,0 @@
-{
- "configurations": [
- {
- "name": "CUDA C++: Launch",
- "type": "cuda-gdb",
- "request": "launch",
- "program": "${workspaceFolder}/simpleVoteIntrinsics_nvrtc"
- }
- ]
-}
diff --git a/Samples/0_Introduction/simpleVoteIntrinsics_nvrtc/.vscode/tasks.json b/Samples/0_Introduction/simpleVoteIntrinsics_nvrtc/.vscode/tasks.json
deleted file mode 100644
index 4509aeb1..00000000
--- a/Samples/0_Introduction/simpleVoteIntrinsics_nvrtc/.vscode/tasks.json
+++ /dev/null
@@ -1,15 +0,0 @@
-{
- "version": "2.0.0",
- "tasks": [
- {
- "label": "sample",
- "type": "shell",
- "command": "make dbg=1",
- "problemMatcher": ["$nvcc"],
- "group": {
- "kind": "build",
- "isDefault": true
- }
- }
- ]
-}
diff --git a/Samples/0_Introduction/simpleVoteIntrinsics_nvrtc/Makefile b/Samples/0_Introduction/simpleVoteIntrinsics_nvrtc/Makefile
deleted file mode 100644
index 4999d92b..00000000
--- a/Samples/0_Introduction/simpleVoteIntrinsics_nvrtc/Makefile
+++ /dev/null
@@ -1,409 +0,0 @@
-################################################################################
-# Copyright (c) 2022, NVIDIA CORPORATION. All rights reserved.
-#
-# Redistribution and use in source and binary forms, with or without
-# modification, are permitted provided that the following conditions
-# are met:
-# * Redistributions of source code must retain the above copyright
-# notice, this list of conditions and the following disclaimer.
-# * Redistributions in binary form must reproduce the above copyright
-# notice, this list of conditions and the following disclaimer in the
-# documentation and/or other materials provided with the distribution.
-# * Neither the name of NVIDIA CORPORATION nor the names of its
-# contributors may be used to endorse or promote products derived
-# from this software without specific prior written permission.
-#
-# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
-# EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
-# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
-# PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
-# CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
-# EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
-# PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
-# PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
-# OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
-# (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
-# OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-#
-################################################################################
-#
-# Makefile project only supported on Mac OS X and Linux Platforms)
-#
-################################################################################
-
-# Location of the CUDA Toolkit
-CUDA_PATH ?= /usr/local/cuda
-
-##############################
-# start deprecated interface #
-##############################
-ifeq ($(x86_64),1)
- $(info WARNING - x86_64 variable has been deprecated)
- $(info WARNING - please use TARGET_ARCH=x86_64 instead)
- TARGET_ARCH ?= x86_64
-endif
-ifeq ($(ARMv7),1)
- $(info WARNING - ARMv7 variable has been deprecated)
- $(info WARNING - please use TARGET_ARCH=armv7l instead)
- TARGET_ARCH ?= armv7l
-endif
-ifeq ($(aarch64),1)
- $(info WARNING - aarch64 variable has been deprecated)
- $(info WARNING - please use TARGET_ARCH=aarch64 instead)
- TARGET_ARCH ?= aarch64
-endif
-ifeq ($(ppc64le),1)
- $(info WARNING - ppc64le variable has been deprecated)
- $(info WARNING - please use TARGET_ARCH=ppc64le instead)
- TARGET_ARCH ?= ppc64le
-endif
-ifneq ($(GCC),)
- $(info WARNING - GCC variable has been deprecated)
- $(info WARNING - please use HOST_COMPILER=$(GCC) instead)
- HOST_COMPILER ?= $(GCC)
-endif
-ifneq ($(abi),)
- $(error ERROR - abi variable has been removed)
-endif
-############################
-# end deprecated interface #
-############################
-
-# architecture
-HOST_ARCH := $(shell uname -m)
-TARGET_ARCH ?= $(HOST_ARCH)
-ifneq (,$(filter $(TARGET_ARCH),x86_64 aarch64 sbsa ppc64le armv7l))
- ifneq ($(TARGET_ARCH),$(HOST_ARCH))
- ifneq (,$(filter $(TARGET_ARCH),x86_64 aarch64 sbsa ppc64le))
- TARGET_SIZE := 64
- else ifneq (,$(filter $(TARGET_ARCH),armv7l))
- TARGET_SIZE := 32
- endif
- else
- TARGET_SIZE := $(shell getconf LONG_BIT)
- endif
-else
- $(error ERROR - unsupported value $(TARGET_ARCH) for TARGET_ARCH!)
-endif
-
-# sbsa and aarch64 systems look similar. Need to differentiate them at host level for now.
-ifeq ($(HOST_ARCH),aarch64)
- ifeq ($(CUDA_PATH)/targets/sbsa-linux,$(shell ls -1d $(CUDA_PATH)/targets/sbsa-linux 2>/dev/null))
- HOST_ARCH := sbsa
- TARGET_ARCH := sbsa
- endif
-endif
-
-ifneq ($(TARGET_ARCH),$(HOST_ARCH))
- ifeq (,$(filter $(HOST_ARCH)-$(TARGET_ARCH),aarch64-armv7l x86_64-armv7l x86_64-aarch64 x86_64-sbsa x86_64-ppc64le))
- $(error ERROR - cross compiling from $(HOST_ARCH) to $(TARGET_ARCH) is not supported!)
- endif
-endif
-
-# When on native aarch64 system with userspace of 32-bit, change TARGET_ARCH to armv7l
-ifeq ($(HOST_ARCH)-$(TARGET_ARCH)-$(TARGET_SIZE),aarch64-aarch64-32)
- TARGET_ARCH = armv7l
-endif
-
-# operating system
-HOST_OS := $(shell uname -s 2>/dev/null | tr "[:upper:]" "[:lower:]")
-TARGET_OS ?= $(HOST_OS)
-ifeq (,$(filter $(TARGET_OS),linux darwin qnx android))
- $(error ERROR - unsupported value $(TARGET_OS) for TARGET_OS!)
-endif
-
-# host compiler
-ifdef HOST_COMPILER
- CUSTOM_HOST_COMPILER = 1
-endif
-
-ifeq ($(TARGET_OS),darwin)
- ifeq ($(shell expr `xcodebuild -version | grep -i xcode | awk '{print $$2}' | cut -d'.' -f1` \>= 5),1)
- HOST_COMPILER ?= clang++
- endif
-else ifneq ($(TARGET_ARCH),$(HOST_ARCH))
- ifeq ($(HOST_ARCH)-$(TARGET_ARCH),x86_64-armv7l)
- ifeq ($(TARGET_OS),linux)
- HOST_COMPILER ?= arm-linux-gnueabihf-g++
- else ifeq ($(TARGET_OS),qnx)
- ifeq ($(QNX_HOST),)
- $(error ERROR - QNX_HOST must be passed to the QNX host toolchain)
- endif
- ifeq ($(QNX_TARGET),)
- $(error ERROR - QNX_TARGET must be passed to the QNX target toolchain)
- endif
- export QNX_HOST
- export QNX_TARGET
- HOST_COMPILER ?= $(QNX_HOST)/usr/bin/arm-unknown-nto-qnx6.6.0eabi-g++
- else ifeq ($(TARGET_OS),android)
- HOST_COMPILER ?= arm-linux-androideabi-g++
- endif
- else ifeq ($(TARGET_ARCH),aarch64)
- ifeq ($(TARGET_OS), linux)
- HOST_COMPILER ?= aarch64-linux-gnu-g++
- else ifeq ($(TARGET_OS),qnx)
- ifeq ($(QNX_HOST),)
- $(error ERROR - QNX_HOST must be passed to the QNX host toolchain)
- endif
- ifeq ($(QNX_TARGET),)
- $(error ERROR - QNX_TARGET must be passed to the QNX target toolchain)
- endif
- export QNX_HOST
- export QNX_TARGET
- HOST_COMPILER ?= $(QNX_HOST)/usr/bin/q++
- else ifeq ($(TARGET_OS), android)
- HOST_COMPILER ?= aarch64-linux-android-clang++
- endif
- else ifeq ($(TARGET_ARCH),sbsa)
- HOST_COMPILER ?= aarch64-linux-gnu-g++
- else ifeq ($(TARGET_ARCH),ppc64le)
- HOST_COMPILER ?= powerpc64le-linux-gnu-g++
- endif
-endif
-HOST_COMPILER ?= g++
-NVCC := $(CUDA_PATH)/bin/nvcc -ccbin $(HOST_COMPILER)
-
-# internal flags
-NVCCFLAGS := -m${TARGET_SIZE}
-CCFLAGS :=
-LDFLAGS :=
-
-# build flags
-
-# Link flag for customized HOST_COMPILER with gcc realpath
-GCC_PATH := $(shell which gcc)
-ifeq ($(CUSTOM_HOST_COMPILER),1)
- ifneq ($(filter /%,$(HOST_COMPILER)),)
- ifneq ($(findstring gcc,$(HOST_COMPILER)),)
- ifneq ($(GCC_PATH),$(HOST_COMPILER))
- LDFLAGS += -lstdc++
- endif
- endif
- endif
-endif
-
-ifeq ($(TARGET_OS),darwin)
- LDFLAGS += -rpath $(CUDA_PATH)/lib
- CCFLAGS += -arch $(HOST_ARCH)
-else ifeq ($(HOST_ARCH)-$(TARGET_ARCH)-$(TARGET_OS),x86_64-armv7l-linux)
- LDFLAGS += --dynamic-linker=/lib/ld-linux-armhf.so.3
- CCFLAGS += -mfloat-abi=hard
-else ifeq ($(TARGET_OS),android)
- LDFLAGS += -pie
- CCFLAGS += -fpie -fpic -fexceptions
-endif
-
-ifneq ($(TARGET_ARCH),$(HOST_ARCH))
- ifeq ($(TARGET_ARCH)-$(TARGET_OS),armv7l-linux)
- ifneq ($(TARGET_FS),)
- GCCVERSIONLTEQ46 := $(shell expr `$(HOST_COMPILER) -dumpversion` \<= 4.6)
- ifeq ($(GCCVERSIONLTEQ46),1)
- CCFLAGS += --sysroot=$(TARGET_FS)
- endif
- LDFLAGS += --sysroot=$(TARGET_FS)
- LDFLAGS += -rpath-link=$(TARGET_FS)/lib
- LDFLAGS += -rpath-link=$(TARGET_FS)/usr/lib
- LDFLAGS += -rpath-link=$(TARGET_FS)/usr/lib/arm-linux-gnueabihf
- endif
- endif
- ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-linux)
- ifneq ($(TARGET_FS),)
- GCCVERSIONLTEQ46 := $(shell expr `$(HOST_COMPILER) -dumpversion` \<= 4.6)
- ifeq ($(GCCVERSIONLTEQ46),1)
- CCFLAGS += --sysroot=$(TARGET_FS)
- endif
- LDFLAGS += --sysroot=$(TARGET_FS)
- LDFLAGS += -rpath-link=$(TARGET_FS)/lib -L$(TARGET_FS)/lib
- LDFLAGS += -rpath-link=$(TARGET_FS)/lib/aarch64-linux-gnu -L$(TARGET_FS)/lib/aarch64-linux-gnu
- LDFLAGS += -rpath-link=$(TARGET_FS)/usr/lib -L$(TARGET_FS)/usr/lib
- LDFLAGS += -rpath-link=$(TARGET_FS)/usr/lib/aarch64-linux-gnu -L$(TARGET_FS)/usr/lib/aarch64-linux-gnu
- LDFLAGS += --unresolved-symbols=ignore-in-shared-libs
- CCFLAGS += -isystem=$(TARGET_FS)/usr/include -I$(TARGET_FS)/usr/include -I$(TARGET_FS)/usr/include/libdrm
- CCFLAGS += -isystem=$(TARGET_FS)/usr/include/aarch64-linux-gnu -I$(TARGET_FS)/usr/include/aarch64-linux-gnu
- endif
- endif
- ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-qnx)
- NVCCFLAGS += -D_QNX_SOURCE
- NVCCFLAGS += --qpp-config 8.3.0,gcc_ntoaarch64le
- CCFLAGS += -DWIN_INTERFACE_CUSTOM -I/usr/include/aarch64-qnx-gnu
- LDFLAGS += -lsocket
- LDFLAGS += -L/usr/lib/aarch64-qnx-gnu
- CCFLAGS += "-Wl\,-rpath-link\,/usr/lib/aarch64-qnx-gnu"
- ifdef TARGET_OVERRIDE
- LDFLAGS += -lslog2
- endif
-
- ifneq ($(TARGET_FS),)
- LDFLAGS += -L$(TARGET_FS)/usr/lib
- CCFLAGS += "-Wl\,-rpath-link\,$(TARGET_FS)/usr/lib"
- LDFLAGS += -L$(TARGET_FS)/usr/libnvidia
- CCFLAGS += "-Wl\,-rpath-link\,$(TARGET_FS)/usr/libnvidia"
- CCFLAGS += -I$(TARGET_FS)/../include
- endif
- endif
-endif
-
-ifdef TARGET_OVERRIDE # cuda toolkit targets override
- NVCCFLAGS += -target-dir $(TARGET_OVERRIDE)
-endif
-
-# Install directory of different arch
-CUDA_INSTALL_TARGET_DIR :=
-ifeq ($(TARGET_ARCH)-$(TARGET_OS),armv7l-linux)
- CUDA_INSTALL_TARGET_DIR = targets/armv7-linux-gnueabihf/
-else ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-linux)
- CUDA_INSTALL_TARGET_DIR = targets/aarch64-linux/
-else ifeq ($(TARGET_ARCH)-$(TARGET_OS),sbsa-linux)
- CUDA_INSTALL_TARGET_DIR = targets/sbsa-linux/
-else ifeq ($(TARGET_ARCH)-$(TARGET_OS),armv7l-android)
- CUDA_INSTALL_TARGET_DIR = targets/armv7-linux-androideabi/
-else ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-android)
- CUDA_INSTALL_TARGET_DIR = targets/aarch64-linux-androideabi/
-else ifeq ($(TARGET_ARCH)-$(TARGET_OS),armv7l-qnx)
- CUDA_INSTALL_TARGET_DIR = targets/ARMv7-linux-QNX/
-else ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-qnx)
- CUDA_INSTALL_TARGET_DIR = targets/aarch64-qnx/
-else ifeq ($(TARGET_ARCH),ppc64le)
- CUDA_INSTALL_TARGET_DIR = targets/ppc64le-linux/
-endif
-
-# Debug build flags
-ifeq ($(dbg),1)
- NVCCFLAGS += -g -G
- BUILD_TYPE := debug
-else
- BUILD_TYPE := release
-endif
-
-ALL_CCFLAGS :=
-ALL_CCFLAGS += $(NVCCFLAGS)
-ALL_CCFLAGS += $(EXTRA_NVCCFLAGS)
-ALL_CCFLAGS += $(addprefix -Xcompiler ,$(CCFLAGS))
-ALL_CCFLAGS += $(addprefix -Xcompiler ,$(EXTRA_CCFLAGS))
-
-UBUNTU = $(shell lsb_release -i -s 2>/dev/null | grep -i ubuntu)
-
-SAMPLE_ENABLED := 1
-
-# This sample is not supported on ARMv7
-ifeq ($(TARGET_ARCH),armv7l)
- $(info >>> WARNING - simpleVoteIntrinsics_nvrtc is not supported on ARMv7 - waiving sample <<<)
- SAMPLE_ENABLED := 0
-endif
-
-ALL_LDFLAGS :=
-ALL_LDFLAGS += $(ALL_CCFLAGS)
-ALL_LDFLAGS += $(addprefix -Xlinker ,$(LDFLAGS))
-ALL_LDFLAGS += $(addprefix -Xlinker ,$(EXTRA_LDFLAGS))
-
-# Common includes and paths for CUDA
-INCLUDES := -I../../../Common
-LIBRARIES :=
-
-################################################################################
-
-# libNVRTC specific libraries
-ifeq ($(TARGET_OS),darwin)
- LDFLAGS += -L$(CUDA_PATH)/lib -F/Library/Frameworks -framework CUDA
-endif
-
-ifeq ($(TARGET_OS),darwin)
- ALL_LDFLAGS += -Xcompiler -F/Library/Frameworks -Xlinker -framework -Xlinker CUDA
-else
- ifeq ($(TARGET_ARCH),x86_64)
- CUDA_SEARCH_PATH ?= $(CUDA_PATH)/lib64/stubs
- CUDA_SEARCH_PATH += $(CUDA_PATH)/lib/stubs
- CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/x86_64-linux/lib/stubs
- endif
-
- ifeq ($(TARGET_ARCH)-$(TARGET_OS),armv7l-linux)
- CUDA_SEARCH_PATH ?= $(CUDA_PATH)/targets/armv7-linux-gnueabihf/lib/stubs
- endif
-
- ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-linux)
- CUDA_SEARCH_PATH ?= $(CUDA_PATH)/targets/aarch64-linux/lib/stubs
- endif
-
- ifeq ($(TARGET_ARCH)-$(TARGET_OS),sbsa-linux)
- CUDA_SEARCH_PATH ?= $(CUDA_PATH)/targets/sbsa-linux/lib/stubs
- endif
-
- ifeq ($(TARGET_ARCH)-$(TARGET_OS),armv7l-android)
- CUDA_SEARCH_PATH ?= $(CUDA_PATH)/targets/armv7-linux-androideabi/lib/stubs
- endif
-
- ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-android)
- CUDA_SEARCH_PATH ?= $(CUDA_PATH)/targets/aarch64-linux-androideabi/lib/stubs
- endif
-
- ifeq ($(TARGET_ARCH)-$(TARGET_OS),armv7l-qnx)
- CUDA_SEARCH_PATH ?= $(CUDA_PATH)/targets/ARMv7-linux-QNX/lib/stubs
- endif
-
- ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-qnx)
- CUDA_SEARCH_PATH ?= $(CUDA_PATH)/targets/aarch64-qnx/lib/stubs
- ifdef TARGET_OVERRIDE
- CUDA_SEARCH_PATH := $(CUDA_PATH)/targets/$(TARGET_OVERRIDE)/lib/stubs
- endif
- endif
-
- ifeq ($(TARGET_ARCH),ppc64le)
- CUDA_SEARCH_PATH ?= $(CUDA_PATH)/targets/ppc64le-linux/lib/stubs
- endif
-
- ifeq ($(HOST_ARCH),ppc64le)
- CUDA_SEARCH_PATH += $(CUDA_PATH)/lib64/stubs
- endif
-
- CUDALIB ?= $(shell find -L $(CUDA_SEARCH_PATH) -maxdepth 1 -name libcuda.so 2> /dev/null)
- ifeq ("$(CUDALIB)","")
- $(info >>> WARNING - libcuda.so not found, CUDA Driver is not installed. Please re-install the driver. <<<)
- SAMPLE_ENABLED := 0
- else
- CUDALIB := $(shell echo $(CUDALIB) | sed "s/ .*//" | sed "s/\/libcuda.so//" )
- LIBRARIES += -L$(CUDALIB) -lcuda
- endif
-endif
-
-ALL_CCFLAGS += --threads 0 --std=c++11
-
-INCLUDES += -I$(CUDA_PATH)/include
-
-LIBRARIES += -lnvrtc
-
-ifeq ($(SAMPLE_ENABLED),0)
-EXEC ?= @echo "[@]"
-endif
-
-################################################################################
-
-# Target rules
-all: build
-
-build: simpleVoteIntrinsics_nvrtc
-
-check.deps:
-ifeq ($(SAMPLE_ENABLED),0)
- @echo "Sample will be waived due to the above missing dependencies"
-else
- @echo "Sample is ready - all dependencies have been met"
-endif
-
-simpleVoteIntrinsics.o:simpleVoteIntrinsics.cpp
- $(EXEC) $(NVCC) $(INCLUDES) $(ALL_CCFLAGS) $(GENCODE_FLAGS) -o $@ -c $<
-
-simpleVoteIntrinsics_nvrtc: simpleVoteIntrinsics.o
- $(EXEC) $(NVCC) $(ALL_LDFLAGS) $(GENCODE_FLAGS) -o $@ $+ $(LIBRARIES)
- $(EXEC) mkdir -p ../../../bin/$(TARGET_ARCH)/$(TARGET_OS)/$(BUILD_TYPE)
- $(EXEC) cp $@ ../../../bin/$(TARGET_ARCH)/$(TARGET_OS)/$(BUILD_TYPE)
-
-run: build
- $(EXEC) ./simpleVoteIntrinsics_nvrtc
-
-testrun: build
-
-clean:
- rm -f simpleVoteIntrinsics_nvrtc simpleVoteIntrinsics.o
- rm -rf ../../../bin/$(TARGET_ARCH)/$(TARGET_OS)/$(BUILD_TYPE)/simpleVoteIntrinsics_nvrtc
-
-clobber: clean
diff --git a/Samples/0_Introduction/simpleVoteIntrinsics_nvrtc/README.md b/Samples/0_Introduction/simpleVoteIntrinsics_nvrtc/README.md
deleted file mode 100644
index a3487a98..00000000
--- a/Samples/0_Introduction/simpleVoteIntrinsics_nvrtc/README.md
+++ /dev/null
@@ -1,74 +0,0 @@
-# simpleVoteIntrinsics_nvrtc - Simple Vote Intrinsics with libNVRTC
-
-## Description
-
-Simple program which demonstrates how to use the Vote (any, all) intrinsic instruction in a CUDA kernel with runtime compilation using NVRTC APIs. Requires Compute Capability 2.0 or higher.
-
-## Key Concepts
-
-Vote Intrinsics, CUDA Driver API, Runtime Compilation
-
-## Supported SM Architectures
-
-[SM 5.0 ](https://developer.nvidia.com/cuda-gpus) [SM 5.2 ](https://developer.nvidia.com/cuda-gpus) [SM 5.3 ](https://developer.nvidia.com/cuda-gpus) [SM 6.0 ](https://developer.nvidia.com/cuda-gpus) [SM 6.1 ](https://developer.nvidia.com/cuda-gpus) [SM 7.0 ](https://developer.nvidia.com/cuda-gpus) [SM 7.2 ](https://developer.nvidia.com/cuda-gpus) [SM 7.5 ](https://developer.nvidia.com/cuda-gpus) [SM 8.0 ](https://developer.nvidia.com/cuda-gpus) [SM 8.6 ](https://developer.nvidia.com/cuda-gpus) [SM 8.7 ](https://developer.nvidia.com/cuda-gpus) [SM 8.9 ](https://developer.nvidia.com/cuda-gpus) [SM 9.0 ](https://developer.nvidia.com/cuda-gpus)
-
-## Supported OSes
-
-Linux, Windows, QNX
-
-## Supported CPU Architecture
-
-x86_64, ppc64le, aarch64
-
-## CUDA APIs involved
-
-### [CUDA Driver API](http://docs.nvidia.com/cuda/cuda-driver-api/index.html)
-cuMemcpyDtoH, cuLaunchKernel, cuMemcpyHtoD, cuCtxSynchronize, cuMemAlloc, cuMemFree, cuModuleGetFunction
-
-## Dependencies needed to build/run
-[NVRTC](../../../README.md#nvrtc)
-
-## Prerequisites
-
-Download and install the [CUDA Toolkit 12.5](https://developer.nvidia.com/cuda-downloads) for your corresponding platform.
-Make sure the dependencies mentioned in [Dependencies]() section above are installed.
-
-## Build and Run
-
-### Windows
-The Windows samples are built using the Visual Studio IDE. Solution files (.sln) are provided for each supported version of Visual Studio, using the format:
-```
-*_vs.sln - for Visual Studio
-```
-Each individual sample has its own set of solution files in its directory:
-
-To build/examine all the samples at once, the complete solution files should be used. To build/examine a single sample, the individual sample solution files should be used.
-> **Note:** Some samples require that the Microsoft DirectX SDK (June 2010 or newer) be installed and that the VC++ directory paths are properly set up (**Tools > Options...**). Check DirectX Dependencies section for details."
-
-### Linux
-The Linux samples are built using makefiles. To use the makefiles, change the current directory to the sample directory you wish to build, and run make:
-```
-$ cd
-$ make
-```
-The samples makefiles can take advantage of certain options:
-* **TARGET_ARCH=** - cross-compile targeting a specific architecture. Allowed architectures are x86_64, ppc64le, aarch64.
- By default, TARGET_ARCH is set to HOST_ARCH. On a x86_64 machine, not setting TARGET_ARCH is the equivalent of setting TARGET_ARCH=x86_64.
-`$ make TARGET_ARCH=x86_64`
`$ make TARGET_ARCH=ppc64le`
`$ make TARGET_ARCH=aarch64`
- See [here](http://docs.nvidia.com/cuda/cuda-samples/index.html#cross-samples) for more details.
-* **dbg=1** - build with debug symbols
- ```
- $ make dbg=1
- ```
-* **SMS="A B ..."** - override the SM architectures for which the sample will be built, where `"A B ..."` is a space-delimited list of SM architectures. For example, to generate SASS for SM 50 and SM 60, use `SMS="50 60"`.
- ```
- $ make SMS="50 60"
- ```
-
-* **HOST_COMPILER=** - override the default g++ host compiler. See the [Linux Installation Guide](http://docs.nvidia.com/cuda/cuda-installation-guide-linux/index.html#system-requirements) for a list of supported host compilers.
-```
- $ make HOST_COMPILER=g++
-```
-
-## References (for more details)
-
diff --git a/Samples/0_Introduction/simpleVoteIntrinsics_nvrtc/simpleVoteIntrinsics.cpp b/Samples/0_Introduction/simpleVoteIntrinsics_nvrtc/simpleVoteIntrinsics.cpp
deleted file mode 100644
index 10f05ad4..00000000
--- a/Samples/0_Introduction/simpleVoteIntrinsics_nvrtc/simpleVoteIntrinsics.cpp
+++ /dev/null
@@ -1,341 +0,0 @@
-/* Copyright (c) 2022, NVIDIA CORPORATION. All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- * * Redistributions of source code must retain the above copyright
- * notice, this list of conditions and the following disclaimer.
- * * Redistributions in binary form must reproduce the above copyright
- * notice, this list of conditions and the following disclaimer in the
- * documentation and/or other materials provided with the distribution.
- * * Neither the name of NVIDIA CORPORATION nor the names of its
- * contributors may be used to endorse or promote products derived
- * from this software without specific prior written permission.
- *
- * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
- * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
- * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
- * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
- * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
- * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
- * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
- * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
- * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-// System includes
-#include
-#include
-
-// CUDA runtime
-#include
-#include "nvrtc_helper.h"
-
-// helper functions and utilities to work with CUDA
-#include
-
-#ifndef MAX
-#define MAX(a, b) (a > b ? a : b)
-#endif
-
-static const char *sSDKsample = "[simpleVoteIntrinsics_nvrtc]\0";
-
-////////////////////////////////////////////////////////////////////////////////
-// Global types and parameters
-////////////////////////////////////////////////////////////////////////////////
-
-#define VOTE_DATA_GROUP 4
-
-////////////////////////////////////////////////////////////////////////////////
-// CUDA Voting Kernel functions
-////////////////////////////////////////////////////////////////////////////////
-
-// Generate the test pattern for Tests 1 and 2
-void genVoteTestPattern(unsigned int *VOTE_PATTERN, int size) {
- // For testing VOTE.Any (all of these threads will return 0)
- for (int i = 0; i < size / 4; i++) {
- VOTE_PATTERN[i] = 0x00000000;
- }
-
- // For testing VOTE.Any (1/2 these threads will return 1)
- for (int i = 2 * size / 8; i < 4 * size / 8; i++) {
- VOTE_PATTERN[i] = (i & 0x01) ? i : 0;
- }
-
- // For testing VOTE.all (1/2 of these threads will return 0)
- for (int i = 2 * size / 4; i < 3 * size / 4; i++) {
- VOTE_PATTERN[i] = (i & 0x01) ? 0 : i;
- }
-
- // For testing VOTE.all (all of these threads will return 1)
- for (int i = 3 * size / 4; i < 4 * size / 4; i++) {
- VOTE_PATTERN[i] = 0xffffffff;
- }
-}
-
-int checkErrors1(unsigned int *h_result, int start, int end, int warp_size,
- const char *voteType) {
- int i, sum = 0;
-
- for (sum = 0, i = start; i < end; i++) {
- sum += h_result[i];
- }
-
- if (sum > 0) {
- printf("\t<%s>[%d - %d] = ", voteType, start, end - 1);
-
- for (i = start; i < end; i++) {
- printf("%d", h_result[i]);
- }
-
- printf("%d values FAILED\n", sum);
- }
-
- return (sum > 0);
-}
-
-int checkErrors2(unsigned int *h_result, int start, int end, int warp_size,
- const char *voteType) {
- int i, sum = 0;
-
- for (sum = 0, i = start; i < end; i++) {
- sum += h_result[i];
- }
-
- if (sum != warp_size) {
- printf("\t<%s>[%d - %d] = ", voteType, start, end - 1);
-
- for (i = start; i < end; i++) {
- printf("%d", h_result[i]);
- }
-
- printf(" - FAILED\n");
- }
-
- return (sum != warp_size);
-}
-
-// Verification code for Kernel #1
-int checkResultsVoteAnyKernel1(unsigned int *h_result, int size,
- int warp_size) {
- int error_count = 0;
-
- error_count += checkErrors1(h_result, 0, VOTE_DATA_GROUP * warp_size / 4,
- warp_size, "Vote.Any");
- error_count +=
- checkErrors2(h_result, VOTE_DATA_GROUP * warp_size / 4,
- 2 * VOTE_DATA_GROUP * warp_size / 4, warp_size, "Vote.Any");
- error_count +=
- checkErrors2(h_result, 2 * VOTE_DATA_GROUP * warp_size / 4,
- 3 * VOTE_DATA_GROUP * warp_size / 4, warp_size, "Vote.Any");
- error_count +=
- checkErrors2(h_result, 3 * VOTE_DATA_GROUP * warp_size / 4,
- 4 * VOTE_DATA_GROUP * warp_size / 4, warp_size, "Vote.Any");
-
- printf((error_count == 0) ? "\tOK\n" : "\tERROR\n");
-
- return error_count;
-}
-
-// Verification code for Kernel #2
-int checkResultsVoteAllKernel2(unsigned int *h_result, int size,
- int warp_size) {
- int error_count = 0;
-
- error_count += checkErrors1(h_result, 0, VOTE_DATA_GROUP * warp_size / 4,
- warp_size, "Vote.All");
- error_count +=
- checkErrors1(h_result, VOTE_DATA_GROUP * warp_size / 4,
- 2 * VOTE_DATA_GROUP * warp_size / 4, warp_size, "Vote.All");
- error_count +=
- checkErrors1(h_result, 2 * VOTE_DATA_GROUP * warp_size / 4,
- 3 * VOTE_DATA_GROUP * warp_size / 4, warp_size, "Vote.All");
- error_count +=
- checkErrors2(h_result, 3 * VOTE_DATA_GROUP * warp_size / 4,
- 4 * VOTE_DATA_GROUP * warp_size / 4, warp_size, "Vote.All");
-
- printf((error_count == 0) ? "\tOK\n" : "\tERROR\n");
-
- return error_count;
-}
-
-// Verification code for Kernel #3
-int checkResultsVoteAnyKernel3(bool *hinfo, int size) {
- int i, error_count = 0;
-
- for (i = 0; i < size * 3; i++) {
- switch (i % 3) {
- case 0:
- // First warp should be all zeros.
- if (hinfo[i] != (i >= size * 1)) {
- error_count++;
- }
- break;
-
- case 1:
- // First warp and half of second should be all zeros.
- if (hinfo[i] != (i >= size * 3 / 2)) {
- error_count++;
- }
- break;
-
- case 2:
- // First two warps should be all zeros.
- if (hinfo[i] != (i >= size * 2)) {
- error_count++;
- }
- break;
- }
- }
-
- printf((error_count == 0) ? "\tOK\n" : "\tERROR\n");
-
- return error_count;
-}
-
-int main(int argc, char **argv) {
- unsigned int *h_input, *h_result;
- CUdeviceptr d_input, d_result;
-
- char *cubin, *kernel_file;
- size_t cubinSize;
- kernel_file = sdkFindFilePath("simpleVote_kernel.cuh", argv[0]);
- compileFileToCUBIN(kernel_file, argc, argv, &cubin, &cubinSize, 0);
- CUmodule module = loadCUBIN(cubin, argc, argv);
-
- bool *hinfo = NULL;
- CUdeviceptr dinfo;
-
- int error_count[3] = {0, 0, 0};
- int warp_size = 32;
-
- printf("%s\n", sSDKsample);
-
- h_input = (unsigned int *)malloc(VOTE_DATA_GROUP * warp_size *
- sizeof(unsigned int));
- h_result = (unsigned int *)malloc(VOTE_DATA_GROUP * warp_size *
- sizeof(unsigned int));
-
- checkCudaErrors(
- cuMemAlloc(&d_input, VOTE_DATA_GROUP * warp_size * sizeof(unsigned int)));
- checkCudaErrors(cuMemAlloc(
- &d_result, VOTE_DATA_GROUP * warp_size * sizeof(unsigned int)));
-
- genVoteTestPattern(h_input, VOTE_DATA_GROUP * warp_size);
-
- checkCudaErrors(cuMemcpyHtoD(
- d_input, h_input, VOTE_DATA_GROUP * warp_size * sizeof(unsigned int)));
-
- // Start of Vote Any Test Kernel #1
- printf("[VOTE Kernel Test 1/3]\n");
- printf("\tRunning <> kernel1 ...\n");
- {
- dim3 gridBlock(1, 1);
- dim3 threadBlock(VOTE_DATA_GROUP * warp_size, 1);
- CUfunction kernel_addr;
- checkCudaErrors(
- cuModuleGetFunction(&kernel_addr, module, "VoteAnyKernel1"));
-
- int size = VOTE_DATA_GROUP * warp_size;
- void *arr[] = {(void *)&d_input, (void *)&d_result, (void *)&size};
-
- checkCudaErrors(cuLaunchKernel(
- kernel_addr, gridBlock.x, gridBlock.y, gridBlock.z, /* grid dim */
- threadBlock.x, threadBlock.y, threadBlock.z, /* block dim */
- 0, 0, /* shared mem, stream */
- &arr[0], /* arguments */
- 0));
-
- checkCudaErrors(cuCtxSynchronize());
- }
-
- checkCudaErrors(cuMemcpyDtoH(
- h_result, d_result, VOTE_DATA_GROUP * warp_size * sizeof(unsigned int)));
-
- error_count[0] += checkResultsVoteAnyKernel1(
- h_result, VOTE_DATA_GROUP * warp_size, warp_size);
-
- // Start of Vote All Test Kernel #2
-
- printf("\n[VOTE Kernel Test 2/3]\n");
-
- printf("\tRunning <> kernel2 ...\n");
- {
- dim3 gridBlock(1, 1);
- dim3 threadBlock(VOTE_DATA_GROUP * warp_size, 1);
-
- CUfunction kernel_addr;
- checkCudaErrors(
- cuModuleGetFunction(&kernel_addr, module, "VoteAllKernel2"));
-
- int size = VOTE_DATA_GROUP * warp_size;
- void *arr[] = {(void *)&d_input, (void *)&d_result, (void *)&size};
-
- checkCudaErrors(cuLaunchKernel(
- kernel_addr, gridBlock.x, gridBlock.y, gridBlock.z, /* grid dim */
- threadBlock.x, threadBlock.y, threadBlock.z, /* block dim */
- 0, 0, /* shared mem, stream */
- &arr[0], /* arguments */
- 0));
-
- checkCudaErrors(cuCtxSynchronize());
- }
-
- checkCudaErrors(cuMemcpyDtoH(
- h_result, d_result, VOTE_DATA_GROUP * warp_size * sizeof(unsigned int)));
-
- error_count[1] += checkResultsVoteAllKernel2(
- h_result, VOTE_DATA_GROUP * warp_size, warp_size);
-
- // Second Vote Kernel Test #3 (both Any/All)
- hinfo = (bool *)calloc(warp_size * 3 * 3, sizeof(bool));
-
- checkCudaErrors(cuMemAlloc(&dinfo, warp_size * 3 * 3 * sizeof(bool)));
- checkCudaErrors(cuMemcpyHtoD(dinfo, hinfo, warp_size * 3 * 3 * sizeof(bool)));
-
- printf("\n[VOTE Kernel Test 3/3]\n");
- printf("\tRunning <> kernel3 ...\n");
- {
- dim3 gridBlock(1, 1);
- dim3 threadBlock(warp_size * 3, 1);
-
- CUfunction kernel_addr;
-
- checkCudaErrors(
- cuModuleGetFunction(&kernel_addr, module, "VoteAnyKernel3"));
-
- int size = warp_size;
- void *arr[] = {(void *)&dinfo, (void *)&size};
-
- checkCudaErrors(cuLaunchKernel(
- kernel_addr, gridBlock.x, gridBlock.y, gridBlock.z, /* grid dim */
- threadBlock.x, threadBlock.y, threadBlock.z, /* block dim */
- 0, 0, /* shared mem, stream */
- &arr[0], /* arguments */
- 0));
-
- checkCudaErrors(cuCtxSynchronize());
- }
-
- checkCudaErrors(cuMemcpyDtoH(hinfo, dinfo, warp_size * 3 * 3 * sizeof(bool)));
-
- error_count[2] = checkResultsVoteAnyKernel3(hinfo, warp_size * 3);
-
- // Now free these resources for Test #1,2
- checkCudaErrors(cuMemFree(d_input));
- checkCudaErrors(cuMemFree(d_result));
- free(h_input);
- free(h_result);
-
- // Free resources from Test #3
- free(hinfo);
- checkCudaErrors(cuMemFree(dinfo));
-
- printf("\tShutting down...\n");
-
- return (error_count[0] == 0 && error_count[1] == 0 && error_count[2] == 0)
- ? EXIT_SUCCESS
- : EXIT_FAILURE;
-}
diff --git a/Samples/0_Introduction/simpleVoteIntrinsics_nvrtc/simpleVote_kernel.cuh b/Samples/0_Introduction/simpleVoteIntrinsics_nvrtc/simpleVote_kernel.cuh
deleted file mode 100644
index b488d15d..00000000
--- a/Samples/0_Introduction/simpleVoteIntrinsics_nvrtc/simpleVote_kernel.cuh
+++ /dev/null
@@ -1,80 +0,0 @@
-/* Copyright (c) 2022, NVIDIA CORPORATION. All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- * * Redistributions of source code must retain the above copyright
- * notice, this list of conditions and the following disclaimer.
- * * Redistributions in binary form must reproduce the above copyright
- * notice, this list of conditions and the following disclaimer in the
- * documentation and/or other materials provided with the distribution.
- * * Neither the name of NVIDIA CORPORATION nor the names of its
- * contributors may be used to endorse or promote products derived
- * from this software without specific prior written permission.
- *
- * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
- * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
- * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
- * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
- * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
- * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
- * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
- * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
- * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#ifndef SIMPLEVOTE_KERNEL_CU
-#define SIMPLEVOTE_KERNEL_CU
-
-////////////////////////////////////////////////////////////////////////////////
-// Vote Any/All intrinsic kernel function tests are supported only by CUDA
-// capable devices that are CUDA hardware that has SM1.2 or later
-// Vote Functions (refer to section 4.4.5 in the CUDA Programming Guide)
-////////////////////////////////////////////////////////////////////////////////
-
-// Kernel #1 tests the across-the-warp vote(any) intrinsic.
-// If ANY one of the threads (within the warp) of the predicated condition
-// returns a non-zero value, then all threads within this warp will return a
-// non-zero value
-extern "C" __global__ void VoteAnyKernel1(unsigned int *input,
- unsigned int *result, int size) {
- int tx = threadIdx.x;
-
- int mask = 0xffffffff;
- result[tx] = __any_sync(mask, input[tx]);
-}
-
-// Kernel #2 tests the across-the-warp vote(all) intrinsic.
-// If ALL of the threads (within the warp) of the predicated condition returns
-// a non-zero value, then all threads within this warp will return a non-zero
-// value
-extern "C" __global__ void VoteAllKernel2(unsigned int *input,
- unsigned int *result, int size) {
- int tx = threadIdx.x;
-
- int mask = 0xffffffff;
- result[tx] = __all_sync(mask, input[tx]);
-}
-
-// Kernel #3 is a directed test for the across-the-warp vote(all) intrinsic.
-// This kernel will test for conditions across warps, and within half warps
-extern "C" __global__ void VoteAnyKernel3(bool *info, int warp_size) {
- int tx = threadIdx.x;
- unsigned int mask = 0xffffffff;
- bool *offs = info + (tx * 3);
-
- // The following should hold true for the second and third warp
- *offs = __any_sync(mask, (tx >= (warp_size * 3) / 2));
- // The following should hold true for the "upper half" of the second warp,
- // and all of the third warp
- *(offs + 1) = (tx >= (warp_size * 3) / 2 ? true : false);
-
- // The following should hold true for the third warp only
- if (__all_sync(mask, (tx >= (warp_size * 3) / 2))) {
- *(offs + 2) = true;
- }
-}
-
-#endif
diff --git a/Samples/0_Introduction/template/NsightEclipse.xml b/Samples/0_Introduction/template/NsightEclipse.xml
deleted file mode 100644
index ccf76b74..00000000
--- a/Samples/0_Introduction/template/NsightEclipse.xml
+++ /dev/null
@@ -1,72 +0,0 @@
-
-
-
- template
-
- cudaMalloc
- cudaMemcpy
- cudaFree
-
-
- whole
-
- ./
- ../
- ../../../Common
-
-
- Device Memory Allocation
-
-
-
-
-
-
-
- true
- template.cu
-
- 1:CUDA Basic Topics
-
- sm50
- sm52
- sm53
- sm60
- sm61
- sm70
- sm72
- sm75
- sm80
- sm86
- sm87
- sm89
- sm90
-
-
- x86_64
- linux
-
-
- windows7
-
-
- x86_64
- macosx
-
-
- arm
-
-
- sbsa
-
-
- ppc64le
- linux
-
-
-
- all
-
- Template
- exe
-