CUDA Runtime API. API Reference Manual (vRelease Version, January 2024)

Modules

pNodeParams

- Updated Parameters to set

Returns

cudaSuccess, cudaErrorInvalidValue,

Description

Updates the work represented by node in hGraphExec as though node had contained

pNodeParams at instantiation. node must remain in the graph which was used to instantiate

hGraphExec. Changed edges to and from node are ignored.

The source and destination memory in pNodeParams must be allocated from the same contexts as

the original source and destination memory. Both the instantiation-time memory operands and the

memory operands in pNodeParams must be 1-dimensional. Zero-length operations are not supported.

The modifications only affect future launches of hGraphExec. Already enqueued or running

launches of hGraphExec are not affected by this call. node is also not modified by this call.

Returns cudaErrorInvalidValue if the memory operands' mappings changed or either the original or

new memory operands are multidimensional.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphExecNodeSetParams, cudaGraphAddMemcpyNode,

cudaGraphMemcpyNodeSetParams, cudaGraphExecMemcpyNodeSetParamsToSymbol,

cudaGraphExecMemcpyNodeSetParamsFromSymbol, cudaGraphExecMemcpyNodeSetParams1D,

cudaGraphExecKernelNodeSetParams, cudaGraphExecMemsetNodeSetParams,

cudaGraphExecHostNodeSetParams, cudaGraphExecChildGraphNodeSetParams,

cudaGraphExecEventRecordNodeSetEvent, cudaGraphExecEventWaitNodeSetEvent,

cudaGraphExecExternalSemaphoresSignalNodeSetParams,

cudaGraphExecExternalSemaphoresWaitNodeSetParams, cudaGraphExecUpdate,

cudaGraphInstantiate

343

Modules

__host__cudaError_t

cudaGraphExecMemcpyNodeSetParams1D

(cudaGraphExec_t hGraphExec, cudaGraphNode_t node,

void *dst, const void *src, size_t count, cudaMemcpyKind

kind)

Sets the parameters for a memcpy node in the given graphExec to perform a 1-dimensional copy.

Parameters

hGraphExec

- The executable graph in which to set the specified node

node

- Memcpy node from the graph which was used to instantiate graphExec

dst

- Destination memory address

src

- Source memory address

count

- Size in bytes to copy

kind

- Type of transfer

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Updates the work represented by node in hGraphExec as though node had contained the given

params at instantiation. node must remain in the graph which was used to instantiate hGraphExec.

Changed edges to and from node are ignored.

src and dst must be allocated from the same contexts as the original source and destination memory.

The instantiation-time memory operands must be 1-dimensional. Zero-length operations are not

supported.

The modifications only affect future launches of hGraphExec. Already enqueued or running

launches of hGraphExec are not affected by this call. node is also not modified by this call.

Returns cudaErrorInvalidValue if the memory operands' mappings changed or the original memory

operands are multidimensional.

344

Modules

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphAddMemcpyNode, cudaGraphAddMemcpyNode1D,

cudaGraphMemcpyNodeSetParams, cudaGraphMemcpyNodeSetParams1D,

cudaGraphExecMemcpyNodeSetParams, cudaGraphExecKernelNodeSetParams,

cudaGraphExecMemsetNodeSetParams, cudaGraphExecHostNodeSetParams,

cudaGraphExecChildGraphNodeSetParams, cudaGraphExecEventRecordNodeSetEvent,

cudaGraphExecEventWaitNodeSetEvent, cudaGraphExecExternalSemaphoresSignalNodeSetParams,

cudaGraphExecExternalSemaphoresWaitNodeSetParams, cudaGraphExecUpdate,

cudaGraphInstantiate

__host__cudaError_t

cudaGraphExecMemcpyNodeSetParamsFromSymbol

(cudaGraphExec_t hGraphExec, cudaGraphNode_t node,

void *dst, const void *symbol, size_t count, size_t offset,

cudaMemcpyKind kind)

Sets the parameters for a memcpy node in the given graphExec to copy from a symbol on the device.

Parameters

hGraphExec

- The executable graph in which to set the specified node

node

- Memcpy node from the graph which was used to instantiate graphExec

dst

- Destination memory address

symbol

- Device symbol address

count

- Size in bytes to copy

345

Modules

offset

- Offset from start of symbol in bytes

kind

- Type of transfer

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Updates the work represented by node in hGraphExec as though node had contained the given

params at instantiation. node must remain in the graph which was used to instantiate hGraphExec.

Changed edges to and from node are ignored.

symbol and dst must be allocated from the same contexts as the original source and destination

memory. The instantiation-time memory operands must be 1-dimensional. Zero-length operations are

not supported.

The modifications only affect future launches of hGraphExec. Already enqueued or running

launches of hGraphExec are not affected by this call. node is also not modified by this call.

Returns cudaErrorInvalidValue if the memory operands' mappings changed or the original memory

operands are multidimensional.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphAddMemcpyNode, cudaGraphAddMemcpyNodeFromSymbol,

cudaGraphMemcpyNodeSetParams, cudaGraphMemcpyNodeSetParamsFromSymbol,

cudaGraphExecMemcpyNodeSetParams, cudaGraphExecMemcpyNodeSetParamsToSymbol,

cudaGraphExecKernelNodeSetParams, cudaGraphExecMemsetNodeSetParams,

cudaGraphExecHostNodeSetParams, cudaGraphExecChildGraphNodeSetParams,

cudaGraphExecEventRecordNodeSetEvent, cudaGraphExecEventWaitNodeSetEvent,

cudaGraphExecExternalSemaphoresSignalNodeSetParams,

cudaGraphExecExternalSemaphoresWaitNodeSetParams, cudaGraphExecUpdate,

cudaGraphInstantiate

346

Modules

__host__cudaError_t

cudaGraphExecMemcpyNodeSetParamsToSymbol

(cudaGraphExec_t hGraphExec, cudaGraphNode_t node,

const void *symbol, const void *src, size_t count, size_t

offset, cudaMemcpyKind kind)

Sets the parameters for a memcpy node in the given graphExec to copy to a symbol on the device.

Parameters

hGraphExec

- The executable graph in which to set the specified node

node

- Memcpy node from the graph which was used to instantiate graphExec

symbol

- Device symbol address

src

- Source memory address

count

- Size in bytes to copy

offset

- Offset from start of symbol in bytes

kind

- Type of transfer

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Updates the work represented by node in hGraphExec as though node had contained the given

params at instantiation. node must remain in the graph which was used to instantiate hGraphExec.

Changed edges to and from node are ignored.

src and symbol must be allocated from the same contexts as the original source and destination

memory. The instantiation-time memory operands must be 1-dimensional. Zero-length operations are

not supported.

The modifications only affect future launches of hGraphExec. Already enqueued or running

launches of hGraphExec are not affected by this call. node is also not modified by this call.

Returns cudaErrorInvalidValue if the memory operands' mappings changed or the original memory

operands are multidimensional.

347

Modules

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphAddMemcpyNode, cudaGraphAddMemcpyNodeToSymbol,

cudaGraphMemcpyNodeSetParams, cudaGraphMemcpyNodeSetParamsToSymbol,

cudaGraphExecMemcpyNodeSetParams, cudaGraphExecMemcpyNodeSetParamsFromSymbol,

cudaGraphExecKernelNodeSetParams, cudaGraphExecMemsetNodeSetParams,

cudaGraphExecHostNodeSetParams, cudaGraphExecChildGraphNodeSetParams,

cudaGraphExecEventRecordNodeSetEvent, cudaGraphExecEventWaitNodeSetEvent,

cudaGraphExecExternalSemaphoresSignalNodeSetParams,

cudaGraphExecExternalSemaphoresWaitNodeSetParams, cudaGraphExecUpdate,

cudaGraphInstantiate

__host__cudaError_t

cudaGraphExecMemsetNodeSetParams (cudaGraphExec_t

hGraphExec, cudaGraphNode_t node, const

cudaMemsetParams *pNodeParams)

Sets the parameters for a memset node in the given graphExec.

Parameters

hGraphExec

- The executable graph in which to set the specified node

node

- Memset node from the graph which was used to instantiate graphExec

pNodeParams

- Updated Parameters to set

Returns

cudaSuccess, cudaErrorInvalidValue,

348

Modules

Description

Updates the work represented by node in hGraphExec as though node had contained

pNodeParams at instantiation. node must remain in the graph which was used to instantiate

hGraphExec. Changed edges to and from node are ignored.

Zero sized operations are not supported.

The new destination pointer in pNodeParams must be to the same kind of allocation as the original

destination pointer and have the same context association and device mapping as the original

destination pointer.

Both the value and pointer address may be updated. Changing other aspects of the memset (width,

height, element size or pitch) may cause the update to be rejected. Specifically, for 2d memsets, all

dimension changes are rejected. For 1d memsets, changes in height are explicitly rejected and other

changes are opportunistically allowed if the resulting work maps onto the work resources already

allocated for the node.

The modifications only affect future launches of hGraphExec. Already enqueued or running

launches of hGraphExec are not affected by this call. node is also not modified by this call.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphExecNodeSetParams, cudaGraphAddMemsetNode,

cudaGraphMemsetNodeSetParams, cudaGraphExecKernelNodeSetParams,

cudaGraphExecMemcpyNodeSetParams, cudaGraphExecHostNodeSetParams,

cudaGraphExecChildGraphNodeSetParams, cudaGraphExecEventRecordNodeSetEvent,

cudaGraphExecEventWaitNodeSetEvent, cudaGraphExecExternalSemaphoresSignalNodeSetParams,

cudaGraphExecExternalSemaphoresWaitNodeSetParams, cudaGraphExecUpdate,

cudaGraphInstantiate

349

Modules

__host__cudaError_t cudaGraphExecNodeSetParams

(cudaGraphExec_t graphExec, cudaGraphNode_t node,

cudaGraphNodeParams *nodeParams)

Update's a graph node's parameters in an instantiated graph.

Parameters

graphExec

- The executable graph in which to update the specified node

node

- Corresponding node from the graph from which graphExec was instantiated

nodeParams

- Updated Parameters to set

Returns

cudaSuccess, cudaErrorInvalidValue, cudaErrorInvalidDeviceFunction, cudaErrorNotSupported

Description

Sets the parameters of a node in an executable graph graphExec. The node is identified by

the corresponding node node in the non-executable graph from which the executable graph was

instantiated. node must not have been removed from the original graph.

The modifications only affect future launches of graphExec. Already enqueued or running launches

of graphExec are not affected by this call. node is also not modified by this call.

Allowed changes to parameters on executable graphs are as follows:

Node type

Allowed changes

kernel

See cudaGraphExecKernelNodeSetParams

memcpy

Addresses for 1-dimensional copies

if allocated in same context; see

cudaGraphExecMemcpyNodeSetParams

memset

Addresses for 1-dimensional memsets

if allocated in same context; see

cudaGraphExecMemsetNodeSetParams

host

Unrestricted

child graph

Topology must match and restrictions apply

recursively; see cudaGraphExecUpdate

event wait

Unrestricted

350

Modules

Node type

Allowed changes

event record

Unrestricted

external semaphore signal

Number of semaphore operations cannot change

