diff --git a/activity_trace/Makefile b/activity_trace/Makefile index eb402de..604da66 100644 --- a/activity_trace/Makefile +++ b/activity_trace/Makefile @@ -2,7 +2,7 @@ # Copyright 2011-2013 NVIDIA Corporation. All rights reserved # INCLUDES=-I../common -NVCC=/usr/local/cuda-12.9/bin/nvcc +NVCC=$(CUDA_INSTALL_PATH)/bin/nvcc ifndef OS OS := $(shell uname) diff --git a/activity_trace_async/Makefile b/activity_trace_async/Makefile index 3296fad..13d1ee1 100644 --- a/activity_trace_async/Makefile +++ b/activity_trace_async/Makefile @@ -75,10 +75,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 50 52 53 60 61 62 70 72 75 80 86 87 89 90 100 101 120 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk activity_trace_async: activity_trace_async.$(OBJ) $(NVCC) $(NVCC_COMPILER) $(NVCCFLAGS) -o $@ $^ $(LIBS) $(INCLUDES) diff --git a/autorange_profiling/Makefile b/autorange_profiling/Makefile index 98c9bfa..637d3b3 100644 --- a/autorange_profiling/Makefile +++ b/autorange_profiling/Makefile @@ -82,10 +82,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 70 72 75 80 86 87 89 90 100 101 120 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk .DEFAULT: all .PHONY: all diff --git a/autorange_profiling/auto_range_profiling.cu b/autorange_profiling/auto_range_profiling.cu index 74c93c0..3093f71 100644 --- a/autorange_profiling/auto_range_profiling.cu +++ b/autorange_profiling/auto_range_profiling.cu @@ -426,7 +426,7 @@ main( } CUcontext cuContext; - DRIVER_API_CALL(cuCtxCreate(&cuContext, 0, cuDevice)); + DRIVER_API_CALL(COMPAT_cuCtxCreate(&cuContext, 0, cuDevice)); // Get chip name for the CUDA device. CUpti_Device_GetChipName_Params getChipNameParams = { CUpti_Device_GetChipName_Params_STRUCT_SIZE }; diff --git a/autorange_profiling/simplecuda.cu b/autorange_profiling/simplecuda.cu index e2f35f2..dbf95ac 100644 --- a/autorange_profiling/simplecuda.cu +++ b/autorange_profiling/simplecuda.cu @@ -217,7 +217,7 @@ bool runTest(int deviceNum, DRIVER_API_CALL(cuDeviceGet(&cuDevice, deviceNum)); CUcontext cuContext; - DRIVER_API_CALL(cuCtxCreate(&cuContext, 0, cuDevice)); + DRIVER_API_CALL(COMPAT_cuCtxCreate(&cuContext, 0, cuDevice)); CUpti_Profiler_BeginSession_Params beginSessionParams = {CUpti_Profiler_BeginSession_Params_STRUCT_SIZE}; CUpti_Profiler_SetConfig_Params setConfigParams = {CUpti_Profiler_SetConfig_Params_STRUCT_SIZE}; diff --git a/callback_event/Makefile b/callback_event/Makefile index 2c51ff8..f836f81 100644 --- a/callback_event/Makefile +++ b/callback_event/Makefile @@ -76,10 +76,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 50 52 53 60 61 62 70 72 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk callback_event: callback_event.$(OBJ) $(NVCC) $(NVCC_COMPILER) $(NVCCFLAGS) -o $@ callback_event.$(OBJ) $(LIBS) $(INCLUDES) diff --git a/callback_event/callback_event.cu b/callback_event/callback_event.cu index 0ceef66..a2cbd6d 100644 --- a/callback_event/callback_event.cu +++ b/callback_event/callback_event.cu @@ -225,7 +225,7 @@ main( exit(EXIT_WAIVED); } - DRIVER_API_CALL(cuCtxCreate(&context, 0, dev)); + DRIVER_API_CALL(COMPAT_cuCtxCreate(&context, 0, dev)); // Creating event group for profiling CUPTI_API_CALL(cuptiEventGroupCreate(context, &cuptiEvent.eventGroup, 0)); diff --git a/callback_metric/Makefile b/callback_metric/Makefile index f929f50..cbeeb30 100644 --- a/callback_metric/Makefile +++ b/callback_metric/Makefile @@ -76,10 +76,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 50 52 53 60 61 62 70 72 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk callback_metric: callback_metric.$(OBJ) $(NVCC) $(NVCC_COMPILER) $(NVCCFLAGS) -o $@ callback_metric.$(OBJ) $(LIBS) $(INCLUDES) diff --git a/callback_metric/callback_metric.cu b/callback_metric/callback_metric.cu index 476f023..784ded3 100644 --- a/callback_metric/callback_metric.cu +++ b/callback_metric/callback_metric.cu @@ -372,7 +372,7 @@ main( exit(EXIT_WAIVED); } - DRIVER_API_CALL(cuCtxCreate(&context, 0, device)); + DRIVER_API_CALL(COMPAT_cuCtxCreate(&context, 0, device)); // Get the name of the metric to collect if (argc > 2) diff --git a/callback_profiling/Makefile b/callback_profiling/Makefile index 6c1a3e7..ae8670b 100644 --- a/callback_profiling/Makefile +++ b/callback_profiling/Makefile @@ -82,10 +82,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 70 72 75 80 86 87 89 90 100 101 120 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk .DEFAULT: all .PHONY: all diff --git a/callback_timestamp/Makefile b/callback_timestamp/Makefile index 46b3d04..c0523be 100644 --- a/callback_timestamp/Makefile +++ b/callback_timestamp/Makefile @@ -76,10 +76,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 50 52 53 60 61 62 70 72 75 80 86 87 89 90 100 101 120 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk callback_timestamp: callback_timestamp.$(OBJ) diff --git a/callback_timestamp/callback_timestamp.cu b/callback_timestamp/callback_timestamp.cu index 34500c9..12fb7b0 100644 --- a/callback_timestamp/callback_timestamp.cu +++ b/callback_timestamp/callback_timestamp.cu @@ -282,7 +282,7 @@ main( DRIVER_API_CALL(cuInit(0)); - DRIVER_API_CALL(cuCtxCreate(&context, 0, device)); + DRIVER_API_CALL(COMPAT_cuCtxCreate(&context, 0, device)); // Allocate input vectors pHostA and pHostB in host memory. pHostA = (int *)malloc(size); diff --git a/checkpoint_kernels/Makefile b/checkpoint_kernels/Makefile index f6218ab..76fe3b4 100644 --- a/checkpoint_kernels/Makefile +++ b/checkpoint_kernels/Makefile @@ -81,10 +81,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 50 52 53 60 61 62 70 72 75 80 86 87 89 90 100 101 120 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk checkpoint_kernels: checkpoint_kernels.$(OBJ) $(NVCC) $(NVCC_COMPILER) $(NVCCFLAGS) -o $@ $^ $(LIBS) $(INCLUDES) diff --git a/checkpoint_kernels/checkpoint_kernels.cu b/checkpoint_kernels/checkpoint_kernels.cu index f879758..05c1cbb 100644 --- a/checkpoint_kernels/checkpoint_kernels.cu +++ b/checkpoint_kernels/checkpoint_kernels.cu @@ -69,7 +69,7 @@ main() // Set up a context for device 0. RUNTIME_API_CALL(cudaSetDevice(0)); - DRIVER_API_CALL(cuCtxCreate(&ctx, 0, 0)); + DRIVER_API_CALL(COMPAT_cuCtxCreate(&ctx, 0, 0)); // Allocate host and device arrays and initialize to known values. float *pDeviceA; diff --git a/common/common.mk b/common/common.mk new file mode 100644 index 0000000..068601d --- /dev/null +++ b/common/common.mk @@ -0,0 +1,23 @@ +# Common Makefile configuration for CUPTI samples +# Copyright NVIDIA Corporation + +# Detect CUDA version to determine supported SM architectures +# CUDA 13+ removed support for compute capabilities below sm_75 +CUDA_VERSION := $(shell $(CUDA_INSTALL_PATH)/bin/nvcc --version 2>/dev/null | grep release | sed 's/.*release //' | sed 's/,.*//' | cut -d. -f1) + +# Set default SMS based on CUDA version if not already set +ifndef SMS + ifeq ($(shell test $(CUDA_VERSION) -ge 13 2>/dev/null; echo $$?),0) + # CUDA 13+ only supports sm_75 and above + # Note: CUDA 13 uses sm_100, sm_103 instead of sm_101, sm_102 + SMS := 75 80 86 87 88 89 90 100 103 110 120 + DEFAULT_SM_ARCH := sm_75 + else + # CUDA 12 and below support older architectures + SMS := 50 52 53 60 61 62 70 72 75 80 86 87 89 90 100 101 120 + DEFAULT_SM_ARCH := sm_52 + endif +endif + +# Generate SASS code for each SM architecture listed in $(SMS) +$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) diff --git a/common/helper_cupti.h b/common/helper_cupti.h index 2e292b2..447da34 100644 --- a/common/helper_cupti.h +++ b/common/helper_cupti.h @@ -49,6 +49,16 @@ do } while (0) #endif +// Compatibility wrapper for cuCtxCreate API change in CUDA 13 +// CUDA 13+ requires an additional CUctxCreateParams parameter +#if CUDA_VERSION >= 13000 + #define COMPAT_cuCtxCreate(pctx, flags, dev) \ + cuCtxCreate(pctx, NULL, flags, dev) +#else + #define COMPAT_cuCtxCreate(pctx, flags, dev) \ + cuCtxCreate(pctx, flags, dev) +#endif + #ifndef RUNTIME_API_CALL #define RUNTIME_API_CALL(apiFunctionCall) \ do \ diff --git a/concurrent_profiling/Makefile b/concurrent_profiling/Makefile index 4fda3b3..21c02f6 100644 --- a/concurrent_profiling/Makefile +++ b/concurrent_profiling/Makefile @@ -82,10 +82,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 70 72 75 80 86 87 89 90 100 101 120 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk .DEFAULT: all .PHONY: all diff --git a/concurrent_profiling/concurrent_profiling.cu b/concurrent_profiling/concurrent_profiling.cu index 962975d..29d2a35 100644 --- a/concurrent_profiling/concurrent_profiling.cu +++ b/concurrent_profiling/concurrent_profiling.cu @@ -614,7 +614,7 @@ main( // CUPTI_KernelReplay, CUPTI_UserReplay, or CUPTI_ApplicationReplay. config.replayMode = CUPTI_UserReplay; // Either set to a context, or may be NULL if a default context has been created. - DRIVER_API_CALL(cuCtxCreate(&(config.context), 0, device)); + DRIVER_API_CALL(COMPAT_cuCtxCreate(&(config.context), 0, device)); // Save this device config. deviceData[device].config = config; diff --git a/cuda_graphs_trace/Makefile b/cuda_graphs_trace/Makefile index 30f0b91..c18a8c7 100644 --- a/cuda_graphs_trace/Makefile +++ b/cuda_graphs_trace/Makefile @@ -76,10 +76,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 50 52 53 60 61 62 70 72 75 80 86 87 89 90 100 101 120 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk cuda_graphs_trace: cuda_graphs_trace.$(OBJ) $(NVCC) $(NVCC_COMPILER) $(NVCCFLAGS) -o $@ $^ $(LIBS) $(INCLUDES) diff --git a/cuda_memory_trace/Makefile b/cuda_memory_trace/Makefile index 90e4756..70d85c2 100644 --- a/cuda_memory_trace/Makefile +++ b/cuda_memory_trace/Makefile @@ -75,10 +75,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 50 52 53 60 61 62 70 72 75 80 86 87 89 90 100 101 120 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk cuda_memory_trace: memory_trace.$(OBJ) $(NVCC) $(NVCC_COMPILER) $(NVCCFLAGS) -o $@ $^ $(LIBS) $(INCLUDES) diff --git a/cupti_correlation/Makefile b/cupti_correlation/Makefile index c71814f..dce80f0 100644 --- a/cupti_correlation/Makefile +++ b/cupti_correlation/Makefile @@ -75,10 +75,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 50 52 53 60 61 62 70 72 75 80 86 87 89 90 100 101 120 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk cupti_correlation: cupti_correlation.$(OBJ) $(NVCC) $(NVCC_COMPILER) $(NVCCFLAGS) -o $@ $^ $(LIBS) $(INCLUDES) diff --git a/cupti_external_correlation/Makefile b/cupti_external_correlation/Makefile index 349cd65..0f00657 100644 --- a/cupti_external_correlation/Makefile +++ b/cupti_external_correlation/Makefile @@ -75,10 +75,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 50 52 53 60 61 62 70 72 75 80 86 87 89 90 100 101 120 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk cupti_external_correlation: cupti_external_correlation.$(OBJ) $(NVCC) $(NVCC_COMPILER) $(NVCCFLAGS) -o $@ $^ $(LIBS) $(INCLUDES) diff --git a/cupti_external_correlation/cupti_external_correlation.cu b/cupti_external_correlation/cupti_external_correlation.cu index e804fbd..3d648d1 100644 --- a/cupti_external_correlation/cupti_external_correlation.cu +++ b/cupti_external_correlation/cupti_external_correlation.cu @@ -97,7 +97,7 @@ DoVectorAddition() { // Push external id for the initialization: memory allocation and memcpy operations from host to device. CUPTI_API_CALL_VERBOSE(cuptiActivityPushExternalCorrelationId(CUPTI_EXTERNAL_CORRELATION_KIND_UNKNOWN, static_cast(INITIALIZATION_EXTERNAL_ID))); - DRIVER_API_CALL(cuCtxCreate(&context, 0, device)); + DRIVER_API_CALL(COMPAT_cuCtxCreate(&context, 0, device)); // Allocate vectors in device memory. RUNTIME_API_CALL(cudaMalloc((void **)&pDeviceA, size)); diff --git a/cupti_metric_properties/Makefile b/cupti_metric_properties/Makefile index f8108bf..d356539 100644 --- a/cupti_metric_properties/Makefile +++ b/cupti_metric_properties/Makefile @@ -82,10 +82,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 70 72 75 80 86 87 89 90 100 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk .DEFAULT: all .PHONY: all diff --git a/cupti_metric_properties/cupti_metric_properties.cpp b/cupti_metric_properties/cupti_metric_properties.cpp index e790728..cb05acc 100644 --- a/cupti_metric_properties/cupti_metric_properties.cpp +++ b/cupti_metric_properties/cupti_metric_properties.cpp @@ -454,7 +454,7 @@ CUptiResult GetCounterAvailabilityImage(std::vector& counterAvailabilit { DRIVER_API_CALL(cuInit(0)); CUcontext ctx; - DRIVER_API_CALL(cuCtxCreate(&ctx, 0, 0)); + DRIVER_API_CALL(COMPAT_cuCtxCreate(&ctx, 0, 0)); CUpti_Profiler_GetCounterAvailability_Params getCounterAvailabilityParams = { CUpti_Profiler_GetCounterAvailability_Params_STRUCT_SIZE }; getCounterAvailabilityParams.ctx = ctx; diff --git a/cupti_nvtx/Makefile b/cupti_nvtx/Makefile index b2d1455..30656a2 100644 --- a/cupti_nvtx/Makefile +++ b/cupti_nvtx/Makefile @@ -77,10 +77,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 50 52 53 60 61 62 70 72 75 80 86 87 89 90 100 101 120 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk cupti_nvtx: cupti_nvtx.$(OBJ) $(NVCC) $(NVCC_COMPILER) $(NVCCFLAGS) -o $@ $^ $(LIBS) $(INCLUDES) diff --git a/cupti_nvtx/cupti_nvtx.cu b/cupti_nvtx/cupti_nvtx.cu index 27847ab..286f000 100644 --- a/cupti_nvtx/cupti_nvtx.cu +++ b/cupti_nvtx/cupti_nvtx.cu @@ -114,7 +114,7 @@ DoVectorAddition() InitializeVector(pHostB, N); memset(pHostC, 0, size); - DRIVER_API_CALL(cuCtxCreate(&context, 0, device)); + DRIVER_API_CALL(COMPAT_cuCtxCreate(&context, 0, device)); nvtxNameCuContextA(context, "CUDA Context"); // Push range "Allocate device memory" on domain "Vector Addition". diff --git a/cupti_nvtx_ext_payload/Makefile b/cupti_nvtx_ext_payload/Makefile index 8a25335..a492051 100644 --- a/cupti_nvtx_ext_payload/Makefile +++ b/cupti_nvtx_ext_payload/Makefile @@ -77,10 +77,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 75 80 86 87 89 90 100 103 110 120 121 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk cupti_nvtx_ext_payload: cupti_nvtx_ext_payload.$(OBJ) nvtx_payload_attributes.$(OBJ) nvtx_payload_parser.$(OBJ) $(NVCC) $(NVCC_COMPILER) $(NVCCFLAGS) -o $@ $^ $(LIBS) $(INCLUDES) diff --git a/event_multi_gpu/event_multi_gpu.cu b/event_multi_gpu/event_multi_gpu.cu index 5b7d30c..5722d8d 100644 --- a/event_multi_gpu/event_multi_gpu.cu +++ b/event_multi_gpu/event_multi_gpu.cu @@ -103,7 +103,7 @@ main( { RUNTIME_API_CALL(cudaSetDevice(i)); - DRIVER_API_CALL(cuCtxCreate(&(context[i]), 0, device[i])); + DRIVER_API_CALL(COMPAT_cuCtxCreate(&(context[i]), 0, device[i])); DRIVER_API_CALL(cuCtxPopCurrent(&(context[i]))); } diff --git a/event_sampling/Makefile b/event_sampling/Makefile index 2660824..c424770 100644 --- a/event_sampling/Makefile +++ b/event_sampling/Makefile @@ -76,10 +76,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 50 52 53 60 61 62 70 72 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk event_sampling: event_sampling.$(OBJ) $(NVCC) $(NVCC_COMPILER) $(NVCCFLAGS) -o $@ event_sampling.$(OBJ) $(LIBS) $(INCLUDES) diff --git a/event_sampling/event_sampling.cu b/event_sampling/event_sampling.cu index 6eacf62..43d2ecd 100644 --- a/event_sampling/event_sampling.cu +++ b/event_sampling/event_sampling.cu @@ -311,7 +311,7 @@ main(int argc, char *argv[]) s_pEventName = EVENT_NAME; } - DRIVER_API_CALL(cuCtxCreate(&context, 0, device)); + DRIVER_API_CALL(COMPAT_cuCtxCreate(&context, 0, device)); // Create semaphore. #ifdef _WIN32 diff --git a/nested_range_profiling/Makefile b/nested_range_profiling/Makefile index 2fb510e..a143622 100644 --- a/nested_range_profiling/Makefile +++ b/nested_range_profiling/Makefile @@ -82,10 +82,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 70 72 75 80 86 87 89 90 100 101 120 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk .DEFAULT: all .PHONY: all diff --git a/nested_range_profiling/nested_range_profiling.cu b/nested_range_profiling/nested_range_profiling.cu index df5053e..17dc91a 100644 --- a/nested_range_profiling/nested_range_profiling.cu +++ b/nested_range_profiling/nested_range_profiling.cu @@ -448,7 +448,7 @@ int main( } CUcontext cuContext; - DRIVER_API_CALL(cuCtxCreate(&cuContext, 0, cuDevice)); + DRIVER_API_CALL(COMPAT_cuCtxCreate(&cuContext, 0, cuDevice)); // Get chip name for the cuda device. CUpti_Device_GetChipName_Params getChipNameParams = { CUpti_Device_GetChipName_Params_STRUCT_SIZE }; diff --git a/nvlink_bandwidth/Makefile b/nvlink_bandwidth/Makefile index b5c8023..4bc62d8 100644 --- a/nvlink_bandwidth/Makefile +++ b/nvlink_bandwidth/Makefile @@ -76,10 +76,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 60 70 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk nvlink_bandwidth: nvlink_bandwidth.$(OBJ) $(NVCC) $(NVCC_COMPILER) $(NVCCFLAGS) -o $@ nvlink_bandwidth.$(OBJ) $(LIBS) $(INCLUDES) diff --git a/nvlink_bandwidth/nvlink_bandwidth.cu b/nvlink_bandwidth/nvlink_bandwidth.cu index 0c0cc88..78b2b9d 100644 --- a/nvlink_bandwidth/nvlink_bandwidth.cu +++ b/nvlink_bandwidth/nvlink_bandwidth.cu @@ -472,7 +472,7 @@ main( CUPTI_API_CALL(cuptiMetricGetNumEvents(metricId[i], &numEvents[i])); } - DRIVER_API_CALL(cuCtxCreate(&context, 0, 0)); + DRIVER_API_CALL(COMPAT_cuCtxCreate(&context, 0, 0)); CUPTI_API_CALL(cuptiMetricCreateEventGroupSets(context, (sizeof metricId) ,metricId, &pPasses)); diff --git a/pc_sampling/Makefile b/pc_sampling/Makefile index 75f6517..4cf6e77 100644 --- a/pc_sampling/Makefile +++ b/pc_sampling/Makefile @@ -57,16 +57,14 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 52 53 60 61 62 70 72 75 80 86 87 89 90 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk pc_sampling: pc_sampling.$(OBJ) $(NVCC) $(NVCC_COMPILER) $(NVCCFLAGS) -o $@ pc_sampling.$(OBJ) $(LIBS) pc_sampling.$(OBJ): pc_sampling.cu - $(NVCC) $(NVCC_COMPILER) $(NVCCFLAGS) -arch=sm_52 -lineinfo $(GENCODE_FLAGS) -c $(INCLUDES) $< + $(NVCC) $(NVCC_COMPILER) $(NVCCFLAGS) -arch=$(DEFAULT_SM_ARCH) -lineinfo $(GENCODE_FLAGS) -c $(INCLUDES) $< run: pc_sampling ./$< diff --git a/pc_sampling_start_stop/Makefile b/pc_sampling_start_stop/Makefile index 37c9b63..a61aaa9 100644 --- a/pc_sampling_start_stop/Makefile +++ b/pc_sampling_start_stop/Makefile @@ -57,10 +57,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 70 72 75 80 86 87 89 90 100 101 120 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk pc_sampling_start_stop: pc_sampling_start_stop.$(OBJ) $(NVCC) $(NVCC_COMPILER) $(NVCCFLAGS) -o $@ pc_sampling_start_stop.$(OBJ) $(LIBS) diff --git a/pc_sampling_start_stop/pc_sampling_start_stop.cu b/pc_sampling_start_stop/pc_sampling_start_stop.cu index a98ae46..027bfc6 100644 --- a/pc_sampling_start_stop/pc_sampling_start_stop.cu +++ b/pc_sampling_start_stop/pc_sampling_start_stop.cu @@ -533,7 +533,7 @@ main( } cuCtxGetCurrent(&cuCtx); - DRIVER_API_CALL(cuCtxCreate(&cuCtx, 0, deviceNum)); + DRIVER_API_CALL(COMPAT_cuCtxCreate(&cuCtx, 0, deviceNum)); DoPCSampling(cuCtx); diff --git a/pm_sampling/Makefile b/pm_sampling/Makefile index 65b7362..9276d12 100644 --- a/pm_sampling/Makefile +++ b/pm_sampling/Makefile @@ -58,10 +58,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 75 80 86 87 89 90 100 101 120 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk pm_sampling: pm_sampling.$(OBJ) $(NVCC) $(NVCC_COMPILER) $(NVCCFLAGS) -o $@ pm_sampling.$(OBJ) $(LIBS) diff --git a/profiling_injection/Makefile b/profiling_injection/Makefile index c42c739..e6ce2bf 100644 --- a/profiling_injection/Makefile +++ b/profiling_injection/Makefile @@ -80,11 +80,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 70 72 75 80 86 87 89 90 100 101 120 -# Generate SASS code for each SM architecture listed in $(SMS) -# Needed to support enhanced compatibility -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk .DEFAULT: all .PHONY: all diff --git a/profiling_injection/complex_target.cu b/profiling_injection/complex_target.cu index 02c5621..6ea63e0 100644 --- a/profiling_injection/complex_target.cu +++ b/profiling_injection/complex_target.cu @@ -182,7 +182,7 @@ main( deviceData[device].deviceId = deviceIds[device];// GPU device ID // Either set to a context, or may be NULL if a default context has been created. - DRIVER_API_CALL(cuCtxCreate(&(deviceData[device].context), 0, deviceIds[device])); + DRIVER_API_CALL(COMPAT_cuCtxCreate(&(deviceData[device].context), 0, deviceIds[device])); // Per-stream initialization & memory allocation - copy from constant host array to each device array. deviceData[device].streams.resize(numStreams); diff --git a/range_profiling/Makefile b/range_profiling/Makefile index da855ed..eab58ed 100644 --- a/range_profiling/Makefile +++ b/range_profiling/Makefile @@ -58,10 +58,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 70 72 75 80 86 87 89 90 100 101 120 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk range_profiling: range_profiling.$(OBJ) $(NVCC) $(NVCC_COMPILER) $(NVCCFLAGS) -o $@ range_profiling.$(OBJ) $(LIBS) diff --git a/range_profiling/range_profiling.cu b/range_profiling/range_profiling.cu index 5491e9c..5e58107 100644 --- a/range_profiling/range_profiling.cu +++ b/range_profiling/range_profiling.cu @@ -159,7 +159,7 @@ int main(int argc, char *argv[]) // Create a context CUcontext cuContext; - DRIVER_API_CALL(cuCtxCreate(&cuContext, 0, cuDevice)); + DRIVER_API_CALL(COMPAT_cuCtxCreate(&cuContext, 0, cuDevice)); RangeProfilerTargetPtr pRangeProfilerTarget = std::make_shared(cuContext, config); // Get chip name diff --git a/sass_metrics/Makefile b/sass_metrics/Makefile index eedfc4e..22e9ce9 100644 --- a/sass_metrics/Makefile +++ b/sass_metrics/Makefile @@ -82,10 +82,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 70 72 75 80 86 87 89 90 100 101 120 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk sass_metrics: sass_metrics.$(OBJ) $(NVCC) $(NVCC_COMPILER) $(NVCCFLAGS) -o $@ sass_metrics.$(OBJ) $(LIBS) diff --git a/sass_metrics/sass_metrics.cu b/sass_metrics/sass_metrics.cu index 4761863..fb7c351 100644 --- a/sass_metrics/sass_metrics.cu +++ b/sass_metrics/sass_metrics.cu @@ -429,7 +429,7 @@ main(int argc, char *argv[]) CUcontext cuCtx; - DRIVER_API_CALL(cuCtxCreate(&cuCtx, 0, deviceNum)); + DRIVER_API_CALL(COMPAT_cuCtxCreate(&cuCtx, 0, deviceNum)); if (bListMetrics){ ListSupportedMetrics(deviceNum); diff --git a/sass_source_map/Makefile b/sass_source_map/Makefile index 2ec3779..14cb3e3 100644 --- a/sass_source_map/Makefile +++ b/sass_source_map/Makefile @@ -75,10 +75,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 50 52 53 60 61 62 70 72 75 80 86 87 89 90 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk sass_source_map: sass_source_map.$(OBJ) $(NVCC) $(NVCC_COMPILER) $(NVCCFLAGS) -o $@ sass_source_map.$(OBJ) $(LIBS) $(INCLUDES) diff --git a/unified_memory/Makefile b/unified_memory/Makefile index d3e2f67..8d38e38 100644 --- a/unified_memory/Makefile +++ b/unified_memory/Makefile @@ -75,10 +75,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 50 52 53 60 61 62 70 72 75 80 86 87 89 90 100 101 120 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk ifeq ($(OS), Darwin) diff --git a/userrange_profiling/Makefile b/userrange_profiling/Makefile index caedb2b..2af13e8 100644 --- a/userrange_profiling/Makefile +++ b/userrange_profiling/Makefile @@ -82,10 +82,8 @@ ifneq ($(TARGET_ARCH), $(HOST_ARCH)) endif endif -# Gencode arguments -SMS ?= 70 72 75 80 86 87 89 90 100 101 120 -# Generate SASS code for each SM architecture listed in $(SMS) -$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm))) +# Include common configuration for SMS architecture detection +include ../common/common.mk .DEFAULT: all .PHONY: all diff --git a/userrange_profiling/simplecuda.cu b/userrange_profiling/simplecuda.cu index 01b3e93..8d9fa7a 100644 --- a/userrange_profiling/simplecuda.cu +++ b/userrange_profiling/simplecuda.cu @@ -217,7 +217,7 @@ bool runTest(int deviceNum, DRIVER_API_CALL(cuDeviceGet(&cuDevice, deviceNum)); CUcontext cuContext; - DRIVER_API_CALL(cuCtxCreate(&cuContext, 0, cuDevice)); + DRIVER_API_CALL(COMPAT_cuCtxCreate(&cuContext, 0, cuDevice)); CUpti_Profiler_BeginSession_Params beginSessionParams = {CUpti_Profiler_BeginSession_Params_STRUCT_SIZE}; CUpti_Profiler_SetConfig_Params setConfigParams = {CUpti_Profiler_SetConfig_Params_STRUCT_SIZE}; diff --git a/userrange_profiling/user_range_profiling.cu b/userrange_profiling/user_range_profiling.cu index 54359fd..b65f41a 100644 --- a/userrange_profiling/user_range_profiling.cu +++ b/userrange_profiling/user_range_profiling.cu @@ -417,7 +417,7 @@ main( } CUcontext cuContext; - DRIVER_API_CALL(cuCtxCreate(&cuContext, 0, cuDevice)); + DRIVER_API_CALL(COMPAT_cuCtxCreate(&cuContext, 0, cuDevice)); // Get chip name for the cuda device. CUpti_Device_GetChipName_Params getChipNameParams = { CUpti_Device_GetChipName_Params_STRUCT_SIZE };