CUDA Driver API. API Reference Manual (January 2024) - page 6

 

  Index      Manuals     CUDA Driver API. API Reference Manual (January 2024)

 

Search            copyright infringement  

 

   

 

    

 

Content      ..     4      5      6      7     ..

 

 

 

CUDA Driver API. API Reference Manual (January 2024) - page 6

 

 

Modules
CUresult cuStreamBeginCapture (CUstream hStream,
CUstreamCaptureMode mode)
Begins graph capture on a stream.
Parameters
hStream
- Stream in which to initiate capture
mode
- Controls the interaction of this capture sequence with other API calls that are potentially unsafe.
For more details see cuThreadExchangeStreamCaptureMode.
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_VALUE
Description
Begin graph capture on hStream. When a stream is in capture mode, all operations pushed into
the stream will not be executed, but will instead be captured into a graph, which will be returned via
cuStreamEndCapture. Capture may not be initiated if stream is CU_STREAM_LEGACY. Capture
must be ended on the same stream in which it was initiated, and it may only be initiated if the stream is
not already in capture mode. The capture mode may be queried via cuStreamIsCapturing. A unique id
representing the capture sequence may be queried via cuStreamGetCaptureInfo.
If mode is not CU_STREAM_CAPTURE_MODE_RELAXED, cuStreamEndCapture must be called
on this stream from the same thread.
Note:
Kernels captured using this API must not use texture and surface references. Reading or writing through
any texture or surface reference is undefined behavior. This restriction does not apply to texture and
surface objects.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuStreamCreate, cuStreamIsCapturing, cuStreamEndCapture, cuThreadExchangeStreamCaptureMode
|
337
Modules
CUresult cuStreamBeginCaptureToGraph (CUstream
hStream, CUgraph hGraph, const CUgraphNode
*dependencies, const CUgraphEdgeData *dependencyData,
size_t numDependencies, CUstreamCaptureMode mode)
Begins graph capture on a stream to an existing graph.
Parameters
hStream
- Stream in which to initiate capture.
hGraph
- Graph to capture into.
dependencies
- Dependencies of the first node captured in the stream. Can be NULL if numDependencies is 0.
dependencyData
- Optional array of data associated with each dependency.
numDependencies
- Number of dependencies.
mode
- Controls the interaction of this capture sequence with other API calls that are potentially unsafe.
For more details see cuThreadExchangeStreamCaptureMode.
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_VALUE
Description
Begin graph capture on hStream, placing new nodes into an existing graph. When a stream is in
capture mode, all operations pushed into the stream will not be executed, but will instead be captured
into hGraph. The graph will not be instantiable until the user calls cuStreamEndCapture.
Capture may not be initiated if stream is CU_STREAM_LEGACY. Capture must be ended on
the same stream in which it was initiated, and it may only be initiated if the stream is not already in
capture mode. The capture mode may be queried via cuStreamIsCapturing. A unique id representing
the capture sequence may be queried via cuStreamGetCaptureInfo.
If mode is not CU_STREAM_CAPTURE_MODE_RELAXED, cuStreamEndCapture must be called
on this stream from the same thread.
Note:
|
338
Modules
Kernels captured using this API must not use texture and surface references. Reading or writing through
any texture or surface reference is undefined behavior. This restriction does not apply to texture and
surface objects.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuStreamBeginCapture, cuStreamCreate, cuStreamIsCapturing, cuStreamEndCapture,
cuThreadExchangeStreamCaptureMode, cuGraphAddNode
CUresult cuStreamCopyAttributes (CUstream dst,
CUstream src)
Copies attributes from source stream to destination stream.
Parameters
dst
Destination stream
src
Source stream For list of attributes see CUstreamAttrID
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE
Description
Copies attributes from source stream src to destination stream dst. Both streams must have the same
context.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
CUaccessPolicyWindow
|
339
Modules
CUresult cuStreamCreate (CUstream *phStream, unsigned
int Flags)
Create a stream.
Parameters
phStream
- Returned newly created stream
Flags
- Parameters for stream creation
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_OUT_OF_MEMORY
Description
Creates a stream and returns a handle in phStream. The Flags argument determines behaviors of
the stream.
Valid values for Flags are:
‣ CU_STREAM_DEFAULT: Default stream creation flag.
‣ CU_STREAM_NON_BLOCKING: Specifies that work running in the created stream may run
concurrently with work in stream 0 (the NULL stream), and that the created stream should perform
no implicit synchronization with stream 0.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuStreamDestroy, cuStreamCreateWithPriority, cuGreenCtxStreamCreate, cuStreamGetPriority,
cuStreamGetFlags, cuStreamGetDevice cuStreamWaitEvent, cuStreamQuery, cuStreamSynchronize,
cuStreamAddCallback, cudaStreamCreate, cudaStreamCreateWithFlags
|
340
Modules
CUresult cuStreamCreateWithPriority (CUstream
*phStream, unsigned int flags, int priority)
Create a stream with the given priority.
Parameters
phStream
- Returned newly created stream
flags
- Flags for stream creation. See cuStreamCreate for a list of valid flags
priority
- Stream priority. Lower numbers represent higher priorities. See cuCtxGetStreamPriorityRange for
more information about meaningful stream priorities that can be passed.
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_OUT_OF_MEMORY
Description
Creates a stream with the specified priority and returns a handle in phStream. This affects the
scheduling priority of work in the stream. Priorities provide a hint to preferentially run work with
higher priority when possible, but do not preempt already-running work or provide any other functional
guarantee on execution order.
priority follows a convention where lower numbers represent higher priorities. '0'
represents default priority. The range of meaningful numerical priorities can be queried using
cuCtxGetStreamPriorityRange. If the specified priority is outside the numerical range returned by
cuCtxGetStreamPriorityRange, it will automatically be clamped to the lowest or the highest number in
the range.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Stream priorities are supported only on GPUs with compute capability 3.5 or higher.
‣ In the current implementation, only compute kernels launched in priority streams are affected by
the stream's priority. Stream priorities have no effect on host-to-device and device-to-host memory
operations.
See also:
|
341
Modules
cuStreamDestroy, cuStreamCreate, cuGreenCtxStreamCreate, cuStreamGetPriority,
cuCtxGetStreamPriorityRange, cuStreamGetFlags, cuStreamGetDevice, cuStreamWaitEvent,
cuStreamQuery, cuStreamSynchronize, cuStreamAddCallback, cudaStreamCreateWithPriority
CUresult cuStreamDestroy (CUstream hStream)
Destroys a stream.
Parameters
hStream
- Stream to destroy
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_INVALID_HANDLE
Description
Destroys the stream specified by hStream.
In case the device is still doing work in the stream hStream when cuStreamDestroy() is called,
the function will return immediately and the resources associated with hStream will be released
automatically once the device has completed all work in hStream.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuStreamCreate, cuStreamWaitEvent, cuStreamQuery, cuStreamSynchronize, cuStreamAddCallback,
cudaStreamDestroy
CUresult cuStreamEndCapture (CUstream hStream,
CUgraph *phGraph)
Ends capture on a stream, returning the captured graph.
Parameters
hStream
- Stream to query
phGraph
- The captured graph
|
342
Modules
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_STREAM_CAPTURE_WRONG_THREAD
Description
End capture on hStream, returning the captured graph via phGraph. Capture must have been
initiated on hStream via a call to cuStreamBeginCapture. If capture was invalidated, due to a
violation of the rules of stream capture, then a NULL graph will be returned.
If the mode argument to cuStreamBeginCapture was not
CU_STREAM_CAPTURE_MODE_RELAXED, this call must be from the same thread as
cuStreamBeginCapture.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuStreamCreate, cuStreamBeginCapture, cuStreamIsCapturing, cuGraphDestroy
CUresult cuStreamGetAttribute (CUstream hStream,
CUstreamAttrID attr, CUstreamAttrValue *value_out)
Queries stream attribute.
Parameters
hStream
attr
value_out
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_INVALID_HANDLE
Description
Queries attribute attr from hStream and stores it in corresponding member of value_out.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
|
343
Modules
See also:
CUaccessPolicyWindow
CUresult cuStreamGetCaptureInfo (CUstream hStream,
CUstreamCaptureStatus *captureStatus_out, cuuint64_t
*id_out, CUgraph *graph_out, const CUgraphNode
**dependencies_out, size_t *numDependencies_out)
Query a stream's capture state.
Parameters
hStream
- The stream to query
captureStatus_out
- Location to return the capture status of the stream; required
id_out
- Optional location to return an id for the capture sequence, which is unique over the lifetime of the
process
graph_out
- Optional location to return the graph being captured into. All operations other than destroy and
node removal are permitted on the graph while the capture sequence is in progress. This API does
not transfer ownership of the graph, which is transferred or destroyed at cuStreamEndCapture. Note
that the graph handle may be invalidated before end of capture for certain errors. Nodes that are or
become unreachable from the original stream at cuStreamEndCapture due to direct actions on the
graph do not trigger CUDA_ERROR_STREAM_CAPTURE_UNJOINED.
dependencies_out
- Optional location to store a pointer to an array of nodes. The next node to be captured in the
stream will depend on this set of nodes, absent operations such as event wait which modify this set.
The array pointer is valid until the next API call which operates on the stream or until the capture is
terminated. The node handles may be copied out and are valid until they or the graph is destroyed.
The driver-owned array may also be passed directly to APIs that operate on the graph (not the
stream) without copying.
numDependencies_out
- Optional location to store the size of the array returned in dependencies_out.
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_STREAM_CAPTURE_IMPLICIT
Description
Query stream state related to stream capture.
|
344
Modules
If called on CU_STREAM_LEGACY (the "null stream") while a stream
not created with CU_STREAM_NON_BLOCKING is capturing, returns
CUDA_ERROR_STREAM_CAPTURE_IMPLICIT.
Valid data (other than capture status) is returned only if both of the following are true:
‣ the call returns CUDA_SUCCESS
‣ the returned capture status is CU_STREAM_CAPTURE_STATUS_ACTIVE
Note:
‣ Graph objects are not threadsafe. More here.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuStreamGetCaptureInfo_v3 cuStreamBeginCapture, cuStreamIsCapturing,
cuStreamUpdateCaptureDependencies
CUresult cuStreamGetCaptureInfo_v3 (CUstream hStream,
CUstreamCaptureStatus *captureStatus_out, cuuint64_t
*id_out, CUgraph *graph_out, const CUgraphNode
**dependencies_out, const CUgraphEdgeData
**edgeData_out, size_t *numDependencies_out)
Query a stream's capture state (12.3+).
Parameters
hStream
- The stream to query
captureStatus_out
- Location to return the capture status of the stream; required
id_out
- Optional location to return an id for the capture sequence, which is unique over the lifetime of the
process
graph_out
- Optional location to return the graph being captured into. All operations other than destroy and
node removal are permitted on the graph while the capture sequence is in progress. This API does
not transfer ownership of the graph, which is transferred or destroyed at cuStreamEndCapture. Note
that the graph handle may be invalidated before end of capture for certain errors. Nodes that are or
become unreachable from the original stream at cuStreamEndCapture due to direct actions on the
graph do not trigger CUDA_ERROR_STREAM_CAPTURE_UNJOINED.
|
345
Modules
dependencies_out
- Optional location to store a pointer to an array of nodes. The next node to be captured in the
stream will depend on this set of nodes, absent operations such as event wait which modify this set.
The array pointer is valid until the next API call which operates on the stream or until the capture is
terminated. The node handles may be copied out and are valid until they or the graph is destroyed.
The driver-owned array may also be passed directly to APIs that operate on the graph (not the
stream) without copying.
edgeData_out
- Optional location to store a pointer to an array of graph edge data. This array parallels
dependencies_out; the next node to be added has an edge to dependencies_out[i] with
annotation edgeData_out[i] for each i. The array pointer is valid until the next API call which
operates on the stream or until the capture is terminated.
numDependencies_out
- Optional location to store the size of the array returned in dependencies_out.
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_STREAM_CAPTURE_IMPLICIT, CUDA_ERROR_LOSSY_QUERY
Description
Query stream state related to stream capture.
If called on CU_STREAM_LEGACY (the "null stream") while a stream
not created with CU_STREAM_NON_BLOCKING is capturing, returns
CUDA_ERROR_STREAM_CAPTURE_IMPLICIT.
Valid data (other than capture status) is returned only if both of the following are true:
‣ the call returns CUDA_SUCCESS
‣ the returned capture status is CU_STREAM_CAPTURE_STATUS_ACTIVE
If edgeData_out is non-NULL then dependencies_out must be as well. If
dependencies_out is non-NULL and edgeData_out is NULL, but there is non-
zero edge data for one or more of the current stream dependencies, the call will return
CUDA_ERROR_LOSSY_QUERY.
Note:
‣ Graph objects are not threadsafe. More here.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuStreamGetCaptureInfo, cuStreamBeginCapture, cuStreamIsCapturing,
cuStreamUpdateCaptureDependencies
|
346
Modules
CUresult cuStreamGetCtx (CUstream hStream, CUcontext
*pctx)
Query the context associated with a stream.
Parameters
hStream
- Handle to the stream to be queried
pctx
- Returned context associated with the stream
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE,
CUDA_ERROR_NOT_SUPPORTED
Description
Returns the CUDA context that the stream is associated with.
Note there is a later version of this API, cuStreamGetCtx_v2. It will supplant this version in
CUDA 13.0. It is recommended to use cuStreamGetCtx_v2 till then as this version will return
CUDA_ERROR_NOT_SUPPORTED for streams created via the API cuGreenCtxStreamCreate.
The stream handle hStream can refer to any of the following:
‣ a stream created via any of the CUDA driver APIs such as cuStreamCreate and
cuStreamCreateWithPriority, or their runtime API equivalents such as cudaStreamCreate,
cudaStreamCreateWithFlags and cudaStreamCreateWithPriority. The returned context is the
context that was active in the calling thread when the stream was created. Passing an invalid handle
will result in undefined behavior.
‣ any of the special streams such as the NULL stream, CU_STREAM_LEGACY and
CU_STREAM_PER_THREAD. The runtime API equivalents of these are also accepted, which are
NULL, cudaStreamLegacy and cudaStreamPerThread respectively. Specifying any of the special
handles will return the context current to the calling thread. If no context is current to the calling
thread, CUDA_ERROR_INVALID_CONTEXT is returned.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
|
347
Modules
cuStreamDestroy, cuStreamCreateWithPriority, cuStreamGetPriority, cuStreamGetFlags,
cuStreamGetDevice cuStreamWaitEvent, cuStreamQuery, cuStreamSynchronize,
cuStreamAddCallback, cudaStreamCreate, cuStreamGetCtx_v2, cudaStreamCreateWithFlags
CUresult cuStreamGetCtx_v2 (CUstream hStream,
CUcontext *pCtx, CUgreenCtx *pGreenCtx)
Query the contexts associated with a stream.
Parameters
hStream
- Handle to the stream to be queried
pCtx
- Returned regular context associated with the stream
pGreenCtx
- Returned green context if the stream is associated with a green context or NULL if not
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE
Description
Returns the contexts that the stream is associated with.
If the stream is associated with a green context, the API returns the green context in pGreenCtx and
the primary context of the associated device in pCtx.
If the stream is associated with a regular context, the API returns the regular context in pCtx and
NULL in pGreenCtx.
The stream handle hStream can refer to any of the following:
‣ a stream created via any of the CUDA driver APIs such as cuStreamCreate,
cuStreamCreateWithPriority and cuGreenCtxStreamCreate, or their runtime API equivalents such
as cudaStreamCreate, cudaStreamCreateWithFlags and cudaStreamCreateWithPriority. Passing an
invalid handle will result in undefined behavior.
‣ any of the special streams such as the NULL stream, CU_STREAM_LEGACY and
CU_STREAM_PER_THREAD. The runtime API equivalents of these are also accepted, which
are NULL, cudaStreamLegacy and cudaStreamPerThread respectively. If any of the special
handles are specified, the API will operate on the context current to the calling thread. If a green
context (that was converted via cuCtxFromGreenCtx() before setting it current) is current to the
calling thread, the API will return the green context in pGreenCtx and the primary context of
the associated device in pCtx. If a regular context is current, the API returns the regular context
in pCtx and NULL in pGreenCtx. Note that specifying CU_STREAM_PER_THREAD
|
348
Modules
or cudaStreamPerThread will return CUDA_ERROR_INVALID_HANDLE if a green
context is current to the calling thread. If no context is current to the calling thread,
CUDA_ERROR_INVALID_CONTEXT is returned.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuStreamDestroy, cuStreamCreate cuStreamCreateWithPriority, cuGreenCtxStreamCreate,
cuStreamGetPriority, cuStreamGetFlags, cuStreamGetDevice, cuStreamWaitEvent, cuStreamQuery,
cuStreamSynchronize, cuStreamAddCallback, cudaStreamCreate, cudaStreamCreateWithFlags,
CUresult cuStreamGetDevice (CUstream hStream,
CUdevice *device)
Returns the device handle of the stream.
Parameters
hStream
- Handle to the stream to be queried
device
- Returns the device to which a stream belongs
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_INVALID_HANDLE, CUDA_ERROR_OUT_OF_MEMORY
Description
Returns in *device the device handle of the stream
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuStreamDestroy, cuStreamCreate, cuGreenCtxStreamCreate, cuStreamGetFlags
|
349
Modules
CUresult cuStreamGetFlags (CUstream hStream, unsigned
int *flags)
Query the flags of a given stream.
Parameters
hStream
- Handle to the stream to be queried
flags
- Pointer to an unsigned integer in which the stream's flags are returned The value returned in
flags is a logical 'OR' of all flags that were used while creating this stream. See cuStreamCreate
for the list of valid flags
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_INVALID_HANDLE, CUDA_ERROR_OUT_OF_MEMORY
Description
Query the flags of a stream created using cuStreamCreate, cuStreamCreateWithPriority or
cuGreenCtxStreamCreate and return the flags in flags.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuStreamDestroy, cuStreamCreate, cuGreenCtxStreamCreate, cuStreamGetPriority,
cudaStreamGetFlags, cuStreamGetDevice
CUresult cuStreamGetId (CUstream hStream, unsigned
long long *streamId)
Returns the unique Id associated with the stream handle supplied.
Parameters
hStream
- Handle to the stream to be queried
streamId
- Pointer to store the Id of the stream
|
350
Modules
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_INVALID_HANDLE
Description
Returns in streamId the unique Id which is associated with the given stream handle. The Id is unique
for the life of the program.
The stream handle hStream can refer to any of the following:
‣ a stream created via any of the CUDA driver APIs such as cuStreamCreate and
cuStreamCreateWithPriority, or their runtime API equivalents such as cudaStreamCreate,
cudaStreamCreateWithFlags and cudaStreamCreateWithPriority. Passing an invalid handle will
result in undefined behavior.
‣ any of the special streams such as the NULL stream, CU_STREAM_LEGACY and
CU_STREAM_PER_THREAD. The runtime API equivalents of these are also accepted, which are
NULL, cudaStreamLegacy and cudaStreamPerThread respectively.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuStreamDestroy, cuStreamCreate, cuStreamGetPriority, cudaStreamGetId
CUresult cuStreamGetPriority (CUstream hStream, int
*priority)
Query the priority of a given stream.
Parameters
hStream
- Handle to the stream to be queried
priority
- Pointer to a signed integer in which the stream's priority is returned
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_INVALID_HANDLE, CUDA_ERROR_OUT_OF_MEMORY
|
351
Modules
Description
Query the priority of a stream created using cuStreamCreate, cuStreamCreateWithPriority or
cuGreenCtxStreamCreate and return the priority in priority. Note that if the stream was created
with a priority outside the numerical range returned by cuCtxGetStreamPriorityRange, this function
returns the clamped priority. See cuStreamCreateWithPriority for details about priority clamping.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuStreamDestroy, cuStreamCreate, cuStreamCreateWithPriority, cuGreenCtxStreamCreate,
cuCtxGetStreamPriorityRange, cuStreamGetFlags, cuStreamGetDevice, cudaStreamGetPriority
CUresult cuStreamIsCapturing (CUstream hStream,
CUstreamCaptureStatus *captureStatus)
Returns a stream's capture status.
Parameters
hStream
- Stream to query
captureStatus
- Returns the stream's capture status
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_STREAM_CAPTURE_IMPLICIT
Description
Return the capture status of hStream via captureStatus. After a successful call,
*captureStatus will contain one of the following:
‣ CU_STREAM_CAPTURE_STATUS_NONE: The stream is not capturing.
‣ CU_STREAM_CAPTURE_STATUS_ACTIVE: The stream is capturing.
‣ CU_STREAM_CAPTURE_STATUS_INVALIDATED: The stream was capturing but an
error has invalidated the capture sequence. The capture sequence must be terminated with
cuStreamEndCapture on the stream where it was initiated in order to continue using hStream.
|
352
Modules
Note that, if this is called on CU_STREAM_LEGACY (the "null stream") while a blocking stream in
the same context is capturing, it will return CUDA_ERROR_STREAM_CAPTURE_IMPLICIT and
*captureStatus is unspecified after the call. The blocking stream capture is not invalidated.
When a blocking stream is capturing, the legacy stream is in an unusable state until the blocking stream
capture is terminated. The legacy stream is not supported for stream capture, but attempted use would
have an implicit dependency on the capturing stream(s).
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuStreamCreate, cuStreamBeginCapture, cuStreamEndCapture
CUresult cuStreamQuery (CUstream hStream)
Determine status of a compute stream.
Parameters
hStream
- Stream to query status of
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE,
CUDA_ERROR_NOT_READY
Description
Returns CUDA_SUCCESS if all operations in the stream specified by hStream have completed, or
CUDA_ERROR_NOT_READY if not.
For the purposes of Unified Memory, a return value of CUDA_SUCCESS is equivalent to having
called cuStreamSynchronize().
Note:
‣ This function uses standard default stream semantics.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
|
353
Modules
cuStreamCreate, cuStreamWaitEvent, cuStreamDestroy, cuStreamSynchronize, cuStreamAddCallback,
cudaStreamQuery
CUresult cuStreamSetAttribute (CUstream hStream,
CUstreamAttrID attr, const CUstreamAttrValue *value)
Sets stream attribute.
Parameters
hStream
attr
value
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_INVALID_HANDLE
Description
Sets attribute attr on hStream from corresponding attribute of value. The updated attribute will
be applied to subsequent work submitted to the stream. It will not affect previously submitted work.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
CUaccessPolicyWindow
CUresult cuStreamSynchronize (CUstream hStream)
Wait until a stream's tasks are completed.
Parameters
hStream
- Stream to wait for
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE
|
354
Modules
Description
Waits until the device has completed all operations in the stream specified by hStream. If the context
was created with the CU_CTX_SCHED_BLOCKING_SYNC flag, the CPU thread will block until the
stream is finished with all of its tasks.
Note:
‣ This function uses standard default stream semantics.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuStreamCreate, cuStreamDestroy, cuStreamWaitEvent, cuStreamQuery, cuStreamAddCallback,
cudaStreamSynchronize
CUresult cuStreamUpdateCaptureDependencies
(CUstream hStream, CUgraphNode *dependencies, size_t
numDependencies, unsigned int flags)
Update the set of dependencies in a capturing stream (11.3+).
Parameters
hStream
- The stream to update
dependencies
- The set of dependencies to add
numDependencies
- The size of the dependencies array
flags
- See above
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_ILLEGAL_STATE
Description
Modifies the dependency set of a capturing stream. The dependency set is the set of nodes that the next
captured node in the stream will depend on.
Valid flags are CU_STREAM_ADD_CAPTURE_DEPENDENCIES and
CU_STREAM_SET_CAPTURE_DEPENDENCIES. These control whether the set
|
355
Modules
passed to the API is added to the existing set or replaces it. A flags value of 0 defaults to
CU_STREAM_ADD_CAPTURE_DEPENDENCIES.
Nodes that are removed from the dependency set via this API do not result in
CUDA_ERROR_STREAM_CAPTURE_UNJOINED if they are unreachable from the stream at
cuStreamEndCapture.
Returns CUDA_ERROR_ILLEGAL_STATE if the stream is not capturing.
This API is new in CUDA 11.3. Developers requiring compatibility across minor versions to CUDA
11.0 should not use this API or provide a fallback.
See also:
cuStreamBeginCapture, cuStreamGetCaptureInfo,
CUresult cuStreamUpdateCaptureDependencies_v2
(CUstream hStream, CUgraphNode *dependencies,
const CUgraphEdgeData *dependencyData, size_t
numDependencies, unsigned int flags)
Update the set of dependencies in a capturing stream (12.3+).
Parameters
hStream
- The stream to update
dependencies
- The set of dependencies to add
dependencyData
- Optional array of data associated with each dependency.
numDependencies
- The size of the dependencies array
flags
- See above
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_ILLEGAL_STATE
Description
Modifies the dependency set of a capturing stream. The dependency set is the set of nodes that the next
captured node in the stream will depend on along with the edge data for those dependencies.
Valid flags are CU_STREAM_ADD_CAPTURE_DEPENDENCIES and
CU_STREAM_SET_CAPTURE_DEPENDENCIES. These control whether the set
|
356
Modules
passed to the API is added to the existing set or replaces it. A flags value of 0 defaults to
CU_STREAM_ADD_CAPTURE_DEPENDENCIES.
Nodes that are removed from the dependency set via this API do not result in
CUDA_ERROR_STREAM_CAPTURE_UNJOINED if they are unreachable from the stream at
cuStreamEndCapture.
Returns CUDA_ERROR_ILLEGAL_STATE if the stream is not capturing.
See also:
cuStreamBeginCapture, cuStreamGetCaptureInfo
CUresult cuStreamWaitEvent (CUstream hStream,
CUevent hEvent, unsigned int Flags)
Make a compute stream wait on an event.
Parameters
hStream
- Stream to wait
hEvent
- Event to wait on (may not be NULL)
Flags
- See CUevent_capture_flags
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE,
Description
Makes all future work submitted to hStream wait for all work captured in hEvent. See
cuEventRecord() for details on what is captured by an event. The synchronization will be performed
efficiently on the device when applicable. hEvent may be from a different context or device than
hStream.
flags include:
‣ CU_EVENT_WAIT_DEFAULT: Default event creation flag.
‣ CU_EVENT_WAIT_EXTERNAL: Event is captured in the graph as an external event node when
performing stream capture. This flag is invalid outside of stream capture.
Note:
|
357
Modules
‣ This function uses standard default stream semantics.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuStreamCreate, cuEventRecord, cuStreamQuery, cuStreamSynchronize, cuStreamAddCallback,
cuStreamDestroy, cudaStreamWaitEvent
CUresult cuThreadExchangeStreamCaptureMode
(CUstreamCaptureMode *mode)
Swaps the stream capture interaction mode for a thread.
Parameters
mode
- Pointer to mode value to swap with the current mode
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_VALUE
Description
Sets the calling thread's stream capture interaction mode to the value contained in *mode, and
overwrites *mode with the previous mode for the thread. To facilitate deterministic behavior across
function or module boundaries, callers are encouraged to use this API in a push-pop fashion:
‎
CUstreamCaptureMode mode = desiredMode;
cuThreadExchangeStreamCaptureMode(&mode);
cuThreadExchangeStreamCaptureMode(&mode); // restore previous mode
During stream capture (see cuStreamBeginCapture), some actions, such as a call to cudaMalloc,
may be unsafe. In the case of cudaMalloc, the operation is not enqueued asynchronously to a
stream, and is not observed by stream capture. Therefore, if the sequence of operations captured via
cuStreamBeginCapture depended on the allocation being replayed whenever the graph is launched, the
captured graph would be invalid.
Therefore, stream capture places restrictions on API calls that can be made within or concurrently to a
cuStreamBeginCapture-cuStreamEndCapture sequence. This behavior can be controlled via this API
and flags to cuStreamBeginCapture.
A thread's mode is one of the following:
‣ CU_STREAM_CAPTURE_MODE_GLOBAL: This is the default mode. If the local thread has an
ongoing capture sequence that was not initiated with CU_STREAM_CAPTURE_MODE_RELAXED
at cuStreamBeginCapture, or if any other thread has a concurrent capture sequence initiated
|
358
Modules
with CU_STREAM_CAPTURE_MODE_GLOBAL, this thread is prohibited from potentially unsafe
API calls.
‣ CU_STREAM_CAPTURE_MODE_THREAD_LOCAL: If the local thread has an ongoing capture
sequence not initiated with CU_STREAM_CAPTURE_MODE_RELAXED, it is prohibited from
potentially unsafe API calls. Concurrent capture sequences in other threads are ignored.
‣ CU_STREAM_CAPTURE_MODE_RELAXED: The local thread is not prohibited from potentially
unsafe API calls. Note that the thread is still prohibited from API calls which necessarily conflict
with stream capture, for example, attempting cuEventQuery on an event that was last recorded
inside a capture sequence.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuStreamBeginCapture
6.19. Event Management
This section describes the event management functions of the low-level CUDA driver application
programming interface.
CUresult cuEventCreate (CUevent *phEvent, unsigned int
Flags)
Creates an event.
Parameters
phEvent
- Returns newly created event
Flags
- Event creation flags
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_OUT_OF_MEMORY
|
359
Modules
Description
Creates an event *phEvent for the current context with the flags specified via Flags. Valid flags
include:
‣ CU_EVENT_DEFAULT: Default event creation flag.
‣ CU_EVENT_BLOCKING_SYNC: Specifies that the created event should use blocking
synchronization. A CPU thread that uses cuEventSynchronize() to wait on an event created with
this flag will block until the event has actually been recorded.
‣ CU_EVENT_DISABLE_TIMING: Specifies that the created event does not need to record
timing data. Events created with this flag specified and the CU_EVENT_BLOCKING_SYNC
flag not specified will provide the best performance when used with cuStreamWaitEvent() and
cuEventQuery().
‣ CU_EVENT_INTERPROCESS: Specifies that the created event may be used as an interprocess
event by cuIpcGetEventHandle(). CU_EVENT_INTERPROCESS must be specified along with
CU_EVENT_DISABLE_TIMING.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuEventRecord, cuEventQuery, cuEventSynchronize, cuEventDestroy, cuEventElapsedTime,
cudaEventCreate, cudaEventCreateWithFlags
CUresult cuEventDestroy (CUevent hEvent)
Destroys an event.
Parameters
hEvent
- Event to destroy
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE
Description
Destroys the event specified by hEvent.
|
360
Modules
An event may be destroyed before it is complete (i.e., while cuEventQuery() would return
CUDA_ERROR_NOT_READY). In this case, the call does not block on completion of the event, and
any associated resources will automatically be released asynchronously at completion.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuEventCreate, cuEventRecord, cuEventQuery, cuEventSynchronize, cuEventElapsedTime,
cudaEventDestroy
CUresult cuEventElapsedTime (float *pMilliseconds,
CUevent hStart, CUevent hEnd)
Computes the elapsed time between two events.
Parameters
pMilliseconds
- Time between hStart and hEnd in ms
hStart
- Starting event
hEnd
- Ending event
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE,
CUDA_ERROR_NOT_READY, CUDA_ERROR_UNKNOWN
Description
Computes the elapsed time between two events (in milliseconds with a resolution of around 0.5
microseconds).
If either event was last recorded in a non-NULL stream, the resulting time may be greater than
expected (even if both used the same stream handle). This happens because the cuEventRecord()
operation takes place asynchronously and there is no guarantee that the measured latency is actually
just between the two events. Any number of other different stream operations could execute in between
the two measured events, thus altering the timing in a significant way.
If cuEventRecord() has not been called on either event then CUDA_ERROR_INVALID_HANDLE
is returned. If cuEventRecord() has been called on both events but one or both of them has not
|
361
Modules
yet been completed (that is, cuEventQuery() would return CUDA_ERROR_NOT_READY
on at least one of the events), CUDA_ERROR_NOT_READY is returned. If either event
was created with the CU_EVENT_DISABLE_TIMING flag, then this function will return
CUDA_ERROR_INVALID_HANDLE.
Note there is a later version of this API, cuEventElapsedTime_v2. It will supplant this version in
CUDA 13.0, which is retained for minor version compatibility.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuEventCreate, cuEventRecord, cuEventQuery, cuEventSynchronize, cuEventDestroy,
cudaEventElapsedTime
CUresult cuEventElapsedTime_v2 (float *pMilliseconds,
CUevent hStart, CUevent hEnd)
Computes the elapsed time between two events.
Parameters
pMilliseconds
- Time between hStart and hEnd in ms
hStart
- Starting event
hEnd
- Ending event
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE,
CUDA_ERROR_NOT_READY, CUDA_ERROR_UNKNOWN
Description
Computes the elapsed time between two events (in milliseconds with a resolution of around 0.5
microseconds). Note this API is not guaranteed to return the latest errors for pending work. As such
this API is intended to serve as an elapsed time calculation only and any polling for completion on the
events to be compared should be done with cuEventQuery instead.
If either event was last recorded in a non-NULL stream, the resulting time may be greater than
expected (even if both used the same stream handle). This happens because the cuEventRecord()
|
362
Modules
operation takes place asynchronously and there is no guarantee that the measured latency is actually
just between the two events. Any number of other different stream operations could execute in between
the two measured events, thus altering the timing in a significant way.
If cuEventRecord() has not been called on either event then CUDA_ERROR_INVALID_HANDLE
is returned. If cuEventRecord() has been called on both events but one or both of them has not
yet been completed (that is, cuEventQuery() would return CUDA_ERROR_NOT_READY
on at least one of the events), CUDA_ERROR_NOT_READY is returned. If either event
was created with the CU_EVENT_DISABLE_TIMING flag, then this function will return
CUDA_ERROR_INVALID_HANDLE.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuEventCreate, cuEventRecord, cuEventQuery, cuEventSynchronize, cuEventDestroy,
cudaEventElapsedTime
CUresult cuEventQuery (CUevent hEvent)
Queries an event's status.
Parameters
hEvent
- Event to query
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_HANDLE, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_NOT_READY
Description
Queries the status of all work currently captured by hEvent. See cuEventRecord() for details on what
is captured by an event.
Returns CUDA_SUCCESS if all captured work has been completed, or
CUDA_ERROR_NOT_READY if any captured work is incomplete.
For the purposes of Unified Memory, a return value of CUDA_SUCCESS is equivalent to having
called cuEventSynchronize().
|
363
Modules
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuEventCreate, cuEventRecord, cuEventSynchronize, cuEventDestroy, cuEventElapsedTime,
cudaEventQuery
CUresult cuEventRecord (CUevent hEvent, CUstream
hStream)
Records an event.
Parameters
hEvent
- Event to record
hStream
- Stream to record event for
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE,
CUDA_ERROR_INVALID_VALUE
Description
Captures in hEvent the contents of hStream at the time of this call. hEvent and hStream must
be from the same context otherwise CUDA_ERROR_INVALID_HANDLE is returned. Calls such as
cuEventQuery() or cuStreamWaitEvent() will then examine or wait for completion of the work that was
captured. Uses of hStream after this call do not modify hEvent. See note on default stream behavior
for what is captured in the default case.
cuEventRecord() can be called multiple times on the same event and will overwrite the previously
captured state. Other APIs such as cuStreamWaitEvent() use the most recently captured state at the
time of the API call, and are not affected by later calls to cuEventRecord(). Before the first call to
cuEventRecord(), an event represents an empty set of work, so for example cuEventQuery() would
return CUDA_SUCCESS.
Note:
‣ This function uses standard default stream semantics.
‣ Note that this function may also return error codes from previous, asynchronous launches.
|
364
Modules
See also:
cuEventCreate, cuEventQuery, cuEventSynchronize, cuStreamWaitEvent, cuEventDestroy,
cuEventElapsedTime, cudaEventRecord, cuEventRecordWithFlags
CUresult cuEventRecordWithFlags (CUevent hEvent,
CUstream hStream, unsigned int flags)
Records an event.
Parameters
hEvent
- Event to record
hStream
- Stream to record event for
flags
- See CUevent_capture_flags
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE,
CUDA_ERROR_INVALID_VALUE
Description
Captures in hEvent the contents of hStream at the time of this call. hEvent and hStream must
be from the same context otherwise CUDA_ERROR_INVALID_HANDLE is returned. Calls such as
cuEventQuery() or cuStreamWaitEvent() will then examine or wait for completion of the work that was
captured. Uses of hStream after this call do not modify hEvent. See note on default stream behavior
for what is captured in the default case.
cuEventRecordWithFlags() can be called multiple times on the same event and will overwrite the
previously captured state. Other APIs such as cuStreamWaitEvent() use the most recently captured
state at the time of the API call, and are not affected by later calls to cuEventRecordWithFlags().
Before the first call to cuEventRecordWithFlags(), an event represents an empty set of work, so for
example cuEventQuery() would return CUDA_SUCCESS.
flags include:
‣ CU_EVENT_RECORD_DEFAULT: Default event creation flag.
‣ CU_EVENT_RECORD_EXTERNAL: Event is captured in the graph as an external event node
when performing stream capture. This flag is invalid outside of stream capture.
Note:
|
365
Modules
‣ This function uses standard default stream semantics.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuEventCreate, cuEventQuery, cuEventSynchronize, cuStreamWaitEvent, cuEventDestroy,
cuEventElapsedTime, cuEventRecord, cudaEventRecord
CUresult cuEventSynchronize (CUevent hEvent)
Waits for an event to complete.
Parameters
hEvent
- Event to wait for
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE
Description
Waits until the completion of all work currently captured in hEvent. See cuEventRecord() for details
on what is captured by an event.
Waiting for an event that was created with the CU_EVENT_BLOCKING_SYNC flag will
cause the calling CPU thread to block until the event has been completed by the device. If the
CU_EVENT_BLOCKING_SYNC flag has not been set, then the CPU thread will busy-wait until the
event has been completed by the device.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuEventCreate, cuEventRecord, cuEventQuery, cuEventDestroy, cuEventElapsedTime,
cudaEventSynchronize
6.20. External Resource Interoperability
This section describes the external resource interoperability functions of the low-level CUDA driver
application programming interface.
|
366
Modules
CUresult cuDestroyExternalMemory (CUexternalMemory
extMem)
Destroys an external memory object.
Parameters
extMem
- External memory object to be destroyed
Returns
CUDA_SUCCESS, CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_INVALID_HANDLE
Description
Destroys the specified external memory object. Any existing buffers and CUDA mipmapped arrays
mapped onto this object must no longer be used and must be explicitly freed using cuMemFree and
cuMipmappedArrayDestroy respectively.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuImportExternalMemory, cuExternalMemoryGetMappedBuffer,
cuExternalMemoryGetMappedMipmappedArray
CUresult cuDestroyExternalSemaphore
(CUexternalSemaphore extSem)
Destroys an external semaphore.
Parameters
extSem
- External semaphore to be destroyed
Returns
CUDA_SUCCESS, CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_INVALID_HANDLE
Description
Destroys an external semaphore object and releases any references to the underlying resource. Any
outstanding signals or waits must have completed before the semaphore is destroyed.
|
367
Modules
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuImportExternalSemaphore, cuSignalExternalSemaphoresAsync, cuWaitExternalSemaphoresAsync
CUresult cuExternalMemoryGetMappedBuffer
(CUdeviceptr *devPtr, CUexternalMemory extMem,
const CUDA_EXTERNAL_MEMORY_BUFFER_DESC
*bufferDesc)
Maps a buffer onto an imported memory object.
Parameters
devPtr
- Returned device pointer to buffer
extMem
- Handle to external memory object
bufferDesc
- Buffer descriptor
Returns
CUDA_SUCCESS, CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_INVALID_HANDLE
Description
Maps a buffer onto an imported memory object and returns a device pointer in devPtr.
The properties of the buffer being mapped must be described in bufferDesc. The
CUDA_EXTERNAL_MEMORY_BUFFER_DESC structure is defined as follows:
‎
typedef struct CUDA_EXTERNAL_MEMORY_BUFFER_DESC_st {
unsigned long long offset;
unsigned long long size;
unsigned int flags;
} CUDA_EXTERNAL_MEMORY_BUFFER_DESC;
where CUDA_EXTERNAL_MEMORY_BUFFER_DESC::offset is the offset in the memory object
where the buffer's base address is. CUDA_EXTERNAL_MEMORY_BUFFER_DESC::size is the size
of the buffer. CUDA_EXTERNAL_MEMORY_BUFFER_DESC::flags must be zero.
The offset and size have to be suitably aligned to match the requirements of the external API. Mapping
two buffers whose ranges overlap may or may not result in the same virtual address being returned
|
368
Modules
for the overlapped portion. In such cases, the application must ensure that all accesses to that region
from the GPU are volatile. Otherwise writes made via one address are not guaranteed to be visible via
the other address, even if they're issued by the same thread. It is recommended that applications map
the combined range instead of mapping separate buffers and then apply the appropriate offsets to the
returned pointer to derive the individual buffers.
The returned pointer devPtr must be freed using cuMemFree.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuImportExternalMemory, cuDestroyExternalMemory,
cuExternalMemoryGetMappedMipmappedArray
CUresult cuExternalMemoryGetMappedMipmappedArray
(CUmipmappedArray *mipmap, CUexternalMemory
extMem, const
CUDA_EXTERNAL_MEMORY_MIPMAPPED_ARRAY_DESC
*mipmapDesc)
Maps a CUDA mipmapped array onto an external memory object.
Parameters
mipmap
- Returned CUDA mipmapped array
extMem
- Handle to external memory object
mipmapDesc
- CUDA array descriptor
Returns
CUDA_SUCCESS, CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_INVALID_HANDLE
Description
Maps a CUDA mipmapped array onto an external object and returns a handle to it in mipmap.
|
369
Modules
The properties of the CUDA mipmapped array being mapped must be described in mipmapDesc. The
structure CUDA_EXTERNAL_MEMORY_MIPMAPPED_ARRAY_DESC is defined as follows:
‎
typedef struct CUDA_EXTERNAL_MEMORY_MIPMAPPED_ARRAY_DESC_st {
unsigned long long offset;
CUDA_ARRAY3D_DESCRIPTOR arrayDesc;
unsigned int numLevels;
} CUDA_EXTERNAL_MEMORY_MIPMAPPED_ARRAY_DESC;
where CUDA_EXTERNAL_MEMORY_MIPMAPPED_ARRAY_DESC::offset
is the offset in the memory object where the base level of the mipmap chain is.
CUDA_EXTERNAL_MEMORY_MIPMAPPED_ARRAY_DESC::arrayDesc describes
the format, dimensions and type of the base level of the mipmap chain. For further details
on these parameters, please refer to the documentation for cuMipmappedArrayCreate.
Note that if the mipmapped array is bound as a color target in the graphics API,
then the flag CUDA_ARRAY3D_COLOR_ATTACHMENT must be specified in
CUDA_EXTERNAL_MEMORY_MIPMAPPED_ARRAY_DESC::arrayDesc::Flags.
CUDA_EXTERNAL_MEMORY_MIPMAPPED_ARRAY_DESC::numLevels specifies the total
number of levels in the mipmap chain.
If extMem was imported from a handle of type
CU_EXTERNAL_MEMORY_HANDLE_TYPE_NVSCIBUF, then
CUDA_EXTERNAL_MEMORY_MIPMAPPED_ARRAY_DESC::numLevels must be equal to 1.
The returned CUDA mipmapped array must be freed using cuMipmappedArrayDestroy.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuImportExternalMemory, cuDestroyExternalMemory, cuExternalMemoryGetMappedBuffer
CUresult cuImportExternalMemory
(CUexternalMemory *extMem_out, const
CUDA_EXTERNAL_MEMORY_HANDLE_DESC
*memHandleDesc)
Imports an external memory object.
Parameters
extMem_out
- Returned handle to an external memory object
memHandleDesc
- Memory import handle descriptor
|
370
Modules
Returns
CUDA_SUCCESS, CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_INVALID_HANDLE, CUDA_ERROR_OPERATING_SYSTEM
Description
Imports an externally allocated memory object and returns a handle to that in extMem_out.
The properties of the handle being imported must be described in memHandleDesc. The
CUDA_EXTERNAL_MEMORY_HANDLE_DESC structure is defined as follows:
‎
typedef struct CUDA_EXTERNAL_MEMORY_HANDLE_DESC_st {
CUexternalMemoryHandleType type;
union {
int fd;
struct {
void *handle;
const void *name;
} win32;
const void *nvSciBufObject;
} handle;
unsigned long long size;
unsigned int flags;
} CUDA_EXTERNAL_MEMORY_HANDLE_DESC;
where CUDA_EXTERNAL_MEMORY_HANDLE_DESC::type specifies the type of handle being
imported. CUexternalMemoryHandleType is defined as:
‎
typedef enum CUexternalMemoryHandleType_enum {
CU_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD
= 1,
CU_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32
= 2,
CU_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_KMT
= 3,
CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D12_HEAP
= 4,
CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D12_RESOURCE
= 5,
CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D11_RESOURCE
= 6,
CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D11_RESOURCE_KMT = 7,
CU_EXTERNAL_MEMORY_HANDLE_TYPE_NVSCIBUF
= 8,
} CUexternalMemoryHandleType;
If CUDA_EXTERNAL_MEMORY_HANDLE_DESC::type is
CU_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD, then
CUDA_EXTERNAL_MEMORY_HANDLE_DESC::handle::fd must be a valid file descriptor
referencing a memory object. Ownership of the file descriptor is transferred to the CUDA driver when
the handle is imported successfully. Performing any operations on the file descriptor after it is imported
results in undefined behavior.
If CUDA_EXTERNAL_MEMORY_HANDLE_DESC::type is
CU_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32, then exactly one
of CUDA_EXTERNAL_MEMORY_HANDLE_DESC::handle::win32::handle and
CUDA_EXTERNAL_MEMORY_HANDLE_DESC::handle::win32::name must not be NULL. If
CUDA_EXTERNAL_MEMORY_HANDLE_DESC::handle::win32::handle is not NULL, then it must
represent a valid shared NT handle that references a memory object. Ownership of this handle is not
transferred to CUDA after the import operation, so the application must release the handle using the
appropriate system call. If CUDA_EXTERNAL_MEMORY_HANDLE_DESC::handle::win32::name
|
371
Modules
is not NULL, then it must point to a NULL-terminated array of UTF-16 characters that refers to a
memory object.
If CUDA_EXTERNAL_MEMORY_HANDLE_DESC::type is
CU_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_KMT, then
CUDA_EXTERNAL_MEMORY_HANDLE_DESC::handle::win32::handle must be non-NULL
and CUDA_EXTERNAL_MEMORY_HANDLE_DESC::handle::win32::name must be NULL. The
handle specified must be a globally shared KMT handle. This handle does not hold a reference to the
underlying object, and thus will be invalid when all references to the memory object are destroyed.
If CUDA_EXTERNAL_MEMORY_HANDLE_DESC::type is
CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D12_HEAP, then exactly one
of CUDA_EXTERNAL_MEMORY_HANDLE_DESC::handle::win32::handle and
CUDA_EXTERNAL_MEMORY_HANDLE_DESC::handle::win32::name must not be NULL. If
CUDA_EXTERNAL_MEMORY_HANDLE_DESC::handle::win32::handle is not NULL, then it
must represent a valid shared NT handle that is returned by ID3D12Device::CreateSharedHandle
when referring to a ID3D12Heap object. This handle holds a reference to the underlying object. If
CUDA_EXTERNAL_MEMORY_HANDLE_DESC::handle::win32::name is not NULL, then it must
point to a NULL-terminated array of UTF-16 characters that refers to a ID3D12Heap object.
If CUDA_EXTERNAL_MEMORY_HANDLE_DESC::type is
CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D12_RESOURCE, then exactly
one of CUDA_EXTERNAL_MEMORY_HANDLE_DESC::handle::win32::handle and
CUDA_EXTERNAL_MEMORY_HANDLE_DESC::handle::win32::name must not be NULL. If
CUDA_EXTERNAL_MEMORY_HANDLE_DESC::handle::win32::handle is not NULL, then it
must represent a valid shared NT handle that is returned by ID3D12Device::CreateSharedHandle
when referring to a ID3D12Resource object. This handle holds a reference to the underlying object. If
CUDA_EXTERNAL_MEMORY_HANDLE_DESC::handle::win32::name is not NULL, then it must
point to a NULL-terminated array of UTF-16 characters that refers to a ID3D12Resource object.
If CUDA_EXTERNAL_MEMORY_HANDLE_DESC::type is
CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D11_RESOURCE, then
CUDA_EXTERNAL_MEMORY_HANDLE_DESC::handle::win32::handle must represent a valid
shared NT handle that is returned by IDXGIResource1::CreateSharedHandle when referring to a
ID3D11Resource object. If CUDA_EXTERNAL_MEMORY_HANDLE_DESC::handle::win32::name
is not NULL, then it must point to a NULL-terminated array of UTF-16 characters that refers to a
ID3D11Resource object.
If CUDA_EXTERNAL_MEMORY_HANDLE_DESC::type is
CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D11_RESOURCE_KMT,
then CUDA_EXTERNAL_MEMORY_HANDLE_DESC::handle::win32::handle
must represent a valid shared KMT handle that is returned by
IDXGIResource::GetSharedHandle when referring to a ID3D11Resource object and
CUDA_EXTERNAL_MEMORY_HANDLE_DESC::handle::win32::name must be NULL.
If CUDA_EXTERNAL_MEMORY_HANDLE_DESC::type is
CU_EXTERNAL_MEMORY_HANDLE_TYPE_NVSCIBUF, then
|
372
Modules
CUDA_EXTERNAL_MEMORY_HANDLE_DESC::handle::nvSciBufObject
must be non-NULL and reference a valid NvSciBuf object. If the NvSciBuf object
imported into CUDA is also mapped by other drivers, then the application must
use cuWaitExternalSemaphoresAsync or cuSignalExternalSemaphoresAsync as
appropriate barriers to maintain coherence between CUDA and the other drivers. See
CUDA_EXTERNAL_SEMAPHORE_SIGNAL_SKIP_NVSCIBUF_MEMSYNC and
CUDA_EXTERNAL_SEMAPHORE_WAIT_SKIP_NVSCIBUF_MEMSYNC for memory
synchronization.
The size of the memory object must be specified in
CUDA_EXTERNAL_MEMORY_HANDLE_DESC::size.
Specifying the flag CUDA_EXTERNAL_MEMORY_DEDICATED in
CUDA_EXTERNAL_MEMORY_HANDLE_DESC::flags indicates that the resource is a
dedicated resource. The definition of what a dedicated resource is outside the scope of this
extension. This flag must be set if CUDA_EXTERNAL_MEMORY_HANDLE_DESC::type is
one of the following: CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D12_RESOURCE
CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D11_RESOURCE
CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D11_RESOURCE_KMT
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ If the Vulkan memory imported into CUDA is mapped on the CPU then the application must use
vkInvalidateMappedMemoryRanges/vkFlushMappedMemoryRanges as well as appropriate Vulkan
pipeline barriers to maintain coherence between CPU and GPU. For more information on these
APIs, please refer to "Synchronization and Cache Control" chapter from Vulkan specification.
See also:
cuDestroyExternalMemory, cuExternalMemoryGetMappedBuffer,
cuExternalMemoryGetMappedMipmappedArray
CUresult cuImportExternalSemaphore
(CUexternalSemaphore *extSem_out, const
CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC
*semHandleDesc)
Imports an external semaphore.
Parameters
extSem_out
- Returned handle to an external semaphore
|
373
Modules
semHandleDesc
- Semaphore import handle descriptor
Returns
CUDA_SUCCESS, CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_NOT_SUPPORTED,
CUDA_ERROR_INVALID_HANDLE, CUDA_ERROR_OPERATING_SYSTEM
Description
Imports an externally allocated synchronization object and returns a handle to that in extSem_out.
The properties of the handle being imported must be described in semHandleDesc. The
CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC is defined as follows:
‎
typedef struct CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC_st {
CUexternalSemaphoreHandleType type;
union {
int fd;
struct {
void *handle;
const void *name;
} win32;
const void* NvSciSyncObj;
} handle;
unsigned int flags;
} CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC;
where CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::type specifies the type of handle being
imported. CUexternalSemaphoreHandleType is defined as:
‎
typedef enum CUexternalSemaphoreHandleType_enum {
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD
= 1,
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32
= 2,
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT
= 3,
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D12_FENCE
= 4,
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_FENCE
= 5,
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_NVSCISYNC
= 6,
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_KEYED_MUTEX
= 7,
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_KEYED_MUTEX_KMT
= 8,
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_TIMELINE_SEMAPHORE_FD
= 9,
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_TIMELINE_SEMAPHORE_WIN32 = 10
} CUexternalSemaphoreHandleType;
If CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::type is
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD, then
CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::handle::fd must be a valid file descriptor
referencing a synchronization object. Ownership of the file descriptor is transferred to the CUDA
driver when the handle is imported successfully. Performing any operations on the file descriptor after
it is imported results in undefined behavior.
If CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::type is
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32, then exactly one
of CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::handle::win32::handle and
CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::handle::win32::name must not be
NULL. If CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::handle::win32::handle
is not NULL, then it must represent a valid shared NT handle that references a
|
374
Modules
synchronization object. Ownership of this handle is not transferred to CUDA after the import
operation, so the application must release the handle using the appropriate system call. If
CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::handle::win32::name is not NULL, then it
must name a valid synchronization object.
If CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::type is
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT, then
CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::handle::win32::handle must be non-NULL
and CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::handle::win32::name must be NULL.
The handle specified must be a globally shared KMT handle. This handle does not hold a reference to
the underlying object, and thus will be invalid when all references to the synchronization object are
destroyed.
If CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::type is
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D12_FENCE, then exactly one
of CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::handle::win32::handle and
CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::handle::win32::name must not be NULL.
If CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::handle::win32::handle is not NULL, then
it must represent a valid shared NT handle that is returned by ID3D12Device::CreateSharedHandle
when referring to a ID3D12Fence object. This handle holds a reference to the underlying object. If
CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::handle::win32::name is not NULL, then it
must name a valid synchronization object that refers to a valid ID3D12Fence object.
If CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::type is
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_FENCE, then
CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::handle::win32::handle represents
a valid shared NT handle that is returned by ID3D11Fence::CreateSharedHandle. If
CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::handle::win32::name is not NULL, then it
must name a valid synchronization object that refers to a valid ID3D11Fence object.
If CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::type is
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_NVSCISYNC, then
CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::handle::nvSciSyncObj represents a valid
NvSciSyncObj.
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_KEYED_MUTEX, then
CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::handle::win32::handle represents a valid
shared NT handle that is returned by IDXGIResource1::CreateSharedHandle when referring to a
IDXGIKeyedMutex object. If
CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::handle::win32::name is not NULL, then it
must name a valid synchronization object that refers to a valid IDXGIKeyedMutex object.
If CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::type is
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_KEYED_MUTEX_KMT, then
CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::handle::win32::handle represents a valid
shared KMT handle that is returned by IDXGIResource::GetSharedHandle when referring to a
|
375
Modules
IDXGIKeyedMutex object and
CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::handle::win32::name must be NULL.
If CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::type is
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_TIMELINE_SEMAPHORE_FD, then
CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::handle::fd must be a valid file descriptor
referencing a synchronization object. Ownership of the file descriptor is transferred to the CUDA
driver when the handle is imported successfully. Performing any operations on the file descriptor after
it is imported results in undefined behavior.
If CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::type is
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_TIMELINE_SEMAPHORE_WIN32, then
exactly one of CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::handle::win32::handle
and CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::handle::win32::name must not
be NULL. If CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::handle::win32::handle
is not NULL, then it must represent a valid shared NT handle that references a
synchronization object. Ownership of this handle is not transferred to CUDA after the import
operation, so the application must release the handle using the appropriate system call. If
CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC::handle::win32::name is not NULL, then it
must name a valid synchronization object.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuDestroyExternalSemaphore, cuSignalExternalSemaphoresAsync, cuWaitExternalSemaphoresAsync
CUresult cuSignalExternalSemaphoresAsync
(const CUexternalSemaphore *extSemArray, const
CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS
*paramsArray, unsigned int numExtSems, CUstream
stream)
Signals a set of external semaphore objects.
Parameters
extSemArray
- Set of external semaphores to be signaled
paramsArray
- Array of semaphore parameters
|
376
Modules
numExtSems
- Number of semaphores to signal
stream
- Stream to enqueue the signal operations in
Returns
CUDA_SUCCESS, CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_INVALID_HANDLE,
CUDA_ERROR_NOT_SUPPORTED
Description
Enqueues a signal operation on a set of externally allocated semaphore object in the specified stream.
The operations will be executed when all prior operations in the stream complete.
The exact semantics of signaling a semaphore depends on the type of the object.
If the semaphore object is any one of the following types:
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD,
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32,
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT then signaling the
semaphore will set it to the signaled state.
If the semaphore object is any one of the following types:
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D12_FENCE,
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_FENCE,
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_TIMELINE_SEMAPHORE_FD,
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_TIMELINE_SEMAPHORE_WIN32 then the
semaphore will be set to the value specified in
CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS::params::fence::value.
If the semaphore object is of the type
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_NVSCISYNC this API sets
CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS::params::nvSciSync::fence to a
value that can be used by subsequent waiters of the same NvSciSync object to order operations
with those currently submitted in stream. Such an update will overwrite previous contents
of CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS::params::nvSciSync::fence.
By default, signaling such an external semaphore object causes appropriate memory
synchronization operations to be performed over all external memory objects that are
imported as CU_EXTERNAL_MEMORY_HANDLE_TYPE_NVSCIBUF. This ensures
that any subsequent accesses made by other importers of the same set of NvSciBuf
memory object(s) are coherent. These operations can be skipped by specifying the flag
CUDA_EXTERNAL_SEMAPHORE_SIGNAL_SKIP_NVSCIBUF_MEMSYNC, which can
be used as a performance optimization when data coherency is not required. But specifying
this flag in scenarios where data coherency is required results in undefined behavior. Also, for
semaphore object of the type CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_NVSCISYNC,
if the NvSciSyncAttrList used to create the NvSciSyncObj had not set the flags in
|
377
Modules
cuDeviceGetNvSciSyncAttributes to CUDA_NVSCISYNC_ATTR_SIGNAL, this API will
return CUDA_ERROR_NOT_SUPPORTED. NvSciSyncFence associated with semaphore
object of the type CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_NVSCISYNC can be
deterministic. For this the NvSciSyncAttrList used to create the semaphore object must have value
of NvSciSyncAttrKey_RequireDeterministicFences key set to true. Deterministic fences allow users
to enqueue a wait over the semaphore object even before corresponding signal is enqueued. For such
a semaphore object, CUDA guarantees that each signal operation will increment the fence value by
'1'. Users are expected to track count of signals enqueued on the semaphore object and insert waits
accordingly. When such a semaphore object is signaled from multiple streams, due to concurrent
stream execution, it is possible that the order in which the semaphore gets signaled is indeterministic.
This could lead to waiters of the semaphore getting unblocked incorrectly. Users are expected to handle
such situations, either by not using the same semaphore object with deterministic fence support enabled
in different streams or by adding explicit dependency amongst such streams so that the semaphore is
signaled in order.
If the semaphore object is any one of the following types:
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_KEYED_MUTEX,
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_KEYED_MUTEX_KMT
then the keyed mutex will be released with the key specified in
CUDA_EXTERNAL_SEMAPHORE_PARAMS::params::keyedmutex::key.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuImportExternalSemaphore, cuDestroyExternalSemaphore, cuWaitExternalSemaphoresAsync
CUresult cuWaitExternalSemaphoresAsync (const
CUexternalSemaphore *extSemArray, const
CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS
*paramsArray, unsigned int numExtSems, CUstream
stream)
Waits on a set of external semaphore objects.
Parameters
extSemArray
- External semaphores to be waited on
paramsArray
- Array of semaphore parameters
|
378
Modules
numExtSems
- Number of semaphores to wait on
stream
- Stream to enqueue the wait operations in
Returns
CUDA_SUCCESS, CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_INVALID_HANDLE,
CUDA_ERROR_NOT_SUPPORTED, CUDA_ERROR_TIMEOUT
Description
Enqueues a wait operation on a set of externally allocated semaphore object in the specified stream.
The operations will be executed when all prior operations in the stream complete.
The exact semantics of waiting on a semaphore depends on the type of the object.
If the semaphore object is any one of the following types:
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD,
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32,
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT then waiting on the
semaphore will wait until the semaphore reaches the signaled state. The semaphore will then be reset to
the unsignaled state. Therefore for every signal operation, there can only be one wait operation.
If the semaphore object is any one of the following types:
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D12_FENCE,
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_FENCE,
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_TIMELINE_SEMAPHORE_FD,
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_TIMELINE_SEMAPHORE_WIN32 then
waiting on the semaphore will wait until the value of the semaphore is greater than or equal to
CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS::params::fence::value.
If the semaphore object is of the type
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_NVSCISYNC then, waiting on the semaphore
will wait until the
CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS::params::nvSciSync::fence
is signaled by the signaler of the NvSciSyncObj that was associated with this semaphore
object. By default, waiting on such an external semaphore object causes appropriate
memory synchronization operations to be performed over all external memory objects
that are imported as CU_EXTERNAL_MEMORY_HANDLE_TYPE_NVSCIBUF. This
ensures that any subsequent accesses made by other importers of the same set of NvSciBuf
memory object(s) are coherent. These operations can be skipped by specifying the flag
CUDA_EXTERNAL_SEMAPHORE_WAIT_SKIP_NVSCIBUF_MEMSYNC, which can be used as
a performance optimization when data coherency is not required. But specifying this flag in scenarios
where data coherency is required results in undefined behavior. Also, for semaphore object of the
type CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_NVSCISYNC, if the NvSciSyncAttrList
|
379
Modules
used to create the NvSciSyncObj had not set the flags in cuDeviceGetNvSciSyncAttributes to
CUDA_NVSCISYNC_ATTR_WAIT, this API will return CUDA_ERROR_NOT_SUPPORTED.
If the semaphore object is any one of the following types:
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_KEYED_MUTEX,
CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_KEYED_MUTEX_KMT
then the keyed mutex will be acquired when it is released with the key specified in
CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS::params::keyedmutex::key or until the
timeout specified by
CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS::params::keyedmutex::timeoutMs has lapsed.
The timeout interval can either be a finite value specified in milliseconds or an infinite value. In case
an infinite value is specified the timeout never elapses. The windows INFINITE macro must be used to
specify infinite timeout.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuImportExternalSemaphore, cuDestroyExternalSemaphore, cuSignalExternalSemaphoresAsync
6.21. Stream Memory Operations
This section describes the stream memory operations of the low-level CUDA driver application
programming interface.
Support for the CU_STREAM_WAIT_VALUE_NOR flag can be queried with
CU_DEVICE_ATTRIBUTE_CAN_USE_STREAM_WAIT_VALUE_NOR_V2.
Support for the cuStreamWriteValue64() and cuStreamWaitValue64() functions, as well as for the
CU_STREAM_MEM_OP_WAIT_VALUE_64 and CU_STREAM_MEM_OP_WRITE_VALUE_64
flags, can be queried with CU_DEVICE_ATTRIBUTE_CAN_USE_64_BIT_STREAM_MEM_OPS.
Support for both CU_STREAM_WAIT_VALUE_FLUSH and
CU_STREAM_MEM_OP_FLUSH_REMOTE_WRITES requires dedicated
platform hardware features and can be queried with cuDeviceGetAttribute() and
CU_DEVICE_ATTRIBUTE_CAN_FLUSH_REMOTE_WRITES.
Note that all memory pointers passed as parameters to these operations are device pointers. Where
necessary a device pointer should be obtained, for example with cuMemHostGetDevicePointer().
None of the operations accepts pointers to managed memory buffers (cuMemAllocManaged).
Note:
|
380
Modules
Warning: Improper use of these APIs may deadlock the application. Synchronization ordering
established through these APIs is not visible to CUDA. CUDA tasks that are (even indirectly) ordered by
these APIs should also have that order expressed with CUDA-visible dependencies such as events. This
ensures that the scheduler does not serialize them in an improper order.
CUresult cuStreamBatchMemOp (CUstream stream,
unsigned int count, CUstreamBatchMemOpParams
*paramArray, unsigned int flags)
Batch operations to synchronize the stream via memory operations.
Parameters
stream
The stream to enqueue the operations in.
count
The number of operations in the array. Must be less than 256.
paramArray
The types and parameters of the individual operations.
flags
Reserved for future expansion; must be 0.
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_NOT_SUPPORTED
Description
This is a batch version of cuStreamWaitValue32() and cuStreamWriteValue32(). Batching operations
may avoid some performance overhead in both the API call and the device execution versus adding
them to the stream in separate API calls. The operations are enqueued in the order they appear in the
array.
See CUstreamBatchMemOpType for the full set of supported operations, and cuStreamWaitValue32(),
cuStreamWaitValue64(), cuStreamWriteValue32(), and cuStreamWriteValue64() for details of specific
operations.
See related APIs for details on querying support for specific operations.
Note:
Warning: Improper use of this API may deadlock the application. Synchronization ordering established
through this API is not visible to CUDA. CUDA tasks that are (even indirectly) ordered by this
API should also have that order expressed with CUDA-visible dependencies such as events. This
ensures that the scheduler does not serialize them in an improper order. For more information, see the
|
381
Modules
Stream Memory Operations section in the programming guide(https://docs.nvidia.com/cuda/cuda-c-
programming-guide/index.html).
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuStreamWaitValue32, cuStreamWaitValue64, cuStreamWriteValue32, cuStreamWriteValue64,
cuMemHostRegister
CUresult cuStreamWaitValue32 (CUstream stream,
CUdeviceptr addr, cuuint32_t value, unsigned int flags)
Wait on a memory location.
Parameters
stream
The stream to synchronize on the memory location.
addr
The memory location to wait on.
value
The value to compare with the memory location.
flags
See CUstreamWaitValue_flags.
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_NOT_SUPPORTED
Description
Enqueues a synchronization of the stream on the given memory location. Work ordered after the
operation will block until the given condition on the memory is satisfied. By default, the condition is to
wait for (int32_t)(*addr - value) >= 0, a cyclic greater-or-equal. Other condition types can be specified
via flags.
If the memory was registered via cuMemHostRegister(), the device pointer should be obtained
with cuMemHostGetDevicePointer(). This function cannot be used with managed memory
(cuMemAllocManaged).
|
382
Modules
Support for CU_STREAM_WAIT_VALUE_NOR can be queried with cuDeviceGetAttribute() and
CU_DEVICE_ATTRIBUTE_CAN_USE_STREAM_WAIT_VALUE_NOR_V2.
Note:
Warning: Improper use of this API may deadlock the application. Synchronization ordering established
through this API is not visible to CUDA. CUDA tasks that are (even indirectly) ordered by this API
should also have that order expressed with CUDA-visible dependencies such as events. This ensures that
the scheduler does not serialize them in an improper order.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuStreamWaitValue64, cuStreamWriteValue32, cuStreamWriteValue64, cuStreamBatchMemOp,
cuMemHostRegister, cuStreamWaitEvent
CUresult cuStreamWaitValue64 (CUstream stream,
CUdeviceptr addr, cuuint64_t value, unsigned int flags)
Wait on a memory location.
Parameters
stream
The stream to synchronize on the memory location.
addr
The memory location to wait on.
value
The value to compare with the memory location.
flags
See CUstreamWaitValue_flags.
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_NOT_SUPPORTED
Description
Enqueues a synchronization of the stream on the given memory location. Work ordered after the
operation will block until the given condition on the memory is satisfied. By default, the condition is to
wait for (int64_t)(*addr - value) >= 0, a cyclic greater-or-equal. Other condition types can be specified
via flags.
|
383
Modules
If the memory was registered via cuMemHostRegister(), the device pointer should be obtained with
cuMemHostGetDevicePointer().
Support for this can be queried with cuDeviceGetAttribute() and
CU_DEVICE_ATTRIBUTE_CAN_USE_64_BIT_STREAM_MEM_OPS.
Note:
Warning: Improper use of this API may deadlock the application. Synchronization ordering established
through this API is not visible to CUDA. CUDA tasks that are (even indirectly) ordered by this API
should also have that order expressed with CUDA-visible dependencies such as events. This ensures that
the scheduler does not serialize them in an improper order.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuStreamWaitValue32, cuStreamWriteValue32, cuStreamWriteValue64, cuStreamBatchMemOp,
cuMemHostRegister, cuStreamWaitEvent
CUresult cuStreamWriteValue32 (CUstream stream,
CUdeviceptr addr, cuuint32_t value, unsigned int flags)
Write a value to memory.
Parameters
stream
The stream to do the write in.
addr
The device address to write to.
value
The value to write.
flags
See CUstreamWriteValue_flags.
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_NOT_SUPPORTED
Description
Write a value to memory.
|
384
Modules
If the memory was registered via cuMemHostRegister(), the device pointer should be obtained
with cuMemHostGetDevicePointer(). This function cannot be used with managed memory
(cuMemAllocManaged).
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuStreamWriteValue64, cuStreamWaitValue32, cuStreamWaitValue64, cuStreamBatchMemOp,
cuMemHostRegister, cuEventRecord
CUresult cuStreamWriteValue64 (CUstream stream,
CUdeviceptr addr, cuuint64_t value, unsigned int flags)
Write a value to memory.
Parameters
stream
The stream to do the write in.
addr
The device address to write to.
value
The value to write.
flags
See CUstreamWriteValue_flags.
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_NOT_SUPPORTED
Description
Write a value to memory.
If the memory was registered via cuMemHostRegister(), the device pointer should be obtained with
cuMemHostGetDevicePointer().
Support for this can be queried with cuDeviceGetAttribute() and
CU_DEVICE_ATTRIBUTE_CAN_USE_64_BIT_STREAM_MEM_OPS.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
|
385
Modules
See also:
cuStreamWriteValue32, cuStreamWaitValue32, cuStreamWaitValue64, cuStreamBatchMemOp,
cuMemHostRegister, cuEventRecord
6.22. Execution Control
This section describes the execution control functions of the low-level CUDA driver application
programming interface.
CUresult cuFuncGetAttribute (int *pi,
CUfunction_attribute attrib, CUfunction hfunc)
Returns information about a function.
Parameters
pi
- Returned attribute value
attrib
- Attribute requested
hfunc
- Function to query attribute of
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE,
CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_FUNCTION_NOT_LOADED
Description
Returns in *pi the integer value of the attribute attrib on the kernel given by hfunc. The
supported attributes are:
‣ CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK: The maximum number of threads
per block, beyond which a launch of the function would fail. This number depends on both the
function and the device on which the function is currently loaded.
‣ CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES: The size in bytes of statically-allocated
shared memory per block required by this function. This does not include dynamically-allocated
shared memory requested by the user at runtime.
‣ CU_FUNC_ATTRIBUTE_CONST_SIZE_BYTES: The size in bytes of user-allocated constant
memory required by this function.
|
386
Modules
‣ CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES: The size in bytes of local memory used by
each thread of this function.
‣ CU_FUNC_ATTRIBUTE_NUM_REGS: The number of registers used by each thread of this
function.
‣ CU_FUNC_ATTRIBUTE_PTX_VERSION: The PTX virtual architecture version for which the
function was compiled. This value is the major PTX version * 10 + the minor PTX version, so a
PTX version 1.3 function would return the value 13. Note that this may return the undefined value
of 0 for cubins compiled prior to CUDA 3.0.
‣ CU_FUNC_ATTRIBUTE_BINARY_VERSION: The binary architecture version for which the
function was compiled. This value is the major binary version * 10 + the minor binary version, so
a binary version 1.3 function would return the value 13. Note that this will return a value of 10 for
legacy cubins that do not have a properly-encoded binary architecture version.
‣ CU_FUNC_CACHE_MODE_CA: The attribute to indicate whether the function has been
compiled with user specified option "-Xptxas --dlcm=ca" set .
‣ CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES: The maximum size in
bytes of dynamically-allocated shared memory.
‣ CU_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT: Preferred shared
memory-L1 cache split ratio in percent of total shared memory.
‣ CU_FUNC_ATTRIBUTE_CLUSTER_SIZE_MUST_BE_SET: If this attribute is set, the kernel
must launch with a valid cluster size specified.
‣ CU_FUNC_ATTRIBUTE_REQUIRED_CLUSTER_WIDTH: The required cluster width in
blocks.
‣ CU_FUNC_ATTRIBUTE_REQUIRED_CLUSTER_HEIGHT: The required cluster height in
blocks.
‣ CU_FUNC_ATTRIBUTE_REQUIRED_CLUSTER_DEPTH: The required cluster depth in
blocks.
‣ CU_FUNC_ATTRIBUTE_NON_PORTABLE_CLUSTER_SIZE_ALLOWED: Indicates whether
the function can be launched with non-portable cluster size. 1 is allowed, 0 is disallowed. A
non-portable cluster size may only function on the specific SKUs the program is tested on. The
launch might fail if the program is run on a different hardware platform. CUDA API provides
cudaOccupancyMaxActiveClusters to assist with checking whether the desired size can be
launched on the current device. A portable cluster size is guaranteed to be functional on all
compute capabilities higher than the target compute capability. The portable cluster size for sm_90
is 8 blocks per cluster. This value may increase for future compute capabilities. The specific
hardware unit may support higher cluster sizes that’s not guaranteed to be portable.
‣ CU_FUNC_ATTRIBUTE_CLUSTER_SCHEDULING_POLICY_PREFERENCE: The block
scheduling policy of a function. The value type is CUclusterSchedulingPolicy.
With a few execeptions, function attributes may also be queried on unloaded function handles returned
from cuModuleEnumerateFunctions. CUDA_ERROR_FUNCTION_NOT_LOADED is returned if the
attribute requires a fully loaded function but the function is not loaded. The loading state of a function
|
387
Modules
may be queried using cuFuncIsloaded. cuFuncLoad may be called to explicitly load a function before
querying the following attributes that require the function to be loaded:
‣ CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK
‣ CU_FUNC_ATTRIBUTE_CONST_SIZE_BYTES
‣ CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuCtxGetCacheConfig, cuCtxSetCacheConfig, cuFuncSetCacheConfig, cuLaunchKernel,
cudaFuncGetAttributes, cudaFuncSetAttribute, cuFuncIsLoaded, cuFuncLoad, cuKernelGetAttribute
CUresult cuFuncGetModule (CUmodule *hmod,
CUfunction hfunc)
Returns a module handle.
Parameters
hmod
- Returned module handle
hfunc
- Function to retrieve module for
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_NOT_FOUND
Description
Returns in *hmod the handle of the module that function hfunc is located in. The lifetime of the
module corresponds to the lifetime of the context it was loaded in or until the module is explicitly
unloaded.
The CUDA runtime manages its own modules loaded into the primary context. If the handle returned
by this API refers to a module loaded by the CUDA runtime, calling cuModuleUnload() on that module
will result in undefined behavior.
|
388
Modules
Note:
Note that this function may also return error codes from previous, asynchronous launches.
CUresult cuFuncGetName (const char **name, CUfunction
hfunc)
Returns the function name for a CUfunction handle.
Parameters
name
- The returned name of the function
hfunc
- The function handle to retrieve the name for
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE,
Description
Returns in **name the function name associated with the function handle hfunc . The function name
is returned as a null-terminated string. The returned name is only valid when the function handle is
valid. If the module is unloaded or reloaded, one must call the API again to get the updated name. This
API may return a mangled name if the function is not declared as having C linkage. If either **name
or hfunc is NULL, CUDA_ERROR_INVALID_VALUE is returned.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
CUresult cuFuncGetParamInfo (CUfunction func, size_t
paramIndex, size_t *paramOffset, size_t *paramSize)
Returns the offset and size of a kernel parameter in the device-side parameter layout.
Parameters
func
- The function to query
paramIndex
- The parameter index to query
|
389
Modules
paramOffset
- Returns the offset into the device-side parameter layout at which the parameter resides
paramSize
- Optionally returns the size of the parameter in the device-side parameter layout
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE,
Description
Queries the kernel parameter at paramIndex into func's list of parameters, and returns in
paramOffset and paramSize the offset and size, respectively, where the parameter will reside
in the device-side parameter layout. This information can be used to update kernel node parameters
from the device via cudaGraphKernelNodeSetParam() and cudaGraphKernelNodeUpdatesApply().
paramIndex must be less than the number of parameters that func takes. paramSize can be set to
NULL if only the parameter offset is desired.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuKernelGetParamInfo
CUresult cuFuncIsLoaded (CUfunctionLoadingState *state,
CUfunction function)
Returns if the function is loaded.
Parameters
state
- returned loading state
function
- the function to check
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_HANDLE, CUDA_ERROR_INVALID_VALUE
Description
Returns in state the loading state of function.
|
390
Modules
See also:
cuFuncLoad, cuModuleEnumerateFunctions
CUresult cuFuncLoad (CUfunction function)
Loads a function.
Parameters
function
- the function to load
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_HANDLE, CUDA_ERROR_INVALID_VALUE
Description
Finalizes function loading for function. Calling this API with a fully loaded function has no effect.
See also:
cuModuleEnumerateFunctions, cuFuncIsLoaded
CUresult cuFuncSetAttribute (CUfunction hfunc,
CUfunction_attribute attrib, int value)
Sets information about a function.
Parameters
hfunc
- Function to query attribute of
attrib
- Attribute requested
value
- The value to set
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE,
CUDA_ERROR_INVALID_VALUE
Description
This call sets the value of a specified attribute attrib on the kernel given by hfunc to an
integer value specified by val This function returns CUDA_SUCCESS if the new value of the
|
391
Modules
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
(CUDA_ERROR_INVALID_VALUE)
Supported attributes for the cuFuncSetAttribute call are:
‣ CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES: This maximum
size in bytes of dynamically-allocated shared memory. The value should contain the requested
maximum size of dynamically-allocated shared memory. The sum of this value and the function
attribute CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES cannot exceed the device attribute
CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK_OPTIN. The maximal
size of requestable dynamic shared memory may differ by GPU architecture.
‣ CU_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT: 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
CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_MULTIPROCESSOR This is
only a hint, and the driver can choose a different ratio if required to execute the function.
‣ CU_FUNC_ATTRIBUTE_REQUIRED_CLUSTER_WIDTH: 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 CUDA_ERROR_NOT_PERMITTED.
‣ CU_FUNC_ATTRIBUTE_REQUIRED_CLUSTER_HEIGHT: 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 CUDA_ERROR_NOT_PERMITTED.
‣ CU_FUNC_ATTRIBUTE_REQUIRED_CLUSTER_DEPTH: 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 CUDA_ERROR_NOT_PERMITTED.
‣ CU_FUNC_ATTRIBUTE_NON_PORTABLE_CLUSTER_SIZE_ALLOWED: Indicates whether
the function can be launched with non-portable cluster size. 1 is allowed, 0 is disallowed.
‣ CU_FUNC_ATTRIBUTE_CLUSTER_SCHEDULING_POLICY_PREFERENCE: The block
scheduling policy of a function. The value type is CUclusterSchedulingPolicy.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuCtxGetCacheConfig, cuCtxSetCacheConfig, cuFuncSetCacheConfig, cuLaunchKernel,
cudaFuncGetAttributes, cudaFuncSetAttribute, cuKernelSetAttribute
|
392
Modules
CUresult cuFuncSetCacheConfig (CUfunction hfunc,
CUfunc_cache config)
Sets the preferred cache configuration for a device function.
Parameters
hfunc
- Kernel to configure cache for
config
- Requested cache configuration
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_DEINITIALIZED,
CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_INVALID_CONTEXT
Description
On devices where the L1 cache and shared memory use the same hardware resources, this sets
through config the preferred cache configuration for the device function hfunc. This is only
a preference. The driver will use the requested configuration if possible, but it is free to choose
a different configuration if required to execute hfunc. Any context-wide preference set via
cuCtxSetCacheConfig() will be overridden by this per-function setting unless the per-function setting is
CU_FUNC_CACHE_PREFER_NONE. In that case, the current context-wide setting will be used.
This setting does nothing on devices where the size of the L1 cache and shared memory are fixed.
Launching a kernel with a different preference than the most recent preference setting may insert a
device-side synchronization point.
The supported cache configurations are:
‣ CU_FUNC_CACHE_PREFER_NONE: no preference for shared memory or L1 (default)
‣ CU_FUNC_CACHE_PREFER_SHARED: prefer larger shared memory and smaller L1 cache
‣ CU_FUNC_CACHE_PREFER_L1: prefer larger L1 cache and smaller shared memory
‣ CU_FUNC_CACHE_PREFER_EQUAL: prefer equal sized L1 cache and shared memory
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuCtxGetCacheConfig, cuCtxSetCacheConfig, cuFuncGetAttribute, cuLaunchKernel,
cudaFuncSetCacheConfig, cuKernelSetCacheConfig
|
393
Modules
CUresult cuLaunchCooperativeKernel (CUfunction
f, unsigned int gridDimX, unsigned int gridDimY,
unsigned int gridDimZ, unsigned int blockDimX,
unsigned int blockDimY, unsigned int blockDimZ,
unsigned int sharedMemBytes, CUstream hStream, void
**kernelParams)
Launches a CUDA function CUfunction or a CUDA kernel CUkernel where thread blocks can
cooperate and synchronize as they execute.
Parameters
f
- Function CUfunction or Kernel CUkernel to launch
gridDimX
- Width of grid in blocks
gridDimY
- Height of grid in blocks
gridDimZ
- Depth of grid in blocks
blockDimX
- X dimension of each thread block
blockDimY
- Y dimension of each thread block
blockDimZ
- Z dimension of each thread block
sharedMemBytes
- Dynamic shared-memory size per thread block in bytes
hStream
- Stream identifier
kernelParams
- Array of pointers to kernel parameters
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE,
CUDA_ERROR_INVALID_IMAGE, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_LAUNCH_FAILED, CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES,
CUDA_ERROR_LAUNCH_TIMEOUT,
CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING,
|
394
Modules
CUDA_ERROR_COOPERATIVE_LAUNCH_TOO_LARGE,
CUDA_ERROR_SHARED_OBJECT_INIT_FAILED, CUDA_ERROR_NOT_FOUND
Description
Invokes the function CUfunction or the kernel CUkernel f on a gridDimX x gridDimY x
gridDimZ grid of blocks. Each block contains blockDimX x blockDimY x blockDimZ threads.
sharedMemBytes sets the amount of dynamic shared memory that will be available to each thread
block.
The device on which this kernel is invoked must have a non-zero value for the device attribute
CU_DEVICE_ATTRIBUTE_COOPERATIVE_LAUNCH.
The total number of blocks launched cannot exceed the maximum number of blocks
per multiprocessor as returned by cuOccupancyMaxActiveBlocksPerMultiprocessor (or
cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags) times the number of multiprocessors as
specified by the device attribute CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT.
The kernel cannot make use of CUDA dynamic parallelism.
Kernel parameters must be specified via kernelParams. If f has N parameters, then
kernelParams needs to be an array of N pointers. Each of kernelParams[0] through
kernelParams[N-1] must point to a region of memory from which the actual kernel parameter will
be copied. The number of kernel parameters and their offsets and sizes do not need to be specified as
that information is retrieved directly from the kernel's image.
Calling cuLaunchCooperativeKernel() sets persistent function state that is the same as function state set
through cuLaunchKernel API
When the kernel f is launched via cuLaunchCooperativeKernel(), the previous block shape, shared size
and parameter info associated with f is overwritten.
Note that to use cuLaunchCooperativeKernel(), the kernel f must either have been compiled with
toolchain version 3.2 or later so that it will contain kernel parameter information, or have no kernel
parameters. If either of these conditions is not met, then cuLaunchCooperativeKernel() will return
CUDA_ERROR_INVALID_IMAGE.
Note that the API can also be used to launch context-less kernel CUkernel by querying the handle using
cuLibraryGetKernel() and then passing it to the API by casting to CUfunction. Here, the context to
launch the kernel on will either be taken from the specified stream hStream or the current context in
case of NULL stream.
Note:
‣ This function uses standard default stream semantics.
‣ Note that this function may also return error codes from previous, asynchronous launches.
|
395
Modules
See also:
cuCtxGetCacheConfig, cuCtxSetCacheConfig, cuFuncSetCacheConfig, cuFuncGetAttribute,
cuLaunchCooperativeKernelMultiDevice, cudaLaunchCooperativeKernel, cuLibraryGetKernel,
cuKernelSetCacheConfig, cuKernelGetAttribute, cuKernelSetAttribute
CUresult cuLaunchCooperativeKernelMultiDevice
(CUDA_LAUNCH_PARAMS *launchParamsList,
unsigned int numDevices, unsigned int flags)
Launches CUDA functions on multiple devices where thread blocks can cooperate and synchronize as
they execute.
Parameters
launchParamsList
- List of launch parameters, one per device
numDevices
- Size of the launchParamsList array
flags
- Flags to control launch behavior
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE,
CUDA_ERROR_INVALID_IMAGE, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_LAUNCH_FAILED, CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES,
CUDA_ERROR_LAUNCH_TIMEOUT,
CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING,
CUDA_ERROR_COOPERATIVE_LAUNCH_TOO_LARGE,
CUDA_ERROR_SHARED_OBJECT_INIT_FAILED
Description
Deprecated This function is deprecated as of CUDA 11.3.
Invokes kernels as specified in the launchParamsList array where each element of the array
specifies all the parameters required to perform a single kernel launch. These kernels can cooperate and
synchronize as they execute. The size of the array is specified by numDevices.
No two kernels can be launched on the same device. All the devices targeted by this multi-
device launch must be identical. All devices must have a non-zero value for the device attribute
CU_DEVICE_ATTRIBUTE_COOPERATIVE_MULTI_DEVICE_LAUNCH.
All kernels launched must be identical with respect to the compiled code. Note that any __device__,
__constant__ or __managed__ variables present in the module that owns the kernel launched on each
|
396
Modules
device, are independently instantiated on every device. It is the application's responsibility to ensure
these variables are initialized and used appropriately.
The size of the grids as specified in blocks, the size of the blocks themselves and the amount of shared
memory used by each thread block must also match across all launched kernels.
The streams used to launch these kernels must have been created via either cuStreamCreate
or cuStreamCreateWithPriority. The NULL stream or CU_STREAM_LEGACY or
CU_STREAM_PER_THREAD cannot be used.
The total number of blocks launched per kernel cannot exceed the maximum number of blocks
per multiprocessor as returned by cuOccupancyMaxActiveBlocksPerMultiprocessor (or
cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags) times the number of multiprocessors
as specified by the device attribute CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT.
Since the total number of blocks launched per device has to match across all devices, the maximum
number of blocks that can be launched per device will be limited by the device with the least number of
multiprocessors.
The kernels cannot make use of CUDA dynamic parallelism.
The CUDA_LAUNCH_PARAMS structure is defined as:
‎
typedef struct CUDA_LAUNCH_PARAMS_st
{
CUfunction function;
unsigned int gridDimX;
unsigned int gridDimY;
unsigned int gridDimZ;
unsigned int blockDimX;
unsigned int blockDimY;
unsigned int blockDimZ;
unsigned int sharedMemBytes;
CUstream hStream;
void **kernelParams;
} CUDA_LAUNCH_PARAMS;
where:
‣ CUDA_LAUNCH_PARAMS::function specifies the kernel to be launched. All functions
must be identical with respect to the compiled code. Note that you can also specify context-
less kernel CUkernel by querying the handle using cuLibraryGetKernel() and then casting to
CUfunction. In this case, the context to launch the kernel on be taken from the specified stream
CUDA_LAUNCH_PARAMS::hStream.
‣ CUDA_LAUNCH_PARAMS::gridDimX is the width of the grid in blocks. This must match across
all kernels launched.
‣ CUDA_LAUNCH_PARAMS::gridDimY is the height of the grid in blocks. This must match
across all kernels launched.
‣ CUDA_LAUNCH_PARAMS::gridDimZ is the depth of the grid in blocks. This must match across
all kernels launched.
‣ CUDA_LAUNCH_PARAMS::blockDimX is the X dimension of each thread block. This must
match across all kernels launched.
|
397
Modules
‣ CUDA_LAUNCH_PARAMS::blockDimX is the Y dimension of each thread block. This must
match across all kernels launched.
‣ CUDA_LAUNCH_PARAMS::blockDimZ is the Z dimension of each thread block. This must
match across all kernels launched.
‣ CUDA_LAUNCH_PARAMS::sharedMemBytes is the dynamic shared-memory size per thread
block in bytes. This must match across all kernels launched.
‣ CUDA_LAUNCH_PARAMS::hStream is the handle to the stream to perform the launch in. This
cannot be the NULL stream or CU_STREAM_LEGACY or CU_STREAM_PER_THREAD.
The CUDA context associated with this stream must match that associated with
CUDA_LAUNCH_PARAMS::function.
‣ CUDA_LAUNCH_PARAMS::kernelParams is an array of pointers to kernel
parameters. If CUDA_LAUNCH_PARAMS::function has N parameters,
then CUDA_LAUNCH_PARAMS::kernelParams needs to be an array of N
pointers. Each of CUDA_LAUNCH_PARAMS::kernelParams[0] through
CUDA_LAUNCH_PARAMS::kernelParams[N-1] must point to a region of memory from which
the actual kernel parameter will be copied. The number of kernel parameters and their offsets and
sizes do not need to be specified as that information is retrieved directly from the kernel's image.
By default, the kernel won't begin execution on any GPU until all prior work in all the
specified streams has completed. This behavior can be overridden by specifying the flag
CUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_PRE_LAUNCH_SYNC. When this
flag is specified, each kernel will only wait for prior work in the stream corresponding to that GPU to
complete before it begins execution.
Similarly, by default, any subsequent work pushed in any of the specified streams will not begin
execution until the kernels on all GPUs have completed. This behavior can be overridden by specifying
the flag CUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_POST_LAUNCH_SYNC.
When this flag is specified, any subsequent work pushed in any of the specified streams will only
wait for the kernel launched on the GPU corresponding to that stream to complete before it begins
execution.
Calling cuLaunchCooperativeKernelMultiDevice() sets persistent function state that is the same
as function state set through cuLaunchKernel API when called individually for each element in
launchParamsList.
When kernels are launched via cuLaunchCooperativeKernelMultiDevice(), the previous block shape,
shared size and parameter info associated with each CUDA_LAUNCH_PARAMS::function in
launchParamsList is overwritten.
Note that to use cuLaunchCooperativeKernelMultiDevice(), the kernels must either have
been compiled with toolchain version 3.2 or later so that it will contain kernel parameter
information, or have no kernel parameters. If either of these conditions is not met, then
cuLaunchCooperativeKernelMultiDevice() will return CUDA_ERROR_INVALID_IMAGE.
|
398
Modules
Note:
‣ This function uses standard default stream semantics.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuCtxGetCacheConfig, cuCtxSetCacheConfig, cuFuncSetCacheConfig, cuFuncGetAttribute,
cuLaunchCooperativeKernel, cudaLaunchCooperativeKernelMultiDevice
CUresult cuLaunchHostFunc (CUstream hStream,
CUhostFn fn, void *userData)
Enqueues a host function call in a stream.
Parameters
hStream
- Stream to enqueue function call in
fn
- The function to call once preceding stream operations are complete
userData
- User-specified data to be passed to the function
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE,
CUDA_ERROR_NOT_SUPPORTED
Description
Enqueues a host function to run in a stream. The function will be called after currently enqueued work
and will block work added after it.
The host function must not make any CUDA API calls. Attempting to use a CUDA API may result
in CUDA_ERROR_NOT_PERMITTED, but this is not required. The host function must not perform
any synchronization that may depend on outstanding CUDA work not mandated to run earlier. Host
functions without a mandated order (such as in independent streams) execute in undefined order and
may be serialized.
For the purposes of Unified Memory, execution makes a number of guarantees:
‣ The stream is considered idle for the duration of the function's execution. Thus, for example, the
function may always use memory attached to the stream it was enqueued in.
|
399
Modules
‣ The start of execution of the function has the same effect as synchronizing an event recorded in
the same stream immediately prior to the function. It thus synchronizes streams which have been
"joined" prior to the function.
‣ Adding device work to any stream does not have the effect of making the stream active until all
preceding host functions and stream callbacks have executed. Thus, for example, a function might
use global attached memory even if work has been added to another stream, if the work has been
ordered behind the function call with an event.
‣ Completion of the function does not cause a stream to become active except as described above.
The stream will remain idle if no device work follows the function, and will remain idle across
consecutive host functions or stream callbacks without device work in between. Thus, for example,
stream synchronization can be done by signaling from a host function at the end of the stream.
Note that, in contrast to cuStreamAddCallback, the function will not be called in the event of an error
in the CUDA context.
Note:
‣ This function uses standard default stream semantics.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuStreamCreate, cuStreamQuery, cuStreamSynchronize, cuStreamWaitEvent, cuStreamDestroy,
cuMemAllocManaged, cuStreamAttachMemAsync, cuStreamAddCallback
CUresult cuLaunchKernel (CUfunction f, unsigned int
gridDimX, unsigned int gridDimY, unsigned int gridDimZ,
unsigned int blockDimX, unsigned int blockDimY,
unsigned int blockDimZ, unsigned int sharedMemBytes,
CUstream hStream, void **kernelParams, void **extra)
Launches a CUDA function CUfunction or a CUDA kernel CUkernel.
Parameters
f
- Function CUfunction or Kernel CUkernel to launch
gridDimX
- Width of grid in blocks
gridDimY
- Height of grid in blocks
|
400
Modules
gridDimZ
- Depth of grid in blocks
blockDimX
- X dimension of each thread block
blockDimY
- Y dimension of each thread block
blockDimZ
- Z dimension of each thread block
sharedMemBytes
- Dynamic shared-memory size per thread block in bytes
hStream
- Stream identifier
kernelParams
- Array of pointers to kernel parameters
extra
- Extra options
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE,
CUDA_ERROR_INVALID_IMAGE, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_LAUNCH_FAILED, CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES,
CUDA_ERROR_LAUNCH_TIMEOUT,
CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING,
CUDA_ERROR_SHARED_OBJECT_INIT_FAILED, CUDA_ERROR_NOT_FOUND
Description
Invokes the function CUfunction or the kernel CUkernel f on a gridDimX x gridDimY x
gridDimZ grid of blocks. Each block contains blockDimX x blockDimY x blockDimZ threads.
sharedMemBytes sets the amount of dynamic shared memory that will be available to each thread
block.
Kernel parameters to f can be specified in one of two ways:
1) Kernel parameters can be specified via kernelParams. If f has N parameters, then
kernelParams needs to be an array of N pointers. Each of kernelParams[0] through
kernelParams[N-1] must point to a region of memory from which the actual kernel parameter will
be copied. The number of kernel parameters and their offsets and sizes do not need to be specified as
that information is retrieved directly from the kernel's image.
2) Kernel parameters can also be packaged by the application into a single buffer that is passed in via
the extra parameter. This places the burden on the application of knowing each kernel parameter's
|
401
Modules
size and alignment/padding within the buffer. Here is an example of using the extra parameter in this
manner:
‎
size_t argBufferSize;
char argBuffer[256];
// populate argBuffer and argBufferSize
void *config[] = {
CU_LAUNCH_PARAM_BUFFER_POINTER, argBuffer,
CU_LAUNCH_PARAM_BUFFER_SIZE,
&argBufferSize,
CU_LAUNCH_PARAM_END
};
status = cuLaunchKernel(f, gx, gy, gz, bx, by, bz, sh, s, NULL, config);
The extra parameter exists to allow cuLaunchKernel to take additional less commonly used
arguments. extra specifies a list of names of extra settings and their corresponding values. Each extra
setting name is immediately followed by the corresponding value. The list must be terminated with
either NULL or CU_LAUNCH_PARAM_END.
‣ CU_LAUNCH_PARAM_END, which indicates the end of the extra array;
‣ CU_LAUNCH_PARAM_BUFFER_POINTER, which specifies that the next value in extra will
be a pointer to a buffer containing all the kernel parameters for launching kernel f;
‣ CU_LAUNCH_PARAM_BUFFER_SIZE, which specifies that the next value in
extra will be a pointer to a size_t containing the size of the buffer specified with
CU_LAUNCH_PARAM_BUFFER_POINTER;
The error CUDA_ERROR_INVALID_VALUE will be returned if kernel parameters are specified with
both kernelParams and extra (i.e. both kernelParams and extra are non-NULL).
Calling cuLaunchKernel() invalidates the persistent function state set through the following
deprecated APIs: cuFuncSetBlockShape(), cuFuncSetSharedSize(), cuParamSetSize(), cuParamSeti(),
cuParamSetf(), cuParamSetv().
Note that to use cuLaunchKernel(), the kernel f must either have been compiled with
toolchain version 3.2 or later so that it will contain kernel parameter information, or have no
kernel parameters. If either of these conditions is not met, then cuLaunchKernel() will return
CUDA_ERROR_INVALID_IMAGE.
Note that the API can also be used to launch context-less kernel CUkernel by querying the handle using
cuLibraryGetKernel() and then passing it to the API by casting to CUfunction. Here, the context to
launch the kernel on will either be taken from the specified stream hStream or the current context in
case of NULL stream.
Note:
‣ This function uses standard default stream semantics.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
|
402
Modules
cuCtxGetCacheConfig, cuCtxSetCacheConfig, cuFuncSetCacheConfig, cuFuncGetAttribute,
cudaLaunchKernel, cuLibraryGetKernel, cuKernelSetCacheConfig, cuKernelGetAttribute,
cuKernelSetAttribute
CUresult cuLaunchKernelEx (const CUlaunchConfig
*config, CUfunction f, void **kernelParams, void **extra)
Launches a CUDA function CUfunction or a CUDA kernel CUkernel with launch-time configuration.
Parameters
config
- Config to launch
f
- Function CUfunction or Kernel CUkernel to launch
kernelParams
- Array of pointers to kernel parameters
extra
- Extra options
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE,
CUDA_ERROR_INVALID_IMAGE, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_LAUNCH_FAILED, CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES,
CUDA_ERROR_LAUNCH_TIMEOUT,
CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING,
CUDA_ERROR_COOPERATIVE_LAUNCH_TOO_LARGE,
CUDA_ERROR_SHARED_OBJECT_INIT_FAILED, CUDA_ERROR_NOT_FOUND
Description
Invokes the function CUfunction or the kernel CUkernel f with the specified launch-time configuration
config.
The CUlaunchConfig structure is defined as:
‎
typedef struct CUlaunchConfig_st {
unsigned int gridDimX;
unsigned int gridDimY;
unsigned int gridDimZ;
unsigned int blockDimX;
unsigned int blockDimY;
unsigned int blockDimZ;
unsigned int sharedMemBytes;
CUstream hStream;
CUlaunchAttribute *attrs;
unsigned int numAttrs;
} CUlaunchConfig;
|
403
Modules
where:
‣ CUlaunchConfig::gridDimX is the width of the grid in blocks.
‣ CUlaunchConfig::gridDimY is the height of the grid in blocks.
‣ CUlaunchConfig::gridDimZ is the depth of the grid in blocks.
‣ CUlaunchConfig::blockDimX is the X dimension of each thread block.
‣ CUlaunchConfig::blockDimX is the Y dimension of each thread block.
‣ CUlaunchConfig::blockDimZ is the Z dimension of each thread block.
‣ CUlaunchConfig::sharedMemBytes is the dynamic shared-memory size per thread block in bytes.
‣ CUlaunchConfig::hStream is the handle to the stream to perform the launch in. The CUDA context
associated with this stream must match that associated with function f.
‣ CUlaunchConfig::attrs is an array of CUlaunchConfig::numAttrs continguous CUlaunchAttribute
elements. The value of this pointer is not considered if CUlaunchConfig::numAttrs is zero.
However, in that case, it is recommended to set the pointer to NULL.
‣ CUlaunchConfig::numAttrs is the number of attributes populating the first
CUlaunchConfig::numAttrs positions of the CUlaunchConfig::attrs array.
Launch-time configuration is specified by adding entries to CUlaunchConfig::attrs. Each entry is an
attribute ID and a corresponding attribute value.
The CUlaunchAttribute structure is defined as:
‎
typedef struct CUlaunchAttribute_st {
CUlaunchAttributeID id;
CUlaunchAttributeValue value;
} CUlaunchAttribute;
where:
‣ CUlaunchAttribute::id is a unique enum identifying the attribute.
‣ CUlaunchAttribute::value is a union that hold the attribute value.
An example of using the config parameter:
‎
CUlaunchAttribute coopAttr = {.id = CU_LAUNCH_ATTRIBUTE_COOPERATIVE,
.value = 1};
CUlaunchConfig config = {... // set block and grid dimensions
.attrs = &coopAttr,
.numAttrs = 1};
cuLaunchKernelEx(&config, kernel, NULL, NULL);
The CUlaunchAttributeID enum is defined as:
‎
typedef enum CUlaunchAttributeID_enum {
CU_LAUNCH_ATTRIBUTE_IGNORE = 0,
CU_LAUNCH_ATTRIBUTE_ACCESS_POLICY_WINDOW
= 1,
CU_LAUNCH_ATTRIBUTE_COOPERATIVE
= 2,
CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY = 3,
CU_LAUNCH_ATTRIBUTE_CLUSTER_DIMENSION
= 4,
CU_LAUNCH_ATTRIBUTE_CLUSTER_SCHEDULING_POLICY_PREFERENCE = 5,
CU_LAUNCH_ATTRIBUTE_PROGRAMMATIC_STREAM_SERIALIZATION
= 6,
CU_LAUNCH_ATTRIBUTE_PROGRAMMATIC_EVENT
= 7,
CU_LAUNCH_ATTRIBUTE_PRIORITY
= 8,
CU_LAUNCH_ATTRIBUTE_MEM_SYNC_DOMAIN_MAP
= 9,
CU_LAUNCH_ATTRIBUTE_MEM_SYNC_DOMAIN
= 10,
CU_LAUNCH_ATTRIBUTE_PREFERRED_CLUSTER_DIMENSION = 11,
|
404
Modules
CU_LAUNCH_ATTRIBUTE_LAUNCH_COMPLETION_EVENT = 12,
CU_LAUNCH_ATTRIBUTE_DEVICE_UPDATABLE_KERNEL_NODE = 13,
} CUlaunchAttributeID;
and the corresponding CUlaunchAttributeValue union as :
‎
typedef union CUlaunchAttributeValue_union {
CUaccessPolicyWindow accessPolicyWindow;
int cooperative;
CUsynchronizationPolicy syncPolicy;
struct {
unsigned int x;
unsigned int y;
unsigned int z;
} clusterDim;
CUclusterSchedulingPolicy clusterSchedulingPolicyPreference;
int programmaticStreamSerializationAllowed;
struct {
CUevent event;
int flags;
int triggerAtBlockStart;
} programmaticEvent;
int priority;
CUlaunchMemSyncDomainMap memSyncDomainMap;
CUlaunchMemSyncDomain memSyncDomain;
struct {
unsigned int x;
unsigned int y;
unsigned int z;
} preferredClusterDim;
struct {
CUevent event;
int flags;
} launchCompletionEvent;
struct {
int deviceUpdatable;
CUgraphDeviceNode devNode;
} deviceUpdatableKernelNode;
} CUlaunchAttributeValue;
Setting CU_LAUNCH_ATTRIBUTE_COOPERATIVE to a non-zero value causes the kernel launch
to be a cooperative launch, with exactly the same usage and semantics of cuLaunchCooperativeKernel.
Setting CU_LAUNCH_ATTRIBUTE_PROGRAMMATIC_STREAM_SERIALIZATION to a non-
zero values causes the kernel to use programmatic means to resolve its stream dependency -- enabling
the CUDA runtime to opportunistically allow the grid's execution to overlap with the previous kernel in
the stream, if that kernel requests the overlap.
CU_LAUNCH_ATTRIBUTE_PROGRAMMATIC_EVENT records an event along with the kernel
launch. Event recorded through this launch attribute is guaranteed to only trigger after all block in the
associated kernel trigger the event. A block can trigger the event through PTX launchdep.release or
CUDA builtin function cudaTriggerProgrammaticLaunchCompletion(). A trigger can also be inserted
at the beginning of each block's execution if triggerAtBlockStart is set to non-0. Note that dependents
(including the CPU thread calling cuEventSynchronize()) are not guaranteed to observe the release
precisely when it is released. For example, cuEventSynchronize() may only observe the event trigger
long after the associated kernel has completed. This recording type is primarily meant for establishing
programmatic dependency between device tasks. The event supplied must not be an interprocess or
interop event. The event must disable timing (i.e. created with CU_EVENT_DISABLE_TIMING flag
set).
|
405
Modules
CU_LAUNCH_ATTRIBUTE_LAUNCH_COMPLETION_EVENT records an event along with the
kernel launch. Nominally, the event is triggered once all blocks of the kernel have begun execution.
Currently this is a best effort. If a kernel B has a launch completion dependency on a kernel A, B may
wait until A is complete. Alternatively, blocks of B may begin before all blocks of A have begun, for
example:
‣ If B can claim execution resources unavailable to A, for example if they run on different GPUs.
‣ If B is a higher priority than A.
Exercise caution if such an ordering inversion could lead to deadlock. The event supplied must not
be an interprocess or interop event. The event must disable timing (i.e. must be created with the
CU_EVENT_DISABLE_TIMING flag set).
Setting CU_LAUNCH_ATTRIBUTE_DEVICE_UPDATABLE_KERNEL_NODE to 1 on a captured
launch causes the resulting kernel node to be device-updatable. This attribute is specific to graphs, and
passing it to a launch in a non-capturing stream results in an error. Passing a value other than 0 or 1 is
not allowed.
On success, a handle will be returned via
CUlaunchAttributeValue::deviceUpdatableKernelNode::devNode which can be passed to the various
device-side update functions to update the node's kernel parameters from within another kernel. For
more information on the types of device updates that can be made, as well as the relevant limitations
thereof, see cudaGraphKernelNodeUpdatesApply.
Kernel nodes which are device-updatable have additional restrictions compared to regular kernel
nodes. Firstly, device-updatable nodes cannot be removed from their graph via cuGraphDestroyNode.
Additionally, once opted-in to this functionality, a node cannot opt out, and any attempt to set the
attribute to 0 will result in an error. Graphs containing one or more device-updatable node also do not
allow multiple instantiation.
CU_LAUNCH_ATTRIBUTE_PREFERRED_CLUSTER_DIMENSION allows the kernel launch
to specify a preferred substitute cluster dimension. Blocks may be grouped according to either the
dimensions specified with this attribute (grouped into a "preferred substitute cluster"), or the one
specified with CU_LAUNCH_ATTRIBUTE_CLUSTER_DIMENSION attribute (grouped into a
"regular cluster"). The cluster dimensions of a "preferred substitute cluster" shall be an integer multiple
greater than zero of the regular cluster dimensions. The device will attempt - on a best-effort basis -
to group thread blocks into preferred clusters over grouping them into regular clusters. When it deems
necessary (primarily when the device temporarily runs out of physical resources to launch the larger
preferred clusters), the device may switch to launch the regular clusters instead to attempt to utilize as
much of the physical device resources as possible.
Each type of cluster will have its enumeration / coordinate setup as if the grid consists solely of its
type of cluster. For example, if the preferred substitute cluster dimensions double the regular cluster
dimensions, there might be simultaneously a regular cluster indexed at (1,0,0), and a preferred cluster
indexed at (1,0,0). In this example, the preferred substitute cluster (1,0,0) replaces regular clusters
(2,0,0) and (3,0,0) and groups their blocks.
|
406
Modules
This attribute will only take effect when a regular cluster dimension has been specified. The
preferred substitute The preferred substitute cluster dimension must be an integer multiple greater
than zero of the regular cluster dimension and must divide the grid. It must also be no more than
`maxBlocksPerCluster`, if it is set in the kernel's `__launch_bounds__`. Otherwise it must be less than
the maximum value the driver can support. Otherwise, setting this attribute to a value physically unable
to fit on any particular device is permitted.
The effect of other attributes is consistent with their effect when set via persistent APIs.
See cuStreamSetAttribute for
‣ CU_LAUNCH_ATTRIBUTE_ACCESS_POLICY_WINDOW
‣ CU_LAUNCH_ATTRIBUTE_SYNCHRONIZATION_POLICY
See cuFuncSetAttribute for
‣ CU_LAUNCH_ATTRIBUTE_CLUSTER_DIMENSION
‣ CU_LAUNCH_ATTRIBUTE_CLUSTER_SCHEDULING_POLICY_PREFERENCE
Kernel parameters to f can be specified in the same ways that they can be using cuLaunchKernel.
Note that the API can also be used to launch context-less kernel CUkernel by querying the handle using
cuLibraryGetKernel() and then passing it to the API by casting to CUfunction. Here, the context to
launch the kernel on will either be taken from the specified stream CUlaunchConfig::hStream or the
current context in case of NULL stream.
Note:
‣ This function uses standard default stream semantics.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuCtxGetCacheConfig, cuCtxSetCacheConfig, cuFuncSetCacheConfig, cuFuncGetAttribute,
cudaLaunchKernel, cudaLaunchKernelEx, cuLibraryGetKernel, cuKernelSetCacheConfig,
cuKernelGetAttribute, cuKernelSetAttribute
6.23. Execution Control [DEPRECATED]
This section describes the deprecated execution control functions of the low-level CUDA driver
application programming interface.
|
407
Modules
CUresult cuFuncSetBlockShape (CUfunction hfunc, int x,
int y, int z)
Sets the block-dimensions for the function.
Parameters
hfunc
- Kernel to specify dimensions of
x
- X dimension
y
- Y dimension
z
- Z dimension
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE,
CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Specifies the x, y, and z dimensions of the thread blocks that are created when the kernel given by
hfunc is launched.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuFuncSetSharedSize, cuFuncSetCacheConfig, cuFuncGetAttribute, cuParamSetSize, cuParamSeti,
cuParamSetf, cuParamSetv, cuLaunch, cuLaunchGrid, cuLaunchGridAsync, cuLaunchKernel
|
408
Modules
CUresult cuFuncSetSharedMemConfig (CUfunction hfunc,
CUsharedconfig config)
Sets the shared memory configuration for a device function.
Parameters
hfunc
- kernel to be given a shared memory config
config
- requested shared memory configuration
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_DEINITIALIZED,
CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_INVALID_CONTEXT
Description
Deprecated
On devices with configurable shared memory banks, this function will force all subsequent launches of
the specified device function to have the given shared memory bank size configuration. On any given
launch of the function, the shared memory configuration of the device will be temporarily changed if
needed to suit the function's preferred configuration. Changes in shared memory configuration between
subsequent launches of functions, may introduce a device side synchronization point.
Any per-function setting of shared memory bank size set via cuFuncSetSharedMemConfig will
override the context wide setting set with cuCtxSetSharedMemConfig.
Changing the shared memory bank size will not increase shared memory usage or affect occupancy of
kernels, but may have major effects on performance. Larger bank sizes will allow for greater potential
bandwidth to shared memory, but will change what kinds of accesses to shared memory will result in
bank conflicts.
This function will do nothing on devices with fixed shared memory bank size.
The supported bank configurations are:
‣ CU_SHARED_MEM_CONFIG_DEFAULT_BANK_SIZE: use the context's shared memory
configuration when launching this function.
‣ CU_SHARED_MEM_CONFIG_FOUR_BYTE_BANK_SIZE: set shared memory bank width to
be natively four bytes when launching this function.
‣ CU_SHARED_MEM_CONFIG_EIGHT_BYTE_BANK_SIZE: set shared memory bank width to
be natively eight bytes when launching this function.
Note:
|
409
Modules
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuCtxGetCacheConfig, cuCtxSetCacheConfig, cuCtxGetSharedMemConfig,
cuCtxSetSharedMemConfig, cuFuncGetAttribute, cuLaunchKernel, cudaFuncSetSharedMemConfig
CUresult cuFuncSetSharedSize (CUfunction hfunc,
unsigned int bytes)
Sets the dynamic shared-memory size for the function.
Parameters
hfunc
- Kernel to specify dynamic shared-memory size for
bytes
- Dynamic shared-memory size per thread in bytes
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE,
CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Sets through bytes the amount of dynamic shared memory that will be available to each thread block
when the kernel given by hfunc is launched.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuFuncSetBlockShape, cuFuncSetCacheConfig, cuFuncGetAttribute, cuParamSetSize, cuParamSeti,
cuParamSetf, cuParamSetv, cuLaunch, cuLaunchGrid, cuLaunchGridAsync, cuLaunchKernel
|
410
Modules
CUresult cuLaunch (CUfunction f)
Launches a CUDA function.
Parameters
f
- Kernel to launch
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_LAUNCH_FAILED, CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES,
CUDA_ERROR_LAUNCH_TIMEOUT,
CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING,
CUDA_ERROR_SHARED_OBJECT_INIT_FAILED
Description
Deprecated
Invokes the kernel f on a 1 x 1 x 1 grid of blocks. The block contains the number of threads specified
by a previous call to cuFuncSetBlockShape().
The block shape, dynamic shared memory size, and parameter information must be set using
cuFuncSetBlockShape(), cuFuncSetSharedSize(), cuParamSetSize(), cuParamSeti(), cuParamSetf(),
and cuParamSetv() prior to calling this function.
Launching a function via cuLaunchKernel() invalidates the function's block shape, dynamic shared
memory size, and parameter information. After launching via cuLaunchKernel, this state must be re-
initialized prior to calling this function. Failure to do so results in undefined behavior.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuFuncSetBlockShape, cuFuncSetSharedSize, cuFuncGetAttribute, cuParamSetSize, cuParamSetf,
cuParamSeti, cuParamSetv, cuLaunchGrid, cuLaunchGridAsync, cuLaunchKernel
|
411

 

 

 

 

 

 

 

Content      ..     4      5      6      7     ..