external semaphore wait

Number of semaphore operations cannot change

memory allocation

API unsupported

memory free

API unsupported

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphAddNode, cudaGraphNodeSetParams cudaGraphExecUpdate, cudaGraphInstantiate

__host__cudaError_t cudaGraphExecUpdate

(cudaGraphExec_t hGraphExec, cudaGraph_t hGraph,

cudaGraphExecUpdateResultInfo *resultInfo)

Check whether an executable graph can be updated with a graph and perform the update if possible.

Parameters

hGraphExec

The instantiated graph to be updated

hGraph

The graph containing the updated parameters

resultInfo

the error info structure

Returns

cudaSuccess, cudaErrorGraphExecUpdateFailure,

351

Modules

Description

Updates the node parameters in the instantiated graph specified by hGraphExec with the node

parameters in a topologically identical graph specified by hGraph.

Limitations:

‣ Kernel nodes:

‣ The owning context of the function cannot change.

‣ A node whose function originally did not use CUDA dynamic parallelism cannot be updated to

a function which uses CDP.

‣ A node whose function originally did not make device-side update calls cannot be updated to a

function which makes device-side update calls.

‣ A cooperative node cannot be updated to a non-cooperative node, and vice-versa.

‣ If the graph was instantiated with cudaGraphInstantiateFlagUseNodePriority, the priority

attribute cannot change. Equality is checked on the originally requested priority values, before

they are clamped to the device's supported range.

‣ If hGraphExec was not instantiated for device launch, a node whose function originally did

not use device-side cudaGraphLaunch() cannot be updated to a function which uses device-side

cudaGraphLaunch() unless the node resides on the same device as nodes which contained such

calls at instantiate-time. If no such calls were present at instantiation, these updates cannot be

performed at all.

‣ Neither hGraph nor hGraphExec may contain device-updatable kernel nodes.

‣ Memset and memcpy nodes:

‣ The CUDA device(s) to which the operand(s) was allocated/mapped cannot change.

‣ The source/destination memory must be allocated from the same contexts as the original

source/destination memory.

‣ For 2d memsets, only address and assigned value may be updated.

‣ For 1d memsets, updating dimensions is also allowed, but may fail if the resulting operation

doesn't map onto the work resources already allocated for the node.

‣ Additional memcpy node restrictions:

‣ Changing either the source or destination memory type(i.e. CU_MEMORYTYPE_DEVICE,

CU_MEMORYTYPE_ARRAY, etc.) is not supported.

‣ Conditional nodes:

‣ Changing node parameters is not supported.

‣ Changing parameters of nodes within the conditional body graph is subject to the rules above.

‣ Conditional handle flags and default values are updated as part of the graph update.

Note: The API may add further restrictions in future releases. The return code should always be

checked.

352

Modules

cudaGraphExecUpdate sets the result member of resultInfo to

cudaGraphExecUpdateErrorTopologyChanged under the following conditions:

‣ The count of nodes directly in hGraphExec and hGraph differ, in which case resultInfo-

>errorNode is set to NULL.

‣ hGraph has more exit nodes than hGraph, in which case resultInfo->errorNode is set to one of

the exit nodes in hGraph.

‣ A node in hGraph has a different number of dependencies than the node from hGraphExec it is

paired with, in which case resultInfo->errorNode is set to the node from hGraph.

‣ A node in hGraph has a dependency that does not match with the corresponding dependency of

the paired node from hGraphExec. resultInfo->errorNode will be set to the node from hGraph.

resultInfo->errorFromNode will be set to the mismatched dependency. The dependencies are

paired based on edge order and a dependency does not match when the nodes are already paired

based on other edges examined in the graph.

cudaGraphExecUpdate sets the result member of resultInfo to:

‣ cudaGraphExecUpdateError if passed an invalid value.

‣ cudaGraphExecUpdateErrorTopologyChanged if the graph topology changed

‣ cudaGraphExecUpdateErrorNodeTypeChanged if the type of a node changed, in which case

hErrorNode_out is set to the node from hGraph.

‣ cudaGraphExecUpdateErrorFunctionChanged if the function of a kernel node changed (CUDA

driver < 11.2)

‣ cudaGraphExecUpdateErrorUnsupportedFunctionChange if the func field of a kernel changed in

an unsupported way(see note above), in which case hErrorNode_out is set to the node from

hGraph

‣ cudaGraphExecUpdateErrorParametersChanged if any parameters to a node changed in a way that

is not supported, in which case hErrorNode_out is set to the node from hGraph

‣ cudaGraphExecUpdateErrorAttributesChanged if any attributes of a node changed in a way that is

not supported, in which case hErrorNode_out is set to the node from hGraph

‣ cudaGraphExecUpdateErrorNotSupported if something about a node is unsupported, like the

node's type or configuration, in which case hErrorNode_out is set to the node from hGraph

If the update fails for a reason not listed above, the result member of resultInfo will be

set to cudaGraphExecUpdateError. If the update succeeds, the result member will be set to

cudaGraphExecUpdateSuccess.

cudaGraphExecUpdate returns cudaSuccess when the updated was performed successfully. It returns

cudaErrorGraphExecUpdateFailure if the graph update was not performed because it included changes

which violated constraints specific to instantiated graph update.

Note:

‣ Graph objects are not threadsafe. More here.

353

Modules

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphInstantiate

__host__cudaError_t

cudaGraphExternalSemaphoresSignalNodeGetParams

(cudaGraphNode_t hNode,

cudaExternalSemaphoreSignalNodeParams *params_out)

Returns an external semaphore signal node's parameters.

Parameters

hNode

- Node to get the parameters for

params_out

- Pointer to return the parameters

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Returns the parameters of an external semaphore signal node hNode in params_out. The

extSemArray and paramsArray returned in params_out, are owned by the node. This memory

remains valid until the node is destroyed or its parameters are modified, and should not be modified

directly. Use cudaGraphExternalSemaphoresSignalNodeSetParams to update the parameters of this

node.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

354

Modules

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaLaunchKernel, cudaGraphAddExternalSemaphoresSignalNode,

cudaGraphExternalSemaphoresSignalNodeSetParams, cudaGraphAddExternalSemaphoresWaitNode,

cudaSignalExternalSemaphoresAsync, cudaWaitExternalSemaphoresAsync

__host__cudaError_t

cudaGraphExternalSemaphoresSignalNodeSetParams

(cudaGraphNode_t hNode, const

cudaExternalSemaphoreSignalNodeParams *nodeParams)

Sets an external semaphore signal node's parameters.

Parameters

hNode

- Node to set the parameters for

nodeParams

- Parameters to copy

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Sets the parameters of an external semaphore signal node hNode to nodeParams.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

355

Modules

cudaGraphNodeSetParams, cudaGraphAddExternalSemaphoresSignalNode,

cudaGraphExternalSemaphoresSignalNodeSetParams, cudaGraphAddExternalSemaphoresWaitNode,

cudaSignalExternalSemaphoresAsync, cudaWaitExternalSemaphoresAsync

__host__cudaError_t

cudaGraphExternalSemaphoresWaitNodeGetParams

(cudaGraphNode_t hNode,

cudaExternalSemaphoreWaitNodeParams *params_out)

Returns an external semaphore wait node's parameters.

Parameters

hNode

- Node to get the parameters for

params_out

- Pointer to return the parameters

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Returns the parameters of an external semaphore wait node hNode in params_out. The

extSemArray and paramsArray returned in params_out, are owned by the node. This memory

remains valid until the node is destroyed or its parameters are modified, and should not be modified

directly. Use cudaGraphExternalSemaphoresSignalNodeSetParams to update the parameters of this

node.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

356

Modules

cudaLaunchKernel, cudaGraphAddExternalSemaphoresWaitNode,

cudaGraphExternalSemaphoresWaitNodeSetParams, cudaGraphAddExternalSemaphoresWaitNode,

cudaSignalExternalSemaphoresAsync, cudaWaitExternalSemaphoresAsync

__host__cudaError_t

cudaGraphExternalSemaphoresWaitNodeSetParams

(cudaGraphNode_t hNode, const

cudaExternalSemaphoreWaitNodeParams *nodeParams)

Sets an external semaphore wait node's parameters.

Parameters

hNode

- Node to set the parameters for

nodeParams

- Parameters to copy

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Sets the parameters of an external semaphore wait node hNode to nodeParams.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphNodeSetParams, cudaGraphAddExternalSemaphoresWaitNode,

cudaGraphExternalSemaphoresWaitNodeSetParams, cudaGraphAddExternalSemaphoresWaitNode,

cudaSignalExternalSemaphoresAsync, cudaWaitExternalSemaphoresAsync

357

Modules

__host__cudaError_t cudaGraphGetEdges (cudaGraph_t

graph, cudaGraphNode_t *from, cudaGraphNode_t *to,

size_t *numEdges)

Returns a graph's dependency edges.

Parameters

graph

- Graph to get the edges from

from

- Location to return edge endpoints

numEdges

- See description

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Returns a list of graph's dependency edges. Edges are returned via corresponding indices in from

and to; that is, the node in to[i] has a dependency on the node in from[i]. from and to may both

be NULL, in which case this function only returns the number of edges in numEdges. Otherwise,

numEdges entries will be filled in. If numEdges is higher than the actual number of edges, the

remaining entries in from and to will be set to NULL, and the number of edges actually returned will

be written to numEdges.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

358

Modules

cudaGraphGetNodes, cudaGraphGetRootNodes, cudaGraphAddDependencies,

cudaGraphRemoveDependencies, cudaGraphNodeGetDependencies,

cudaGraphNodeGetDependentNodes

__host__cudaError_t cudaGraphGetEdges_v2

(cudaGraph_t graph, cudaGraphNode_t *from,

cudaGraphNode_t *to, cudaGraphEdgeData *edgeData,

size_t *numEdges)

Returns a graph's dependency edges (12.3+).

Parameters

graph

- Graph to get the edges from

from

- Location to return edge endpoints

edgeData

- Optional location to return edge data

numEdges

- See description

Returns

cudaSuccess, cudaErrorLossyQuery, cudaErrorInvalidValue

Description

Returns a list of graph's dependency edges. Edges are returned via corresponding indices in from,

to and edgeData; that is, the node in to[i] has a dependency on the node in from[i] with data

edgeData[i]. from and to may both be NULL, in which case this function only returns the number

of edges in numEdges. Otherwise, numEdges entries will be filled in. If numEdges is higher

than the actual number of edges, the remaining entries in from and to will be set to NULL, and the

number of edges actually returned will be written to numEdges. edgeData may alone be NULL,

in which case the edges must all have default (zeroed) edge data. Attempting a losst query via NULL

edgeData will result in cudaErrorLossyQuery. If edgeData is non-NULL then from and to must

be as well.

Note:

‣ Graph objects are not threadsafe. More here.

359

Modules

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphGetNodes, cudaGraphGetRootNodes, cudaGraphAddDependencies,

cudaGraphRemoveDependencies, cudaGraphNodeGetDependencies,

cudaGraphNodeGetDependentNodes

__host__cudaError_t cudaGraphGetNodes (cudaGraph_t

graph, cudaGraphNode_t *nodes, size_t *numNodes)

