argonne-lcf · Kerilk · Jul 23, 2026
diff --git a/backends/cuda/Makefile.am b/backends/cuda/Makefile.am
@@ -231,7 +231,7 @@ tracer_cudart.c: $(srcdir)/gen_cudart.rb $(srcdir)/tracer_cudart_helpers.include
 	SRC_DIR=$(srcdir) $(RUBY) $< > $@
 
 tracer_cudart.map: $(srcdir)/gen_cudart_linker_script.rb $(CUDART_MODEL)
-	SRC_DIR=$(srcdir) $(RUBY) $< libcudart.so.12 > $@
+	SRC_DIR=$(srcdir) $(RUBY) $< libcudart.so.13 > $@
 
 EXTRA_DIST += \
 	gen_cuda.rb \
@@ -297,7 +297,7 @@ libcudart_la_DEPENDENCIES = tracer_cudart.map libcudarttracepoints.la
 libcudart_la_CPPFLAGS = -I$(top_srcdir)/utils -I$(top_srcdir)/utils/include -I$(srcdir)/include -I./
 libcudart_la_CFLAGS = -Wall -Wextra $(WERROR) $(LIBFFI_CFLAGS) $(LTTNG_UST_CFLAGS)
 libcudart_la_LDFLAGS = $(LTTNG_UST_LIBS) -ldl -lpthread $(LIBFFI_LIBS)
-libcudart_la_LDFLAGS += -version-number 12:1:55 -Wl,--version-script,tracer_cudart.map
+libcudart_la_LDFLAGS += -version-number 13:0:48 -Wl,--version-script,tracer_cudart.map
 libcudart_la_LIBADD = libcudarttracepoints.la
 
 tmplibdir = $(libdir)/tmp

diff --git a/backends/cuda/cuda_meta_parameters.yaml b/backends/cuda/cuda_meta_parameters.yaml
@@ -838,3 +838,243 @@ meta_parameters:
   cuGetExportTable:
     - [InScalar, pExportTableId]
     - [OutScalar, ppExportTable]
