Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion activity_trace/Makefile
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
6 changes: 2 additions & 4 deletions activity_trace_async/Makefile
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
6 changes: 2 additions & 4 deletions autorange_profiling/Makefile
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
2 changes: 1 addition & 1 deletion autorange_profiling/auto_range_profiling.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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 };
Expand Down
2 changes: 1 addition & 1 deletion autorange_profiling/simplecuda.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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};
Expand Down
6 changes: 2 additions & 4 deletions callback_event/Makefile
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
2 changes: 1 addition & 1 deletion callback_event/callback_event.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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));
Expand Down
6 changes: 2 additions & 4 deletions callback_metric/Makefile
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
2 changes: 1 addition & 1 deletion callback_metric/callback_metric.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
6 changes: 2 additions & 4 deletions callback_profiling/Makefile
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
6 changes: 2 additions & 4 deletions callback_timestamp/Makefile
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
2 changes: 1 addition & 1 deletion callback_timestamp/callback_timestamp.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
6 changes: 2 additions & 4 deletions checkpoint_kernels/Makefile
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
2 changes: 1 addition & 1 deletion checkpoint_kernels/checkpoint_kernels.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
23 changes: 23 additions & 0 deletions common/common.mk
Original file line number Diff line number Diff line change
@@ -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)))
10 changes: 10 additions & 0 deletions common/helper_cupti.h
Original file line number Diff line number Diff line change
Expand Up @@ -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 \
Expand Down
6 changes: 2 additions & 4 deletions concurrent_profiling/Makefile
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
2 changes: 1 addition & 1 deletion concurrent_profiling/concurrent_profiling.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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;

Expand Down
6 changes: 2 additions & 4 deletions cuda_graphs_trace/Makefile
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
6 changes: 2 additions & 4 deletions cuda_memory_trace/Makefile
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
6 changes: 2 additions & 4 deletions cupti_correlation/Makefile
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
6 changes: 2 additions & 4 deletions cupti_external_correlation/Makefile
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
2 changes: 1 addition & 1 deletion cupti_external_correlation/cupti_external_correlation.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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<uint64_t>(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));
Expand Down
6 changes: 2 additions & 4 deletions cupti_metric_properties/Makefile
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
2 changes: 1 addition & 1 deletion cupti_metric_properties/cupti_metric_properties.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -454,7 +454,7 @@ CUptiResult GetCounterAvailabilityImage(std::vector<uint8_t>& 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;
Expand Down
6 changes: 2 additions & 4 deletions cupti_nvtx/Makefile
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
2 changes: 1 addition & 1 deletion cupti_nvtx/cupti_nvtx.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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".
Expand Down
6 changes: 2 additions & 4 deletions cupti_nvtx_ext_payload/Makefile
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
2 changes: 1 addition & 1 deletion event_multi_gpu/event_multi_gpu.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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])));
}
Expand Down
6 changes: 2 additions & 4 deletions event_sampling/Makefile
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
2 changes: 1 addition & 1 deletion event_sampling/event_sampling.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
6 changes: 2 additions & 4 deletions nested_range_profiling/Makefile
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
2 changes: 1 addition & 1 deletion nested_range_profiling/nested_range_profiling.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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 };
Expand Down
Loading