Returns a graph's nodes.

Parameters

graph

- Graph to query

nodes

- Pointer to return the nodes

numNodes

- See description

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Returns a list of graph's nodes. nodes may be NULL, in which case this function will return the

number of nodes in numNodes. Otherwise, numNodes entries will be filled in. If numNodes is

higher than the actual number of nodes, the remaining entries in nodes will be set to NULL, and the

number of nodes actually obtained will be returned in numNodes.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

360

Modules

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphCreate, cudaGraphGetRootNodes, cudaGraphGetEdges, cudaGraphNodeGetType,

cudaGraphNodeGetDependencies, cudaGraphNodeGetDependentNodes

__host__cudaError_t cudaGraphGetRootNodes

(cudaGraph_t graph, cudaGraphNode_t *pRootNodes,

size_t *pNumRootNodes)

Returns a graph's root nodes.

Parameters

graph

- Graph to query

pRootNodes

- Pointer to return the root nodes

pNumRootNodes

- See description

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Returns a list of graph's root nodes. pRootNodes may be NULL, in which case this function will

return the number of root nodes in pNumRootNodes. Otherwise, pNumRootNodes entries will be

filled in. If pNumRootNodes is higher than the actual number of root nodes, the remaining entries

in pRootNodes will be set to NULL, and the number of nodes actually obtained will be returned in

pNumRootNodes.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

361

Modules

See also:

cudaGraphCreate, cudaGraphGetNodes, cudaGraphGetEdges, cudaGraphNodeGetType,

cudaGraphNodeGetDependencies, cudaGraphNodeGetDependentNodes

__host__cudaError_t cudaGraphHostNodeGetParams

(cudaGraphNode_t node, cudaHostNodeParams

*pNodeParams)

Returns a host node's parameters.

Parameters

node

- Node to get the parameters for

pNodeParams

- Pointer to return the parameters

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Returns the parameters of host node node in pNodeParams.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaLaunchHostFunc, cudaGraphAddHostNode, cudaGraphHostNodeSetParams

362

Modules

__host__cudaError_t cudaGraphHostNodeSetParams

(cudaGraphNode_t node, const cudaHostNodeParams

*pNodeParams)

Sets a host node's parameters.

Parameters

node

- Node to set the parameters for

pNodeParams

- Parameters to copy

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Sets the parameters of host node node to nodeParams.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphNodeSetParams, cudaLaunchHostFunc, cudaGraphAddHostNode,

cudaGraphHostNodeGetParams

363

Modules

__host__cudaError_t cudaGraphInstantiate

(cudaGraphExec_t *pGraphExec, cudaGraph_t graph,

unsigned long long flags)

Creates an executable graph from a graph.

Parameters

pGraphExec

- Returns instantiated graph

graph

- Graph to instantiate

flags

- Flags to control instantiation. See CUgraphInstantiate_flags.

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Instantiates graph as an executable graph. The graph is validated for any structural constraints or

intra-node constraints which were not previously validated. If instantiation is successful, a handle to the

instantiated graph is returned in pGraphExec.

The flags parameter controls the behavior of instantiation and subsequent graph launches. Valid

flags are:

‣ cudaGraphInstantiateFlagAutoFreeOnLaunch, which configures a graph containing memory

allocation nodes to automatically free any unfreed memory allocations before the graph is

relaunched.

‣ cudaGraphInstantiateFlagDeviceLaunch, which configures the graph for launch from the

device. If this flag is passed, the executable graph handle returned can be used to launch

the graph from both the host and device. This flag cannot be used in conjunction with

cudaGraphInstantiateFlagAutoFreeOnLaunch.

‣ cudaGraphInstantiateFlagUseNodePriority, which causes the graph to use the priorities from

the per-node attributes rather than the priority of the launch stream during execution. Note that

priorities are only available on kernel nodes, and are copied from stream priority during stream

capture.

If graph contains any allocation or free nodes, there can be at most one executable graph in existence

for that graph at a time. An attempt to instantiate a second executable graph before destroying the

first with cudaGraphExecDestroy will result in an error. The same also applies if graph contains any

device-updatable kernel nodes.

364

Modules

Graphs instantiated for launch on the device have additional restrictions which do not apply to host

graphs:

‣ The graph's nodes must reside on a single device.

‣ The graph can only contain kernel nodes, memcpy nodes, memset nodes, and child graph nodes.

‣ The graph cannot be empty and must contain at least one kernel, memcpy, or memset node.

Operation-specific restrictions are outlined below.

‣ Kernel nodes:

‣ Use of CUDA Dynamic Parallelism is not permitted.

‣ Cooperative launches are permitted as long as MPS is not in use.

‣ Memcpy nodes:

‣ Only copies involving device memory and/or pinned device-mapped host memory are

permitted.

‣ Copies involving CUDA arrays are not permitted.

‣ Both operands must be accessible from the current device, and the current device must match

the device of other nodes in the graph.

If graph is not instantiated for launch on the device but contains kernels which call device-side

cudaGraphLaunch() from multiple devices, this will result in an error.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphInstantiateWithFlags, cudaGraphCreate, cudaGraphUpload, cudaGraphLaunch,

cudaGraphExecDestroy

365

Modules

__host__cudaError_t cudaGraphInstantiateWithFlags

(cudaGraphExec_t *pGraphExec, cudaGraph_t graph,

unsigned long long flags)

Creates an executable graph from a graph.

Parameters

pGraphExec

- Returns instantiated graph

graph

- Graph to instantiate

flags

- Flags to control instantiation. See CUgraphInstantiate_flags.

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Instantiates graph as an executable graph. The graph is validated for any structural constraints or

intra-node constraints which were not previously validated. If instantiation is successful, a handle to the

instantiated graph is returned in pGraphExec.

The flags parameter controls the behavior of instantiation and subsequent graph launches. Valid

flags are:

‣ cudaGraphInstantiateFlagAutoFreeOnLaunch, which configures a graph containing memory

allocation nodes to automatically free any unfreed memory allocations before the graph is

relaunched.

‣ cudaGraphInstantiateFlagDeviceLaunch, which configures the graph for launch from the device. If

this flag is passed, the executable graph handle returned can be used to launch the graph from both

the host and device. This flag can only be used on platforms which support unified addressing. This

flag cannot be used in conjunction with cudaGraphInstantiateFlagAutoFreeOnLaunch.

‣ cudaGraphInstantiateFlagUseNodePriority, which causes the graph to use the priorities from

the per-node attributes rather than the priority of the launch stream during execution. Note that

priorities are only available on kernel nodes, and are copied from stream priority during stream

capture.

If graph contains any allocation or free nodes, there can be at most one executable graph in existence

for that graph at a time. An attempt to instantiate a second executable graph before destroying the

first with cudaGraphExecDestroy will result in an error. The same also applies if graph contains any

device-updatable kernel nodes.

366

Modules

If graph contains kernels which call device-side cudaGraphLaunch() from multiple devices, this will

result in an error.

Graphs instantiated for launch on the device have additional restrictions which do not apply to host

graphs:

‣ The graph's nodes must reside on a single device.

‣ The graph can only contain kernel nodes, memcpy nodes, memset nodes, and child graph nodes.

‣ The graph cannot be empty and must contain at least one kernel, memcpy, or memset node.

Operation-specific restrictions are outlined below.

‣ Kernel nodes:

‣ Use of CUDA Dynamic Parallelism is not permitted.

‣ Cooperative launches are permitted as long as MPS is not in use.

‣ Memcpy nodes:

‣ Only copies involving device memory and/or pinned device-mapped host memory are

permitted.

‣ Copies involving CUDA arrays are not permitted.

‣ Both operands must be accessible from the current device, and the current device must match

the device of other nodes in the graph.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphInstantiate, cudaGraphCreate, cudaGraphUpload, cudaGraphLaunch,

cudaGraphExecDestroy

367

Modules

__host__cudaError_t cudaGraphInstantiateWithParams

(cudaGraphExec_t *pGraphExec, cudaGraph_t graph,

cudaGraphInstantiateParams *instantiateParams)

Creates an executable graph from a graph.

Parameters

pGraphExec

- Returns instantiated graph

graph

- Graph to instantiate

instantiateParams

- Instantiation parameters

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Instantiates graph as an executable graph according to the instantiateParams structure. The

graph is validated for any structural constraints or intra-node constraints which were not previously

validated. If instantiation is successful, a handle to the instantiated graph is returned in pGraphExec.

instantiateParams controls the behavior of instantiation and subsequent graph launches, as well

as returning more detailed information in the event of an error. cudaGraphInstantiateParams is defined

as:

‎

typedef struct {

unsigned long long flags;

cudaStream_t uploadStream;

cudaGraphNode_t errNode_out;

cudaGraphInstantiateResult result_out;

} cudaGraphInstantiateParams;

The flags field controls the behavior of instantiation and subsequent graph launches. Valid flags are:

‣ cudaGraphInstantiateFlagAutoFreeOnLaunch, which configures a graph containing memory

allocation nodes to automatically free any unfreed memory allocations before the graph is

relaunched.

‣ cudaGraphInstantiateFlagUpload, which will perform an upload of the graph into

uploadStream once the graph has been instantiated.

‣ cudaGraphInstantiateFlagDeviceLaunch, which configures the graph for launch from the device. If

this flag is passed, the executable graph handle returned can be used to launch the graph from both

the host and device. This flag can only be used on platforms which support unified addressing. This

flag cannot be used in conjunction with cudaGraphInstantiateFlagAutoFreeOnLaunch.

368

Modules

‣ cudaGraphInstantiateFlagUseNodePriority, which causes the graph to use the priorities from

the per-node attributes rather than the priority of the launch stream during execution. Note that

priorities are only available on kernel nodes, and are copied from stream priority during stream

capture.

If graph contains any allocation or free nodes, there can be at most one executable graph in existence

for that graph at a time. An attempt to instantiate a second executable graph before destroying the

first with cudaGraphExecDestroy will result in an error. The same also applies if graph contains any

device-updatable kernel nodes.

If graph contains kernels which call device-side cudaGraphLaunch() from multiple devices, this will

result in an error.

Graphs instantiated for launch on the device have additional restrictions which do not apply to host

graphs:

‣ The graph's nodes must reside on a single device.

‣ The graph can only contain kernel nodes, memcpy nodes, memset nodes, and child graph nodes.

‣ The graph cannot be empty and must contain at least one kernel, memcpy, or memset node.

Operation-specific restrictions are outlined below.

‣ Kernel nodes:

‣ Use of CUDA Dynamic Parallelism is not permitted.

‣ Cooperative launches are permitted as long as MPS is not in use.

‣ Memcpy nodes:

‣ Only copies involving device memory and/or pinned device-mapped host memory are

permitted.

‣ Copies involving CUDA arrays are not permitted.

‣ Both operands must be accessible from the current device, and the current device must match

the device of other nodes in the graph.

In the event of an error, the result_out and errNode_out fields will contain more information

about the nature of the error. Possible error reporting includes:

‣ cudaGraphInstantiateError, if passed an invalid value or if an unexpected error occurred which is

described by the return value of the function. errNode_out will be set to NULL.