+  cuCheckpointProcessCheckpoint:
+    - [InScalar, args]
+  cuCheckpointProcessGetRestoreThreadId:
+    - [OutScalar, tid]
+  cuCheckpointProcessGetState:
+    - [OutScalar, state]
+  cuCheckpointProcessLock:
+    - [InScalar, args]
+  cuCheckpointProcessRestore:
+    - [InScalar, args]
+    - [InArray, args->gpuPairs, args->gpuPairsCount]
+  cuCheckpointProcessUnlock:
+    - [InScalar, args]
+  cuCtxCreate_v4:
+    - [OutScalar, pctx]
+    - [InScalar, ctxCreateParams]
+    - [InArray, ctxCreateParams->execAffinityParams, ctxCreateParams->numExecAffinityParams]
+    - [InScalar, ctxCreateParams->cigParams]
+  cuCtxFromGreenCtx:
+    - [OutScalar, pContext]
+  cuCtxGetDevice_v2:
+    - [OutScalar, device]
+  cuCtxGetDevResource:
+    - [OutScalar, resource]
+  cuDeviceGetDevResource:
+    - [OutScalar, resource]
+  cuDeviceGetHostAtomicCapabilities:
+    - [OutArray, capabilities, count]
+    - [InArray, operations, count]
+  cuDeviceGetP2PAtomicCapabilities:
+    - [OutArray, capabilities, count]
+    - [InArray, operations, count]
+  cuDeviceRegisterAsyncNotification:
+    - [OutScalar, callback]
+  cuDevResourceGenerateDesc:
+    - [OutScalar, phDesc]
+    - [InArray, resources, nbResources]
+  cuDevSmResourceSplitByCount:
+    - [InOutScalar, nbGroups]
+    - [OutArray, result, nbGroups]
+    - [InScalar, input]
+    - [OutScalar, remaining]
+  cuEventElapsedTime_v2:
+    - [OutScalar, pMilliseconds]
+  cuFuncGetName:
+    - [OutPtrString, name]
+  cuFuncGetParamInfo:
+    - [OutScalar, paramOffset]
+    - [OutScalar, paramSize]
+  cuFuncIsLoaded:
+    - [OutScalar, state]
+  cuGraphAddDependencies_v2:
+    - [InArray, from, numDependencies]
+    - [InArray, to, numDependencies]
+    - [InArray, edgeData, numDependencies]
+  cuGraphAddNode:
+    - [OutScalar, phGraphNode]
+    - [InArray, dependencies, numDependencies]
+    - [InScalar, nodeParams]
+  cuGraphAddNode_v2:
+    - [OutScalar, phGraphNode]
+    - [InArray, dependencies, numDependencies]
+    - [InArray, dependencyData, numDependencies]
+    - [InScalar, nodeParams]
+  cuGraphConditionalHandleCreate:
+    - [OutScalar, pHandle_out]
+  cuGraphExecNodeSetParams:
+    - [InScalar, nodeParams]
+  cuGraphGetEdges_v2:
+    - [InOutScalar, numEdges]
+    - [OutArray, from, numEdges]
+    - [OutArray, to, numEdges]
+    - [OutArray, edgeData, numEdges]
+  cuGraphNodeGetDependencies_v2:
+    - [InOutScalar, numDependencies]
+    - [OutArray, dependencies, numDependencies]
+    - [OutArray, edgeData, numDependencies]
+  cuGraphNodeGetDependentNodes_v2:
+    - [InOutScalar, numDependentNodes]
+    - [OutArray, dependentNodes, numDependentNodes]
+    - [OutArray, edgeData, numDependentNodes]
+  cuGraphNodeSetParams:
+    - [InScalar, nodeParams]
+  cuGraphRemoveDependencies_v2:
+    - [InArray, from, numDependencies]
+    - [InArray, to, numDependencies]
+    - [InArray, edgeData, numDependencies]
+  cuGreenCtxCreate:
+    - [OutScalar, phCtx]
+  cuGreenCtxGetDevResource:
+    - [OutScalar, resource]
+  cuGreenCtxGetId:
+    - [OutScalar, greenCtxId]
+  cuGreenCtxStreamCreate:
+    - [OutScalar, phStream]
+  cuKernelGetLibrary:
+    - [OutScalar, pLib]
+  cuKernelGetName:
+    - [OutPtrString, name]
+  cuKernelGetParamInfo:
+    - [OutScalar, paramOffset]
+    - [OutScalar, paramSize]
+  cuLibraryEnumerateKernels:
+    - [OutArray, kernels, numKernels]
+  cuLibraryGetKernelCount:
+    - [OutScalar, count]
+  cuLogsCurrent:
+    - [OutScalar, iterator_out]
+  cuLogsDumpToFile:
+    - [InScalar, iterator]
+    - [InString, pathToFile]
+  cuLogsDumpToMemory:
+    - [InScalar, iterator]
+    - [InOutScalar, size]
+    - [OutArray, buffer, size]
+  cuLogsRegisterCallback:
+    - [OutScalar, callback_out]
+  cuMemBatchDecompressAsync:
+    - [InArray, paramsArray, count]
+    - [OutScalar, errorIndex]
+  cuMemBatchDecompressAsync_ptsz:
+    - [InArray, paramsArray, count]
+    - [OutScalar, errorIndex]
+  cuMemcpy3DBatchAsync:
+    - [InArray, opList, numOps]
+    - [OutScalar, failIdx]
+  cuMemcpy3DBatchAsync_ptsz:
+    - [InArray, opList, numOps]
+    - [OutScalar, failIdx]
+  cuMemcpy3DBatchAsync_v2:
+    - [InArray, opList, numOps]
+  cuMemcpy3DBatchAsync_v2_ptsz:
+    - [InArray, opList, numOps]
+  cuMemcpyBatchAsync:
+    - [InArray, dsts, count]
+    - [InArray, srcs, count]
+    - [InArray, sizes, count]
+    - [InArray, attrs, numAttrs]
+    - [InArray, attrsIdxs, numAttrs]
+    - [OutScalar, failIdx]
+  cuMemcpyBatchAsync_ptsz:
+    - [InArray, dsts, count]
+    - [InArray, srcs, count]
+    - [InArray, sizes, count]
+    - [InArray, attrs, numAttrs]
+    - [InArray, attrsIdxs, numAttrs]
+    - [OutScalar, failIdx]
+  cuMemcpyBatchAsync_v2:
+    - [InArray, dsts, count]
+    - [InArray, srcs, count]
+    - [InArray, sizes, count]
+    - [InArray, attrs, numAttrs]
+    - [InArray, attrsIdxs, numAttrs]
+  cuMemcpyBatchAsync_v2_ptsz:
+    - [InArray, dsts, count]
+    - [InArray, srcs, count]
+    - [InArray, sizes, count]
+    - [InArray, attrs, numAttrs]
+    - [InArray, attrsIdxs, numAttrs]
+  cuMemDiscardAndPrefetchBatchAsync:
+    - [InArray, dptrs, count]
+    - [InArray, sizes, count]
+    - [InArray, prefetchLocs, numPrefetchLocs]
+    - [InArray, prefetchLocIdxs, numPrefetchLocs]
+  cuMemDiscardAndPrefetchBatchAsync_ptsz:
+    - [InArray, dptrs, count]
+    - [InArray, sizes, count]
+    - [InArray, prefetchLocs, numPrefetchLocs]
+    - [InArray, prefetchLocIdxs, numPrefetchLocs]
+  cuMemDiscardBatchAsync:
+    - [InArray, dptrs, count]
+    - [InArray, sizes, count]
+  cuMemDiscardBatchAsync_ptsz:
+    - [InArray, dptrs, count]
+    - [InArray, sizes, count]
+  cuMemGetDefaultMemPool:
+    - [OutScalar, pool_out]
+    - [InScalar, location]
+  cuMemGetMemPool:
+    - [OutScalar, pool]
+    - [InScalar, location]
+  cuMemPrefetchBatchAsync:
+    - [InArray, dptrs, count]
+    - [InArray, sizes, count]
+    - [InArray, prefetchLocs, numPrefetchLocs]
+    - [InArray, prefetchLocIdxs, numPrefetchLocs]
+  cuMemPrefetchBatchAsync_ptsz:
+    - [InArray, dptrs, count]
+    - [InArray, sizes, count]
+    - [InArray, prefetchLocs, numPrefetchLocs]
+    - [InArray, prefetchLocIdxs, numPrefetchLocs]
+  cuMemSetMemPool:
+    - [InScalar, location]
+  cuModuleEnumerateFunctions:
+    - [OutArray, functions, numFunctions]
+  cuModuleGetFunctionCount:
+    - [OutScalar, count]
+  cuStreamBeginCaptureToGraph:
+    - [InArray, dependencies, numDependencies]
+    - [InArray, dependencyData, numDependencies]
+  cuStreamBeginCaptureToGraph_ptsz: 
+    - [InArray, dependencies, numDependencies]
+    - [InArray, dependencyData, numDependencies]
+  cuStreamGetCaptureInfo_v3:
+    - [OutScalar, captureStatus_out]
+    - [OutScalar, id_out]
+    - [OutScalar, graph_out]
+    - [OutScalar, numDependencies_out]
+    - [OutArrayByRef, dependencies_out, numDependencies_out]
+    - [OutArrayByRef, edgeData_out, numDependencies_out]
+  cuStreamGetCaptureInfo_v3_ptsz:
+    - [OutScalar, captureStatus_out]
+    - [OutScalar, id_out]
+    - [OutScalar, graph_out]
+    - [OutScalar, numDependencies_out]
+    - [OutArrayByRef, dependencies_out, numDependencies_out]
+    - [OutArrayByRef, edgeData_out, numDependencies_out]
+  cuStreamGetCtx_v2:
+    - [OutScalar, pCtx]
+    - [OutScalar, pGreenCtx]
+  cuStreamGetCtx_v2_ptsz:
+    - [OutScalar, pCtx]
+    - [OutScalar, pGreenCtx]
+  cuStreamGetDevice:
+    - [OutScalar, device]
+  cuStreamGetDevice_ptsz:
+    - [OutScalar, device]
+  cuStreamGetGreenCtx:
+    - [OutScalar, phCtx]
+  cuStreamUpdateCaptureDependencies_v2:
+    - [InArray, dependencies, numDependencies]
+    - [InArray, dependencyData, numDependencies]
+  cuStreamUpdateCaptureDependencies_v2_ptsz:
+    - [InArray, dependencies, numDependencies]
+    - [InArray, dependencyData, numDependencies]
+  cuTensorMapEncodeIm2colWide:
+    - [OutScalar, tensorMap]
+    - [InArray, globalDim, tensorRank]
+    - [InArray, globalStrides, tensorRank]
+    - [InArray, elementStrides, tensorRank]
diff --git a/backends/cuda/cuda_model.rb b/backends/cuda/cuda_model.rb
@@ -167,7 +167,8 @@ def upper_snake_case(str)
 }
 
 profiling_start_no_stream = profiling_start.call('NULL')