‣ cudaGraphInstantiateInvalidStructure, if the graph structure is invalid. errNode_out will be set

to one of the offending nodes.

‣ cudaGraphInstantiateNodeOperationNotSupported, if the graph is instantiated for device

launch but contains a node of an unsupported node type, or a node which performs unsupported

operations, such as use of CUDA dynamic parallelism within a kernel node. errNode_out will

be set to this node.

‣ cudaGraphInstantiateMultipleDevicesNotSupported, if the graph is instantiated for device

launch but a node’s device differs from that of another node. This error can also be returned

369

Modules

if a graph is not instantiated for device launch and it contains kernels which call device-side

cudaGraphLaunch() from multiple devices. errNode_out will be set to this node.

If instantiation is successful, result_out will be set to cudaGraphInstantiateSuccess, and

hErrNode_out will be set to NULL.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphCreate, cudaGraphInstantiate, cudaGraphInstantiateWithFlags, cudaGraphExecDestroy

__host__cudaError_t

cudaGraphKernelNodeCopyAttributes (cudaGraphNode_t

hSrc, cudaGraphNode_t hDst)

Copies attributes from source node to destination node.

Returns

cudaSuccess, cudaErrorInvalidContext

Description

Copies attributes from source node src to destination node dst. Both node must have the same

context.

Note:

Note that this function may also return error codes from previous, asynchronous launches.

See also:

cudaAccessPolicyWindow

370

Modules

__host__cudaError_t cudaGraphKernelNodeGetAttribute

(cudaGraphNode_t hNode, cudaKernelNodeAttrID attr,

cudaKernelNodeAttrValue *value_out)

Queries node attribute.

Parameters

hNode

attr

value_out

Returns

cudaSuccess, cudaErrorInvalidValue, cudaErrorInvalidResourceHandle

Description

Queries attribute attr from node hNode and stores it in corresponding member of value_out.

Note:

Note that this function may also return error codes from previous, asynchronous launches.

See also:

cudaAccessPolicyWindow

__host__cudaError_t cudaGraphKernelNodeGetParams

(cudaGraphNode_t node, cudaKernelNodeParams

*pNodeParams)

Returns a kernel node's parameters.

Parameters

node

- Node to get the parameters for

pNodeParams

- Pointer to return the parameters

Returns

cudaSuccess, cudaErrorInvalidValue, cudaErrorInvalidDeviceFunction

371

Modules

Description

Returns the parameters of kernel node node in pNodeParams. The kernelParams or extra

array returned in pNodeParams, as well as the argument values it points to, are owned by the node.

This memory remains valid until the node is destroyed or its parameters are modified, and should not

be modified directly. Use cudaGraphKernelNodeSetParams to update the parameters of this node.

The params will contain either kernelParams or extra, according to which of these was most

recently set on the node.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaLaunchKernel, cudaGraphAddKernelNode, cudaGraphKernelNodeSetParams

__host__cudaError_t cudaGraphKernelNodeSetAttribute

(cudaGraphNode_t hNode, cudaKernelNodeAttrID attr,

const cudaKernelNodeAttrValue *value)

Sets node attribute.

Parameters

hNode

attr

value

Returns

cudaSuccess, cudaErrorInvalidValue, cudaErrorInvalidResourceHandle

Description

Sets attribute attr on node hNode from corresponding attribute of value.

372

Modules

Note:

Note that this function may also return error codes from previous, asynchronous launches.

See also:

cudaAccessPolicyWindow

__device__cudaError_t cudaGraphKernelNodeSetEnabled

(cudaGraphDeviceNode_t node, bool enable)

Enables or disables the given kernel node.

Parameters

node

- The node to update

enable

- Whether to enable or disable the node

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Enables or disables node based upon enable. If enable is true, the node will be enabled; if it is

false, the node will be disabled. Disabled nodes will act as a NOP during execution. node must be

device-updatable, and must reside upon the same device as the calling kernel.

If this function is called for the node's immediate dependent and that dependent is configured for

programmatic dependent launch, then a memory fence must be invoked via __threadfence() before

kickoff of the dependent is triggered via cudaTriggerProgrammaticLaunchCompletion() to ensure that

the update is visible to that dependent node before it is launched.

Note:

Note that this function may also return error codes from previous, asynchronous launches.

See also:

cudaGraphKernelNodeSetParam, cudaGraphKernelNodeSetGridDim,

cudaGraphKernelNodeUpdatesApply

373

Modules

__device__cudaError_t cudaGraphKernelNodeSetGridDim

(cudaGraphDeviceNode_t node, dim3 gridDim)

Updates the grid dimensions of the given kernel node.

Parameters

node

- The node to update

gridDim

- The grid dimensions to set

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Sets the grid dimensions of node to gridDim. node must be device-updatable, and must reside upon

the same device as thecalling kernel.

If this function is called for the node's immediate dependent and that dependent is configured for

programmatic dependent launch, then a memory fence must be invoked via __threadfence() before

kickoff of the dependent is triggered via cudaTriggerProgrammaticLaunchCompletion() to ensure that

the update is visible to that dependent node before it is launched.

Note:

Note that this function may also return error codes from previous, asynchronous launches.

See also:

cudaGraphKernelNodeSetParam, cudaGraphKernelNodeSetEnabled,

cudaGraphKernelNodeUpdatesApply

template < typename T > __device__cudaError_t

cudaGraphKernelNodeSetParam (cudaGraphDeviceNode_t

node, size_t offset, const T value)

Updates the kernel parameters of the given kernel node.

Parameters

node

- The node to update

374

Modules

offset

- The offset into the params at which to make the update

value

- Parameter value to write

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Updates the kernel parameters of node at offset to value. node must be device-updatable, and

must reside upon the same device as the calling kernel.

If this function is called for the node's immediate dependent and that dependent is configured for

programmatic dependent launch, then a memory fence must be invoked via __threadfence() before

kickoff of the dependent is triggered via cudaTriggerProgrammaticLaunchCompletion() to ensure that

the update is visible to that dependent node before it is launched.

Note:

Note that this function may also return error codes from previous, asynchronous launches.

See also:

cudaGraphKernelNodeSetEnabled, cudaGraphKernelNodeSetGridDim,

cudaGraphKernelNodeUpdatesApply

__device__cudaError_t cudaGraphKernelNodeSetParam

(cudaGraphDeviceNode_t node, size_t offset, const void

*value, size_t size)

Updates the kernel parameters of the given kernel node.

Parameters

node

- The node to update

offset

- The offset into the params at which to make the update

value

- Buffer containing the params to write

size

- Size in bytes to update

375

Modules

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Updates size bytes in the kernel parameters of node at offset to the contents of value. node

must be device-updatable, and must reside upon the same device as the calling kernel.

If this function is called for the node's immediate dependent and that dependent is configured for

programmatic dependent launch, then a memory fence must be invoked via __threadfence() before

kickoff of the dependent is triggered via cudaTriggerProgrammaticLaunchCompletion() to ensure that

the update is visible to that dependent node before it is launched.

Note:

Note that this function may also return error codes from previous, asynchronous launches.

See also:

cudaGraphKernelNodeSetEnabled, cudaGraphKernelNodeSetGridDim,

cudaGraphKernelNodeUpdatesApply

__host__cudaError_t cudaGraphKernelNodeSetParams

(cudaGraphNode_t node, const cudaKernelNodeParams

*pNodeParams)

Sets a kernel node's parameters.

Parameters

node

- Node to set the parameters for

pNodeParams

- Parameters to copy

Returns

cudaSuccess, cudaErrorInvalidValue, cudaErrorInvalidResourceHandle, cudaErrorMemoryAllocation

Description

Sets the parameters of kernel node node to pNodeParams.

376

Modules

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

‣ The API can also be used with a kernel cudaKernel_t by querying the handle using

cudaLibraryGetKernel() or cudaGetKernel and then passing it to the API by casting to void*. The

symbol entryFuncAddr passed to cudaGetKernel should be a symbol that is registered with the

same CUDA Runtime instance.

‣ Passing a symbol that belongs that belongs to a different runtime instance will result in undefined

behavior. The only type that can be reliably passed to a different runtime instance is cudaKernel_t

See also:

cudaGraphNodeSetParams, cudaLaunchKernel, cudaGraphAddKernelNode,

cudaGraphKernelNodeGetParams

__device__cudaError_t

cudaGraphKernelNodeUpdatesApply (const

cudaGraphKernelNodeUpdate *updates, size_t

updateCount)

Batch applies multiple kernel node updates.

Parameters

updates

- The updates to apply

updateCount

- The number of updates to apply

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Batch applies one or more kernel node updates based on the information provided in updates.

updateCount specifies the number of updates to apply. Each entry in updates must specify

377

Modules

a node to update, the type of update to apply, and the parameters for that type of update. See the

documentation for cudaGraphKernelNodeUpdate for more detail.

If this function is called for the node's immediate dependent and that dependent is configured for

programmatic dependent launch, then a memory fence must be invoked via __threadfence() before

kickoff of the dependent is triggered via cudaTriggerProgrammaticLaunchCompletion() to ensure that

the update is visible to that dependent node before it is launched.

Note:

Note that this function may also return error codes from previous, asynchronous launches.

See also:

cudaGraphKernelNodeSetParam, cudaGraphKernelNodeSetEnabled,

cudaGraphKernelNodeSetGridDim

__host____

device__cudaError_t cudaGraphLaunch

(cudaGraphExec_t graphExec, cudaStream_t stream)

Launches an executable graph in a stream.

Parameters

graphExec

- Executable graph to launch

stream

- Stream in which to launch the graph

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Executes graphExec in stream. Only one instance of graphExec may be executing at a time.

Each launch is ordered behind both any previous work in stream and any previous launches of

graphExec. To execute a graph concurrently, it must be instantiated multiple times into multiple

executable graphs.

If any allocations created by graphExec remain unfreed (from a previous launch) and graphExec

was not instantiated with cudaGraphInstantiateFlagAutoFreeOnLaunch, the launch will fail with

cudaErrorInvalidValue.

Note:

378

Modules

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphInstantiate, cudaGraphUpload, cudaGraphExecDestroy

__host__cudaError_t cudaGraphMemAllocNodeGetParams

(cudaGraphNode_t node, cudaMemAllocNodeParams

*params_out)

Returns a memory alloc node's parameters.

Parameters

node

- Node to get the parameters for

params_out

- Pointer to return the parameters

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Returns the parameters of a memory alloc node hNode in params_out. The poolProps and

accessDescs returned in params_out, are owned by the node. This memory remains valid until

the node is destroyed. The returned parameters must not be modified.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

379

Modules

See also:

cudaGraphAddMemAllocNode, cudaGraphMemFreeNodeGetParams

__host__cudaError_t cudaGraphMemcpyNodeGetParams

(cudaGraphNode_t node, cudaMemcpy3DParms

*pNodeParams)

Returns a memcpy node's parameters.

Parameters

node

- Node to get the parameters for

pNodeParams

- Pointer to return the parameters

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Returns the parameters of memcpy node node in pNodeParams.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaMemcpy3D, cudaGraphAddMemcpyNode, cudaGraphMemcpyNodeSetParams

380

Modules

__host__cudaError_t cudaGraphMemcpyNodeSetParams

(cudaGraphNode_t node, const cudaMemcpy3DParms

*pNodeParams)

Sets a memcpy node's parameters.

Parameters

node

- Node to set the parameters for

pNodeParams

- Parameters to copy

Returns

cudaSuccess, cudaErrorInvalidValue,

Description

Sets the parameters of memcpy node node to pNodeParams.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphNodeSetParams, cudaMemcpy3D, cudaGraphMemcpyNodeSetParamsToSymbol,

cudaGraphMemcpyNodeSetParamsFromSymbol, cudaGraphMemcpyNodeSetParams1D,

cudaGraphAddMemcpyNode, cudaGraphMemcpyNodeGetParams

381

Modules

__host__cudaError_t

cudaGraphMemcpyNodeSetParams1D (cudaGraphNode_t

node, void *dst, const void *src, size_t count,

cudaMemcpyKind kind)

Sets a memcpy node's parameters to perform a 1-dimensional copy.

Parameters

node

- Node to set the parameters for

dst

- Destination memory address

src

- Source memory address

count

- Size in bytes to copy

kind

- Type of transfer

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Sets the parameters of memcpy node node to the copy described by the provided parameters.

When the graph is launched, the node will copy count bytes from the memory area pointed to by

src to the memory area pointed to by dst, where kind specifies the direction of the copy, and

must be one of cudaMemcpyHostToHost, cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost,

cudaMemcpyDeviceToDevice, or cudaMemcpyDefault. Passing cudaMemcpyDefault is

recommended, in which case the type of transfer is inferred from the pointer values. However,

cudaMemcpyDefault is only allowed on systems that support unified virtual addressing. Launching

a memcpy node with dst and src pointers that do not match the direction of the copy results in an

undefined behavior.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

382

Modules

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaMemcpy, cudaGraphMemcpyNodeSetParams, cudaGraphAddMemcpyNode,

cudaGraphMemcpyNodeGetParams

__host__cudaError_t

cudaGraphMemcpyNodeSetParamsFromSymbol

(cudaGraphNode_t node, void *dst, const void *symbol,

size_t count, size_t offset, cudaMemcpyKind kind)

Sets a memcpy node's parameters to copy from a symbol on the device.

Parameters

node

- Node to set the parameters for

dst

- Destination memory address

symbol

- Device symbol address

count

- Size in bytes to copy

offset

- Offset from start of symbol in bytes

kind

- Type of transfer

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Sets the parameters of memcpy node node to the copy described by the provided parameters.

When the graph is launched, the node will copy count bytes from the memory area pointed to by

offset bytes from the start of symbol symbol to the memory area pointed to by dst. The memory

areas may not overlap. symbol is a variable that resides in global or constant memory space. kind

can be either cudaMemcpyDeviceToHost, cudaMemcpyDeviceToDevice, or cudaMemcpyDefault.

Passing cudaMemcpyDefault is recommended, in which case the type of transfer is inferred from the

pointer values. However, cudaMemcpyDefault is only allowed on systems that support unified virtual

addressing.

383

Modules

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaMemcpyFromSymbol, cudaGraphMemcpyNodeSetParams,

cudaGraphMemcpyNodeSetParamsToSymbol, cudaGraphAddMemcpyNode,

cudaGraphMemcpyNodeGetParams

__host__cudaError_t

cudaGraphMemcpyNodeSetParamsToSymbol

(cudaGraphNode_t node, const void *symbol, const void

*src, size_t count, size_t offset, cudaMemcpyKind kind)

Sets a memcpy node's parameters to copy to a symbol on the device.

Parameters

node

- Node to set the parameters for

symbol

- Device symbol address

src

- Source memory address

count

- Size in bytes to copy

offset

- Offset from start of symbol in bytes

kind

- Type of transfer

Returns

cudaSuccess, cudaErrorInvalidValue

384

Modules

Description

Sets the parameters of memcpy node node to the copy described by the provided parameters.

When the graph is launched, the node will copy count bytes from the memory area pointed to by

src to the memory area pointed to by offset bytes from the start of symbol symbol. The memory

areas may not overlap. symbol is a variable that resides in global or constant memory space. kind

can be either cudaMemcpyHostToDevice, cudaMemcpyDeviceToDevice, or cudaMemcpyDefault.

Passing cudaMemcpyDefault is recommended, in which case the type of transfer is inferred from the

pointer values. However, cudaMemcpyDefault is only allowed on systems that support unified virtual

addressing.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaMemcpyToSymbol, cudaGraphMemcpyNodeSetParams,

cudaGraphMemcpyNodeSetParamsFromSymbol, cudaGraphAddMemcpyNode,

cudaGraphMemcpyNodeGetParams

__host__cudaError_t cudaGraphMemFreeNodeGetParams

(cudaGraphNode_t node, void *dptr_out)

Returns a memory free node's parameters.

Parameters

node

- Node to get the parameters for

dptr_out

- Pointer to return the device address

Returns

cudaSuccess, cudaErrorInvalidValue

385

Modules

Description

Returns the address of a memory free node hNode in dptr_out.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphAddMemFreeNode, cudaGraphMemFreeNodeGetParams

__host__cudaError_t cudaGraphMemsetNodeGetParams

(cudaGraphNode_t node, cudaMemsetParams

*pNodeParams)

Returns a memset node's parameters.

Parameters

node

- Node to get the parameters for

pNodeParams

- Pointer to return the parameters

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Returns the parameters of memset node node in pNodeParams.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

386

Modules

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaMemset2D, cudaGraphAddMemsetNode, cudaGraphMemsetNodeSetParams

__host__cudaError_t cudaGraphMemsetNodeSetParams

(cudaGraphNode_t node, const cudaMemsetParams

*pNodeParams)

Sets a memset node's parameters.

Parameters

node

- Node to set the parameters for

pNodeParams

- Parameters to copy

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Sets the parameters of memset node node to pNodeParams.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphNodeSetParams, cudaMemset2D, cudaGraphAddMemsetNode,

cudaGraphMemsetNodeGetParams

387

Modules

__host__cudaError_t cudaGraphNodeFindInClone

(cudaGraphNode_t *pNode, cudaGraphNode_t

originalNode, cudaGraph_t clonedGraph)

Finds a cloned version of a node.

Parameters

pNode

- Returns handle to the cloned node

originalNode

- Handle to the original node

clonedGraph

- Cloned graph to query

Returns

cudaSuccess, cudaErrorInvalidValue

Description

This function returns the node in clonedGraph corresponding to originalNode in the original

graph.

clonedGraph must have been cloned from originalGraph via cudaGraphClone.

originalNode must have been in originalGraph at the time of the call to cudaGraphClone, and

the corresponding cloned node in clonedGraph must not have been removed. The cloned node is

then returned via pClonedNode.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphClone

388

Modules

__host__cudaError_t cudaGraphNodeGetDependencies

(cudaGraphNode_t node, cudaGraphNode_t

*pDependencies, size_t *pNumDependencies)

Returns a node's dependencies.

Parameters

node

- Node to query

pDependencies

- Pointer to return the dependencies

pNumDependencies

- See description

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Returns a list of node's dependencies. pDependencies may be NULL, in which case this function

will return the number of dependencies in pNumDependencies. Otherwise, pNumDependencies

entries will be filled in. If pNumDependencies is higher than the actual number of dependencies,

the remaining entries in pDependencies will be set to NULL, and the number of nodes actually

obtained will be returned in pNumDependencies.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphNodeGetDependentNodes, cudaGraphGetNodes, cudaGraphGetRootNodes,

cudaGraphGetEdges, cudaGraphAddDependencies, cudaGraphRemoveDependencies

389

Modules

__host__cudaError_t cudaGraphNodeGetDependencies_v2

(cudaGraphNode_t node, cudaGraphNode_t

*pDependencies, cudaGraphEdgeData *edgeData, size_t

*pNumDependencies)

Returns a node's dependencies (12.3+).

Parameters

node

- Node to query

pDependencies

- Pointer to return the dependencies

edgeData

- Optional array to return edge data for each dependency

pNumDependencies

- See description

Returns

cudaSuccess, cudaErrorLossyQuery, cudaErrorInvalidValue

Description

Returns a list of node's dependencies. pDependencies may be NULL, in which case this function

will return the number of dependencies in pNumDependencies. Otherwise, pNumDependencies

entries will be filled in. If pNumDependencies is higher than the actual number of dependencies,

the remaining entries in pDependencies will be set to NULL, and the number of nodes actually

obtained will be returned in pNumDependencies.

Note that if an edge has non-zero (non-default) edge data and edgeData is NULL, this API will

return cudaErrorLossyQuery. If edgeData is non-NULL, then pDependencies must be as well.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

390

Modules

cudaGraphNodeGetDependentNodes, cudaGraphGetNodes, cudaGraphGetRootNodes,

cudaGraphGetEdges, cudaGraphAddDependencies, cudaGraphRemoveDependencies

__host__cudaError_t cudaGraphNodeGetDependentNodes

(cudaGraphNode_t node, cudaGraphNode_t

*pDependentNodes, size_t *pNumDependentNodes)

Returns a node's dependent nodes.

Parameters

node

- Node to query

pDependentNodes

- Pointer to return the dependent nodes

pNumDependentNodes

- See description

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Returns a list of node's dependent nodes. pDependentNodes may be NULL, in which case

this function will return the number of dependent nodes in pNumDependentNodes. Otherwise,

pNumDependentNodes entries will be filled in. If pNumDependentNodes is higher than the

actual number of dependent nodes, the remaining entries in pDependentNodes will be set to NULL,

and the number of nodes actually obtained will be returned in pNumDependentNodes.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphNodeGetDependencies, cudaGraphGetNodes, cudaGraphGetRootNodes,

cudaGraphGetEdges, cudaGraphAddDependencies, cudaGraphRemoveDependencies

391

Modules

__host__cudaError_t

cudaGraphNodeGetDependentNodes_v2

(cudaGraphNode_t node, cudaGraphNode_t

*pDependentNodes, cudaGraphEdgeData *edgeData,

size_t *pNumDependentNodes)

Returns a node's dependent nodes (12.3+).

Parameters

node

- Node to query

pDependentNodes

- Pointer to return the dependent nodes

edgeData

- Optional pointer to return edge data for dependent nodes

pNumDependentNodes

- See description

Returns

cudaSuccess, cudaErrorLossyQuery, cudaErrorInvalidValue

Description

Returns a list of node's dependent nodes. pDependentNodes may be NULL, in which case

this function will return the number of dependent nodes in pNumDependentNodes. Otherwise,

pNumDependentNodes entries will be filled in. If pNumDependentNodes is higher than the

actual number of dependent nodes, the remaining entries in pDependentNodes will be set to NULL,

and the number of nodes actually obtained will be returned in pNumDependentNodes.

Note that if an edge has non-zero (non-default) edge data and edgeData is NULL, this API will

return cudaErrorLossyQuery. If edgeData is non-NULL, then pDependentNodes must be as

well.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

392

Modules

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphNodeGetDependencies, cudaGraphGetNodes, cudaGraphGetRootNodes,

cudaGraphGetEdges, cudaGraphAddDependencies, cudaGraphRemoveDependencies