-profiling_start_stream = profiling_start.call('hStream')
+profiling_start_hstream = profiling_start.call('hStream')
+profiling_start_stream = profiling_start.call('stream')
 
 profiling_start_config = <<EOF
   CUevent _hStart = NULL;
@@ -183,22 +184,24 @@ def upper_snake_case(str)
 }
 
 profiling_stop_no_stream = profiling_stop.call('NULL')
-profiling_stop_stream = profiling_stop.call('hStream')
+profiling_stop_hstream = profiling_stop.call('hStream')
+profiling_stop_stream = profiling_stop.call('stream')
 
 profiling_stop_config = <<EOF
   if (_do_profile && config)
     _event_profile(_retval, _hStart, config->hStream);
 EOF
 
 stream_commands = []
+hstream_commands = []
 no_stream_commands = []
 mem_commands = $cuda_commands.select { |c| c.name.match(/cuMemcpy|cuMemset/) }
 mem_stream_commands = mem_commands.select { |c| c.name.match(/Async/) }
 mem_no_stream_commands = mem_commands - mem_stream_commands
-stream_commands += mem_stream_commands.collect(&:name)
+hstream_commands += mem_stream_commands.collect(&:name)
 no_stream_commands += mem_no_stream_commands.collect(&:name)
 
-stream_commands += %w[
+hstream_commands += %w[
   cuMemPrefetchAsync
   cuMemPrefetchAsync_ptsz
   cuLaunchGridAsync
@@ -208,6 +211,37 @@ def upper_snake_case(str)
   cuLaunchKernel_ptsz
   cuLaunchCooperativeKernel_ptsz
   cuLaunchHostFunc_ptsz
+  cuMemAllocAsync
+  cuMemAllocAsync_ptsz
+  cuMemAllocFromPoolAsync
+  cuMemAllocFromPoolAsync_ptsz
+  cuStreamAttachMemAsync
+  cuStreamAttachMemAsync_ptsz
+  cuMemFreeAsync
+  cuMemFreeAsync_ptsz
+  cuMemPrefetchBatchAsync
+  cuMemPrefetchBatchAsync_ptsz
+  cuMemDiscardAndPrefetchBatchAsync
+  cuMemDiscardAndPrefetchBatchAsync_ptsz
+  cuMemDiscardBatchAsync
+  cuMemDiscardBatchAsync_ptsz
+  cuMemPrefetchAsync_v2
+  cuMemPrefetchAsync_v2_ptsz
+]
+
+stream_commands += $cuda_commands.select { |c| c.name.match(/cuStreamWaitValue|cuStreamWriteValue/) }.collect(&:name)
+
+stream_commands += %w[
+  cuMemBatchDecompressAsync
+  cuMemBatchDecompressAsync_ptsz
+  cuSignalExternalSemaphoresAsync
+  cuSignalExternalSemaphoresAsync_ptsz
+  cuWaitExternalSemaphoresAsync
+  cuWaitExternalSemaphoresAsync_ptsz
+  cuStreamBatchMemOp
+  cuStreamBatchMemOp_ptsz
+  cuStreamBatchMemOp_v2
+  cuStreamBatchMemOp_v2_ptsz
 ]
 
 no_stream_commands += %w[
@@ -220,6 +254,11 @@ def upper_snake_case(str)
   cuLaunchKernelEx_ptsz
 ]
 
+hstream_commands.each do |m|
+  register_prologue m, profiling_start_hstream
+  register_epilogue m, profiling_stop_hstream
+end
+
 stream_commands.each do |m|
   register_prologue m, profiling_start_stream
   register_epilogue m, profiling_stop_stream

diff --git a/backends/cuda/gen_cuda_library.rb b/backends/cuda/gen_cuda_library.rb
@@ -1,7 +1,11 @@
 require_relative 'gen_cuda_library_base'
 
 def print_enum(name, enum)
-  print_enum_with_namespace(:CUDA, name, enum)
+  print_enum_with_namespace(
+    :CUDA, name, enum,
+    filter_members: ->(m) { m.name != 'CU_COREDUMP_LIGHTWEIGHT_FLAGS' },
+    fix_values: ->(v) { v.gsub('1u <<', '1 <<') }
+  )
 end
 
 def print_cuda_object(object)
@@ -13,10 +17,17 @@ def print_cuda_object(object)
 puts <<~EOF
   module CUDA
     CU_COMPUTE_ACCELERATED_TARGET_BASE = 0x10000
+    CU_COMPUTE_FAMILY_TARGET_BASE = 0x20000
     CU_TARGET_COMPUTE_90 = 90
+    CU_TARGET_COMPUTE_100 = 100
+    CU_TARGET_COMPUTE_103 = 103
+    CU_TARGET_COMPUTE_110 = 110
+    CU_TARGET_COMPUTE_120 = 120
+    CU_TARGET_COMPUTE_121 = 121
     CU_TENSOR_MAP_NUM_QWORDS = 16
     CU_IPC_HANDLE_SIZE = 64
     CUDA_IPC_HANDLE_SIZE = 64
+    RESOURCE_ABI_EXTERNAL_BYTES = 48
     extend FFI::Library
 
     module Handle
@@ -76,11 +87,13 @@ def print_struct(name, struct)
   typedef :uint64, #{to_ffi_name('CUsurfObject')}
   typedef :uint64, #{to_ffi_name('CUmemGenericAllocationHandle_v1')}
   typedef #{to_ffi_name('CUmemGenericAllocationHandle_v1')}, #{to_ffi_name('CUmemGenericAllocationHandle')}
+  typedef :uint64, #{to_ffi_name('CUgraphConditionalHandle')}
 EOF
 
 $all_types.each do |t|
   if t.type.is_a? YAMLCAst::Enum
     enum = $all_enums.find { |e| t.type.name == e.name }
+    enum ||= t.type
     print_enum(t.name, enum)
   elsif $objects.include?(t.name)
     print_cuda_object(t.name)