__host__cudaError_t cudaGraphNodeGetEnabled

(cudaGraphExec_t hGraphExec, cudaGraphNode_t hNode,

unsigned int *isEnabled)

Query whether a node in the given graphExec is enabled.

Parameters

hGraphExec

- The executable graph in which to set the specified node

hNode

- Node from the graph from which graphExec was instantiated

isEnabled

- Location to return the enabled status of the node

Returns

cudaSuccess, cudaErrorInvalidValue,

Description

Sets isEnabled to 1 if hNode is enabled, or 0 if hNode is disabled.

The node is identified by the corresponding node hNode in the non-executable graph, from which the

executable graph was instantiated.

hNode must not have been removed from the original graph.

Note:

Currently only kernel, memset and memcpy nodes are supported.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

393

Modules

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphNodeSetEnabled, cudaGraphExecUpdate, cudaGraphInstantiate cudaGraphLaunch

__host__cudaError_t cudaGraphNodeGetType

(cudaGraphNode_t node, cudaGraphNodeType *pType)

Returns a node's type.

Parameters

node

- Node to query

pType

- Pointer to return the node type

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Returns the node type of node in pType.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphGetNodes, cudaGraphGetRootNodes, cudaGraphChildGraphNodeGetGraph,

cudaGraphKernelNodeGetParams, cudaGraphKernelNodeSetParams, cudaGraphHostNodeGetParams,

cudaGraphHostNodeSetParams, cudaGraphMemcpyNodeGetParams,

cudaGraphMemcpyNodeSetParams, cudaGraphMemsetNodeGetParams,

cudaGraphMemsetNodeSetParams

394

Modules

__host__cudaError_t cudaGraphNodeSetEnabled

(cudaGraphExec_t hGraphExec, cudaGraphNode_t hNode,

unsigned int isEnabled)

Enables or disables the specified node in the given graphExec.

Parameters

hGraphExec

- The executable graph in which to set the specified node

hNode

- Node from the graph from which graphExec was instantiated

isEnabled

- Node is enabled if != 0, otherwise the node is disabled

Returns

cudaSuccess, cudaErrorInvalidValue,

Description

Sets hNode to be either enabled or disabled. Disabled nodes are functionally equivalent to empty

nodes until they are reenabled. Existing node parameters are not affected by disabling/enabling the

node.

The node is identified by the corresponding node hNode in the non-executable graph, from which the

executable graph was instantiated.

hNode must not have been removed from the original graph.

The modifications only affect future launches of hGraphExec. Already enqueued or running

launches of hGraphExec are not affected by this call. hNode is also not modified by this call.

Note:

Currently only kernel, memset and memcpy nodes are supported.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

395

Modules

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphNodeGetEnabled, cudaGraphExecUpdate, cudaGraphInstantiate cudaGraphLaunch

__host__cudaError_t cudaGraphNodeSetParams

(cudaGraphNode_t node, cudaGraphNodeParams

*nodeParams)

Update's a graph node's parameters.

Parameters

node

- Node to set the parameters for

nodeParams

- Parameters to copy

Returns

cudaSuccess, cudaErrorInvalidValue, cudaErrorInvalidDeviceFunction, cudaErrorNotSupported

Description

Sets the parameters of graph node node to nodeParams. The node type specified by

nodeParams->type must match the type of node. nodeParams must be fully initialized and all

unused bytes (reserved, padding) zeroed.

Modifying parameters is not supported for node types cudaGraphNodeTypeMemAlloc and

cudaGraphNodeTypeMemFree.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphAddNode, cudaGraphExecNodeSetParams

396

Modules

__host__cudaError_t cudaGraphReleaseUserObject

(cudaGraph_t graph, cudaUserObject_t object, unsigned int

count)

Release a user object reference from a graph.

Parameters

graph

- The graph that will release the reference

object

- The user object to release a reference for

count

- The number of references to release, typically 1. Must be nonzero and not larger than INT_MAX.

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Releases user object references owned by a graph.

See CUDA User Objects in the CUDA C++ Programming Guide for more information on user objects.

See also:

cudaUserObjectCreate cudaUserObjectRetain, cudaUserObjectRelease, cudaGraphRetainUserObject,

cudaGraphCreate

__host__cudaError_t cudaGraphRemoveDependencies

(cudaGraph_t graph, const cudaGraphNode_t *from, const

cudaGraphNode_t *to, size_t numDependencies)

Removes dependency edges from a graph.

Parameters

graph

- Graph from which to remove dependencies

from

- Array of nodes that provide the dependencies

- Array of dependent nodes

397

Modules

numDependencies

- Number of dependencies to be removed

Returns

cudaSuccess, cudaErrorInvalidValue

Description

The number of pDependencies to be removed is defined by numDependencies. Elements in

pFrom and pTo at corresponding indices define a dependency. Each node in pFrom and pTo must

belong to graph.

If numDependencies is 0, elements in pFrom and pTo will be ignored. Specifying a non-existing

dependency will return an error.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphAddDependencies, cudaGraphGetEdges, cudaGraphNodeGetDependencies,

cudaGraphNodeGetDependentNodes

__host__cudaError_t cudaGraphRemoveDependencies_v2

(cudaGraph_t graph, const cudaGraphNode_t *from,

const cudaGraphNode_t *to, const cudaGraphEdgeData

*edgeData, size_t numDependencies)

Removes dependency edges from a graph. (12.3+).

Parameters

graph

- Graph from which to remove dependencies

from

- Array of nodes that provide the dependencies

398

Modules

- Array of dependent nodes

edgeData

- Optional array of edge data. If NULL, edge data is assumed to be default (zeroed).

numDependencies

- Number of dependencies to be removed

Returns

cudaSuccess, cudaErrorInvalidValue

Description

The number of pDependencies to be removed is defined by numDependencies. Elements in

pFrom and pTo at corresponding indices define a dependency. Each node in pFrom and pTo must

belong to graph.

If numDependencies is 0, elements in pFrom and pTo will be ignored. Specifying an edge that

does not exist in the graph, with data matching edgeData, results in an error. edgeData is nullable,

which is equivalent to passing default (zeroed) data for each edge.

Note:

‣ Graph objects are not threadsafe. More here.

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cudaGraphAddDependencies, cudaGraphGetEdges, cudaGraphNodeGetDependencies,

cudaGraphNodeGetDependentNodes

__host__cudaError_t cudaGraphRetainUserObject

(cudaGraph_t graph, cudaUserObject_t object, unsigned int

count, unsigned int flags)

Retain a reference to a user object from a graph.

Parameters

graph

- The graph to associate the reference with

399

Modules

object

- The user object to retain a reference for

count

- The number of references to add to the graph, typically 1. Must be nonzero and not larger than

INT_MAX.

flags

- The optional flag cudaGraphUserObjectMove transfers references from the calling thread, rather

than create new references. Pass 0 to create new references.

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Creates or moves user object references that will be owned by a CUDA graph.

See CUDA User Objects in the CUDA C++ Programming Guide for more information on user objects.

See also:

cudaUserObjectCreate cudaUserObjectRetain, cudaUserObjectRelease, cudaGraphReleaseUserObject,

cudaGraphCreate

__device__ void cudaGraphSetConditional

(cudaGraphConditionalHandle handle, unsigned int value)

Sets the condition value associated with a conditional node.

Description

Sets the condition value associated with a conditional node.

Note: handle must be associated with the same context as the kernel calling this function.

See also:

cudaGraphConditionalHandleCreate

__host__cudaError_t cudaGraphUpload (cudaGraphExec_t

graphExec, cudaStream_t stream)

Uploads an executable graph in a stream.

Returns

cudaSuccess, cudaErrorInvalidValue,

400

Modules

Description

Uploads hGraphExec to the device in hStream without executing it. Uploads of the same

hGraphExec will be serialized. Each upload is ordered behind both any previous work in hStream

and any previous launches of hGraphExec. Uses memory cached by stream to back the allocations

owned by graphExec.

Note:

‣ Note that this function may also return error codes from previous, asynchronous launches.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

See also:

cudaGraphInstantiate, cudaGraphLaunch, cudaGraphExecDestroy

__host__cudaError_t cudaUserObjectCreate

(cudaUserObject_t *object_out, void *ptr, cudaHostFn_t

destroy, unsigned int initialRefcount, unsigned int flags)

Create a user object.

Parameters

object_out

- Location to return the user object handle

ptr

- The pointer to pass to the destroy function

destroy

- Callback to free the user object when it is no longer in use

initialRefcount

- The initial refcount to create the object with, typically 1. The initial references are owned by the

calling thread.

flags

- Currently it is required to pass cudaUserObjectNoDestructorSync, which is the only defined flag.

This indicates that the destroy callback cannot be waited on by any CUDA API. Users requiring

synchronization of the callback should signal its completion manually.

Returns

cudaSuccess, cudaErrorInvalidValue

401

Modules

Description

Create a user object with the specified destructor callback and initial reference count. The initial

references are owned by the caller.

Destructor callbacks cannot make CUDA API calls and should avoid blocking behavior, as they are

executed by a shared internal thread. Another thread may be signaled to perform such actions, if it does

not block forward progress of tasks scheduled through CUDA.

See CUDA User Objects in the CUDA C++ Programming Guide for more information on user objects.

See also:

cudaUserObjectRetain, cudaUserObjectRelease, cudaGraphRetainUserObject,

cudaGraphReleaseUserObject, cudaGraphCreate

__host__cudaError_t cudaUserObjectRelease

(cudaUserObject_t object, unsigned int count)

Release a reference to a user object.

Parameters

object

- The object to release

count

- The number of references to release, typically 1. Must be nonzero and not larger than INT_MAX.

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Releases user object references owned by the caller. The object's destructor is invoked if the reference

count reaches zero.

It is undefined behavior to release references not owned by the caller, or to use a user object handle

after all references are released.

See CUDA User Objects in the CUDA C++ Programming Guide for more information on user objects.

See also:

cudaUserObjectCreate, cudaUserObjectRetain, cudaGraphRetainUserObject,

cudaGraphReleaseUserObject, cudaGraphCreate

402

Modules

__host__cudaError_t cudaUserObjectRetain

(cudaUserObject_t object, unsigned int count)

Retain a reference to a user object.

Parameters

object

- The object to retain

count

- The number of references to retain, typically 1. Must be nonzero and not larger than INT_MAX.

Returns

cudaSuccess, cudaErrorInvalidValue

Description

Retains new references to a user object. The new references are owned by the caller.

See CUDA User Objects in the CUDA C++ Programming Guide for more information on user objects.

See also:

cudaUserObjectCreate, cudaUserObjectRelease, cudaGraphRetainUserObject,

cudaGraphReleaseUserObject, cudaGraphCreate

6.31. Driver Entry Point Access

This section describes the driver entry point access functions of CUDA runtime application

programming interface.

__host__cudaError_t cudaGetDriverEntryPoint (const

char *symbol, void **funcPtr, unsigned long long flags,

cudaDriverEntryPointQueryResult *driverStatus)

Returns the requested driver API function pointer.

Parameters

symbol

- The base name of the driver API function to look for. As an example, for the driver API

cuMemAlloc_v2, symbol would be cuMemAlloc. Note that the API will use the CUDA runtime

403

Modules

version to return the address to the most recent ABI compatible driver symbol, cuMemAlloc or

cuMemAlloc_v2.

funcPtr

- Location to return the function pointer to the requested driver function

flags

- Flags to specify search options.

driverStatus

- Optional location to store the status of finding the symbol from the driver. See

cudaDriverEntryPointQueryResult for possible values.

Returns

cudaSuccess, cudaErrorInvalidValue, cudaErrorNotSupported

Description

Returns in **funcPtr the address of the CUDA driver function for the requested flags.

For a requested driver symbol, if the CUDA version in which the driver symbol was introduced

is less than or equal to the CUDA runtime version, the API will return the function pointer to the

corresponding versioned driver function.

The pointer returned by the API should be cast to a function pointer matching the requested driver

function's definition in the API header file. The function pointer typedef can be picked up from the

corresponding typedefs header file. For example, cudaTypedefs.h consists of function pointer typedefs

for driver APIs defined in cuda.h.

The API will return cudaSuccess and set the returned funcPtr if the requested driver function is valid

and supported on the platform.

The API will return cudaSuccess and set the returned funcPtr to NULL if the requested driver

function is not supported on the platform, no ABI compatible driver function exists for the CUDA

runtime version or if the driver symbol is invalid.

It will also set the optional driverStatus to one of the values in cudaDriverEntryPointQueryResult

with the following meanings:

‣ cudaDriverEntryPointSuccess - The requested symbol was succesfully found based on input

arguments and pfn is valid

‣ cudaDriverEntryPointSymbolNotFound - The requested symbol was not found

‣ cudaDriverEntryPointVersionNotSufficent - The requested symbol was found but is not supported

by the current runtime version (CUDART_VERSION)

The requested flags can be:

‣ cudaEnableDefault: This is the default mode. This is equivalent to

cudaEnablePerThreadDefaultStream if the code is compiled with --default-stream per-thread

compilation flag or the macro CUDA_API_PER_THREAD_DEFAULT_STREAM is defined;

cudaEnableLegacyStream otherwise.

404

Modules

‣ cudaEnableLegacyStream: This will enable the search for all driver symbols that match the

requested driver symbol name except the corresponding per-thread versions.

‣ cudaEnablePerThreadDefaultStream: This will enable the search for all driver symbols that match

the requested driver symbol name including the per-thread versions. If a per-thread version is not

found, the API will return the legacy version of the driver function.

Note:

‣ Version mixing among CUDA-defined types and driver API versions is strongly discouraged and

doing so can result in an undefined behavior. More here.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cuGetProcAddress

__host__cudaError_t cudaGetDriverEntryPointByVersion

(const char *symbol, void **funcPtr, unsigned

int cudaVersion, unsigned long long flags,

cudaDriverEntryPointQueryResult *driverStatus)

Returns the requested driver API function pointer by CUDA version.

Parameters

symbol

- The base name of the driver API function to look for. As an example, for the driver API

cuMemAlloc_v2, symbol would be cuMemAlloc.

funcPtr

- Location to return the function pointer to the requested driver function

cudaVersion

- The CUDA version to look for the requested driver symbol

flags

- Flags to specify search options.

driverStatus

- Optional location to store the status of finding the symbol from the driver. See

cudaDriverEntryPointQueryResult for possible values.

405

Modules

Returns

cudaSuccess, cudaErrorInvalidValue, cudaErrorNotSupported

Description

Returns in **funcPtr the address of the CUDA driver function for the requested flags and CUDA

driver version.

The CUDA version is specified as (1000 * major + 10 * minor), so CUDA 11.2 should be specified

as 11020. For a requested driver symbol, if the specified CUDA version is greater than or equal to the

CUDA version in which the driver symbol was introduced, this API will return the function pointer to

the corresponding versioned function.

The pointer returned by the API should be cast to a function pointer matching the requested driver

function's definition in the API header file. The function pointer typedef can be picked up from the

corresponding typedefs header file. For example, cudaTypedefs.h consists of function pointer typedefs

for driver APIs defined in cuda.h.

For the case where the CUDA version requested is greater than the CUDA Toolkit installed, there may

not be an appropriate function pointer typedef in the corresponding header file and may need a custom

typedef to match the driver function signature returned. This can be done by getting the typedefs from a

later toolkit or creating appropriately matching custom function typedefs.

The API will return cudaSuccess and set the returned funcPtr if the requested driver function is valid

and supported on the platform.

The API will return cudaSuccess and set the returned funcPtr to NULL if the requested driver

function is not supported on the platform, no ABI compatible driver function exists for the requested

version or if the driver symbol is invalid.

It will also set the optional driverStatus to one of the values in cudaDriverEntryPointQueryResult

with the following meanings:

‣ cudaDriverEntryPointSuccess - The requested symbol was succesfully found based on input

arguments and pfn is valid

‣ cudaDriverEntryPointSymbolNotFound - The requested symbol was not found

‣ cudaDriverEntryPointVersionNotSufficent - The requested symbol was found but is not supported

by the specified version cudaVersion

The requested flags can be:

‣ cudaEnableDefault: This is the default mode. This is equivalent to

cudaEnablePerThreadDefaultStream if the code is compiled with --default-stream per-thread

compilation flag or the macro CUDA_API_PER_THREAD_DEFAULT_STREAM is defined;

cudaEnableLegacyStream otherwise.

‣ cudaEnableLegacyStream: This will enable the search for all driver symbols that match the

requested driver symbol name except the corresponding per-thread versions.

406

Modules

‣ cudaEnablePerThreadDefaultStream: This will enable the search for all driver symbols that match

the requested driver symbol name including the per-thread versions. If a per-thread version is not

found, the API will return the legacy version of the driver function.

Note:

‣ Version mixing among CUDA-defined types and driver API versions is strongly discouraged and

doing so can result in an undefined behavior. More here.

‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or

cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.

‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.

cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.

See also:

cuGetProcAddress

6.32. Library Management

This section describes the library management functions of the CUDA runtime application

programming interface.

__host__cudaError_t cudaKernelSetAttributeForDevice

(cudaKernel_t kernel, cudaFuncAttribute attr, int value, int

device)

Sets information about a kernel.

Parameters

kernel

- Kernel to set attribute of

attr

- Attribute requested

value

- Value to set

device

- Device to set attribute of

Returns

cudaSuccess, cudaErrorInvalidDeviceFunction, cudaErrorInvalidValue

407

Modules

Description

This call sets the value of a specified attribute attr on the kernel kernel for the requested device

device to an integer value specified by value. This function returns cudaSuccess if the new

value of the attribute could be successfully set. If the set fails, this call will return an error. Not all

attributes can have values set. Attempting to set a value on a read-only attribute will result in an error

(cudaErrorInvalidValue)

Note that attributes set using cudaFuncSetAttribute() will override the attribute set by this API

irrespective of whether the call to cudaFuncSetAttribute() is made before or after this API call. Because

of this and the stricter locking requirements mentioned below it is suggested that this call be used

during the initialization path and not on each thread accessing kernel such as on kernel launches or

on the critical path.

Valid values for attr are:

‣ cudaFuncAttributeMaxDynamicSharedMemorySize - The requested maximum

size in bytes of dynamically-allocated shared memory. The sum of this value

and the function attribute sharedSizeBytes cannot exceed the device attribute

cudaDevAttrMaxSharedMemoryPerBlockOptin. The maximal size of requestable dynamic shared

memory may differ by GPU architecture.

‣ cudaFuncAttributePreferredSharedMemoryCarveout - On devices where the L1 cache and shared

memory use the same hardware resources, this sets the shared memory carveout preference, in

percent of the total shared memory. See cudaDevAttrMaxSharedMemoryPerMultiprocessor. This

is only a hint, and the driver can choose a different ratio if required to execute the function.

‣ cudaFuncAttributeRequiredClusterWidth: The required cluster width in blocks. The width, height,

and depth values must either all be 0 or all be positive. The validity of the cluster dimensions is

checked at launch time. If the value is set during compile time, it cannot be set at runtime. Setting it

at runtime will return cudaErrorNotPermitted.

‣ cudaFuncAttributeRequiredClusterHeight: The required cluster height in blocks. The width, height,

and depth values must either all be 0 or all be positive. The validity of the cluster dimensions is

checked at launch time. If the value is set during compile time, it cannot be set at runtime. Setting it

at runtime will return cudaErrorNotPermitted.

‣ cudaFuncAttributeRequiredClusterDepth: The required cluster depth in blocks. The width, height,

and depth values must either all be 0 or all be positive. The validity of the cluster dimensions is

checked at launch time. If the value is set during compile time, it cannot be set at runtime. Setting it

at runtime will return cudaErrorNotPermitted.

‣ cudaFuncAttributeNonPortableClusterSizeAllowed: Indicates whether the function can be

launched with non-portable cluster size. 1 is allowed, 0 is disallowed.

‣ cudaFuncAttributeClusterSchedulingPolicyPreference: The block scheduling policy of a function.

The value type is cudaClusterSchedulingPolicy.

Note:

408

Modules

The API has stricter locking requirements in comparison to its legacy counterpart

cudaFuncSetAttribute() due to device-wide semantics. If multiple threads are trying to set the same

attribute on the same device simultaneously, the attribute setting will depend on the interleavings chosen

by the OS scheduler and memory consistency.

See also:

cudaLibraryLoadData, cudaLibraryLoadFromFile, cudaLibraryUnload, cudaLibraryGetKernel,

cudaLaunchKernel, cudaFuncSetAttribute, cuKernelSetAttribute

__host__cudaError_t cudaLibraryEnumerateKernels

(cudaKernel_t *kernels, unsigned int numKernels,

cudaLibrary_t lib)

Retrieve the kernel handles within a library.

Parameters

kernels

- Buffer where the kernel handles are returned to

numKernels

- Maximum number of kernel handles may be returned to the buffer

lib

- Library to query from

Returns

cudaSuccess, cudaErrorCudartUnloading, cudaErrorInitializationError, cudaErrorInvalidValue,

cudaErrorInvalidResourceHandle

Description

Returns in kernels a maximum number of numKernels kernel handles within lib. The returned

kernel handle becomes invalid when the library is unloaded.

See also:

cudaLibraryGetKernelCount, cuLibraryEnumerateKernels

409

Modules

__host__cudaError_t cudaLibraryGetGlobal (void **dptr,

size_t *bytes, cudaLibrary_t library, const char *name)

Returns a global device pointer.

Parameters

dptr

- Returned global device pointer for the requested library

bytes

- Returned global size in bytes

library

- Library to retrieve global from

name

- Name of global to retrieve

Returns

cudaSuccess, cudaErrorCudartUnloading, cudaErrorInitializationError, cudaErrorInvalidValue,

cudaErrorInvalidResourceHandle, cudaErrorSymbolNotFound cudaErrorDeviceUninitialized,

cudaErrorContextIsDestroyed

Description

Returns in *dptr and *bytes the base pointer and size of the global with name name for the

requested library library and the current device. If no global for the requested name name exists,

the call returns cudaErrorSymbolNotFound. One of the parameters dptr or bytes (not both) can be

NULL in which case it is ignored. The returned dptr cannot be passed to the Symbol APIs such as

cudaMemcpyToSymbol, cudaMemcpyFromSymbol, cudaGetSymbolAddress, or cudaGetSymbolSize.

See also:

cudaLibraryLoadData, cudaLibraryLoadFromFile, cudaLibraryUnload, cudaLibraryGetManaged,

cuLibraryGetGlobal

__host__cudaError_t cudaLibraryGetKernel (cudaKernel_t

*pKernel, cudaLibrary_t library, const char *name)

Returns a kernel handle.

Parameters

pKernel

- Returned kernel handle

library

- Library to retrieve kernel from

410

Modules

name

- Name of kernel to retrieve

Returns

cudaSuccess, cudaErrorCudartUnloading, cudaErrorInitializationError, cudaErrorInvalidValue,

cudaErrorInvalidResourceHandle, cudaErrorSymbolNotFound

Description

Returns in pKernel the handle of the kernel with name name located in library library. If kernel

handle is not found, the call returns cudaErrorSymbolNotFound.

See also:

cudaLibraryLoadData, cudaLibraryLoadFromFile, cudaLibraryUnload, cuLibraryGetKernel

__host__cudaError_t cudaLibraryGetKernelCount

(unsigned int *count, cudaLibrary_t lib)

Returns the number of kernels within a library.

Parameters

count

- Number of kernels found within the library

lib

- Library to query

Returns

cudaSuccess, cudaErrorCudartUnloading, cudaErrorInitializationError, cudaErrorInvalidValue,

cudaErrorInvalidResourceHandle

Description

Returns in count the number of kernels in lib.

See also:

cudaLibraryEnumerateKernels, cudaLibraryLoadFromFile, cudaLibraryLoadData,

cuLibraryGetKernelCount

411

Modules

__host__cudaError_t cudaLibraryGetManaged (void

**dptr, size_t *bytes, cudaLibrary_t library, const char

*name)

Returns a pointer to managed memory.

Parameters

dptr

- Returned pointer to the managed memory

bytes

- Returned memory size in bytes

library

- Library to retrieve managed memory from

name

- Name of managed memory to retrieve

Returns

cudaSuccess, cudaErrorCudartUnloading, cudaErrorInitializationError, cudaErrorInvalidValue,

cudaErrorInvalidResourceHandle, cudaErrorSymbolNotFound

Description

Returns in *dptr and *bytes the base pointer and size of the managed memory with name name

for the requested library library. If no managed memory with the requested name name exists,

the call returns cudaErrorSymbolNotFound. One of the parameters dptr or bytes (not both) can

be NULL in which case it is ignored. Note that managed memory for library library is shared

across devices and is registered when the library is loaded. The returned dptr cannot be passed to the

Symbol APIs such as cudaMemcpyToSymbol, cudaMemcpyFromSymbol, cudaGetSymbolAddress, or

cudaGetSymbolSize.

See also:

cudaLibraryLoadData, cudaLibraryLoadFromFile, cudaLibraryUnload, cudaLibraryGetGlobal,

cuLibraryGetManaged

412

Modules

__host__cudaError_t cudaLibraryGetUnifiedFunction (void

**fptr, cudaLibrary_t library, const char *symbol)

Returns a pointer to a unified function.

Parameters

fptr

- Returned pointer to a unified function

library

- Library to retrieve function pointer memory from

symbol

- Name of function pointer to retrieve

Returns

cudaSuccess, cudaErrorCudartUnloading, cudaErrorInitializationError, cudaErrorInvalidValue,

cudaErrorInvalidResourceHandle, cudaErrorSymbolNotFound

Description

Returns in *fptr the function pointer to a unified function denoted by symbol. If no unified

function with name symbol exists, the call returns cudaErrorSymbolNotFound. If there is no device

with attribute cudaDeviceProp::unifiedFunctionPointers present in the system, the call may return

cudaErrorSymbolNotFound.

See also:

cudaLibraryLoadData, cudaLibraryLoadFromFile, cudaLibraryUnload, cuLibraryGetUnifiedFunction

__host__cudaError_t cudaLibraryLoadData

(cudaLibrary_t *library, const void *code, cudaJitOption

*jitOptions, void **jitOptionsValues, unsigned int

numJitOptions, cudaLibraryOption *libraryOptions, void

**libraryOptionValues, unsigned int numLibraryOptions)

Load a library with specified code and options.

Parameters

library

- Returned library

code

- Code to load

413

Modules

jitOptions

- Options for JIT

jitOptionsValues

- Option values for JIT

numJitOptions

- Number of options

libraryOptions

- Options for loading

libraryOptionValues

- Option values for loading

numLibraryOptions

- Number of options for loading

Returns

cudaSuccess, cudaErrorInvalidValue, cudaErrorMemoryAllocation, cudaErrorInitializationError,

cudaErrorCudartUnloading, cudaErrorInvalidPtx, cudaErrorUnsupportedPtxVersion,

cudaErrorNoKernelImageForDevice, cudaErrorSharedObjectSymbolNotFound,

cudaErrorSharedObjectInitFailed, cudaErrorJitCompilerNotFound

Description

Takes a pointer code and loads the corresponding library library based on the application defined

library loading mode:

‣ If module loading is set to EAGER, via the environment variables described in "Module loading",

library is loaded eagerly into all contexts at the time of the call and future contexts at the time

of creation until the library is unloaded with cudaLibraryUnload().

‣ If the environment variables are set to LAZY, library is not immediately loaded onto all

existent contexts and will only be loaded when a function is needed for that context, such as a

kernel launch.

These environment variables are described in the CUDA programming guide under the "CUDA

environment variables" section.

The code may be a cubin or fatbin as output by nvcc, or a NULL-terminated PTX, either as output by

nvcc or hand-written. A fatbin should also contain relocatable code when doing separate compilation.

Please also see the documentation for nvrtc (https://docs.nvidia.com/cuda/nvrtc/index.html ), nvjitlink

(https://docs.nvidia.com/cuda/nvjitlink/index.html), and nvfatbin (https://docs.nvidia.com/cuda/

nvfatbin/index.html) for more information on generating loadable code at runtime.

Options are passed as an array via jitOptions and any corresponding parameters are passed in

jitOptionsValues. The number of total JIT options is supplied via numJitOptions. Any

outputs will be returned via jitOptionsValues.

414

Modules

Library load options are passed as an array via libraryOptions and any corresponding parameters

are passed in libraryOptionValues. The number of total library load options is supplied via

numLibraryOptions.

See also:

cudaLibraryLoadFromFile, cudaLibraryUnload, cuLibraryLoadData

__host__cudaError_t cudaLibraryLoadFromFile

(cudaLibrary_t *library, const char *fileName,

cudaJitOption *jitOptions, void **jitOptionsValues,

unsigned int numJitOptions, cudaLibraryOption

*libraryOptions, void **libraryOptionValues, unsigned int

numLibraryOptions)

Load a library with specified file and options.

Parameters

library

- Returned library

fileName

- File to load from

jitOptions

- Options for JIT

jitOptionsValues

- Option values for JIT

numJitOptions

- Number of options

libraryOptions

- Options for loading

libraryOptionValues

- Option values for loading

numLibraryOptions

- Number of options for loading

Returns

cudaSuccess, cudaErrorInvalidValue, cudaErrorMemoryAllocation, cudaErrorInitializationError,

cudaErrorCudartUnloading, cudaErrorInvalidPtx, cudaErrorUnsupportedPtxVersion,

cudaErrorNoKernelImageForDevice, cudaErrorSharedObjectSymbolNotFound,

cudaErrorSharedObjectInitFailed, cudaErrorJitCompilerNotFound

415

Modules

Description

Takes a pointer code and loads the corresponding library library based on the application defined

library loading mode:

‣ If module loading is set to EAGER, via the environment variables described in "Module loading",

library is loaded eagerly into all contexts at the time of the call and future contexts at the time

of creation until the library is unloaded with cudaLibraryUnload().

‣ If the environment variables are set to LAZY, library is not immediately loaded onto all

existent contexts and will only be loaded when a function is needed for that context, such as a

kernel launch.

These environment variables are described in the CUDA programming guide under the "CUDA

environment variables" section.

The file should be a cubin file as output by nvcc, or a PTX file either as output by nvcc or handwritten,

or a fatbin file as output by nvcc. A fatbin should also contain relocatable code when doing separate

compilation. Please also see the documentation for nvrtc (https://docs.nvidia.com/cuda/nvrtc/

index.html), nvjitlink (https://docs.nvidia.com/cuda/nvjitlink/index.html), and nvfatbin (https://

docs.nvidia.com/cuda/nvfatbin/index.html) for more information on generating loadable code at

runtime.

Options are passed as an array via jitOptions and any corresponding parameters are passed in

jitOptionsValues. The number of total options is supplied via numJitOptions. Any outputs

will be returned via jitOptionsValues.

Library load options are passed as an array via libraryOptions and any corresponding parameters

are passed in libraryOptionValues. The number of total library load options is supplied via

numLibraryOptions.

See also:

cudaLibraryLoadData, cudaLibraryUnload, cuLibraryLoadFromFile

__host__cudaError_t cudaLibraryUnload (cudaLibrary_t

library)

Unloads a library.

Parameters

library

- Library to unload

Returns

cudaSuccess, cudaErrorCudartUnloading, cudaErrorInitializationError, cudaErrorInvalidValue

416

Modules

Description

Unloads the library specified with library

See also:

cudaLibraryLoadData, cudaLibraryLoadFromFile, cuLibraryUnload

6.33. C++ API Routines

C++-style interface built on top of CUDA runtime API.

This section describes the C++ high level API functions of the CUDA runtime application

programming interface. To use these functions, your application needs to be compiled with the nvcc

compiler.

__cudaOccupancyB2DHelper

cppClassifierVisibility: visibility=public

template < class T > __host__cudaCreateChannelDesc

(void)

[C++ API] Returns a channel descriptor using the specified format

Returns

Channel descriptor with format f

Description

Returns a channel descriptor with format f and number of bits of each component x, y, z, and w. The

cudaChannelFormatDesc is defined as:

‎ struct cudaChannelFormatDesc {

int x, y, z, w;

enum cudaChannelFormatKind

};

where cudaChannelFormatKind is one of cudaChannelFormatKindSigned,

cudaChannelFormatKindUnsigned, cudaChannelFormatKindFloat,

cudaChannelFormatKindSignedNormalized8X1, cudaChannelFormatKindSignedNormalized8X2,

cudaChannelFormatKindSignedNormalized8X4, cudaChannelFormatKindUnsignedNormalized8X1,

cudaChannelFormatKindUnsignedNormalized8X2, cudaChannelFormatKindUnsignedNormalized8X4,

cudaChannelFormatKindSignedNormalized16X1, cudaChannelFormatKindSignedNormalized16X2,

cudaChannelFormatKindSignedNormalized16X4, cudaChannelFormatKindUnsignedNormalized16X1,

cudaChannelFormatKindUnsignedNormalized16X2,

417

CUDA Runtime API. API Reference Manual (vRelease Version, January 2024) - page 6