|
|
Modules
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaStreamCreate, cudaStreamCreateWithFlags, cudaStreamQuery, cudaStreamWaitEvent,
cudaStreamAddCallback, cudaStreamDestroy, cuStreamSynchronize
__host__cudaError_t
cudaStreamUpdateCaptureDependencies (cudaStream_t
stream, cudaGraphNode_t *dependencies, size_t
numDependencies, unsigned int flags)
Update the set of dependencies in a capturing stream (11.3+).
Parameters
stream
- The stream to update
dependencies
- The set of dependencies to add
numDependencies
- The size of the dependencies array
flags
- See above
Returns
cudaSuccess, cudaErrorInvalidValue, cudaErrorIllegalState
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 cudaStreamAddCaptureDependencies and cudaStreamSetCaptureDependencies. These
control whether the set passed to the API is added to the existing set or replaces it. A flags value of 0
defaults to cudaStreamAddCaptureDependencies.
Nodes that are removed from the dependency set via this API do not result in
cudaErrorStreamCaptureUnjoined if they are unreachable from the stream at cudaStreamEndCapture.
Returns cudaErrorIllegalState if the stream is not capturing.
This API is new in CUDA 11.3. Developers requiring compatibility across minor versions of the
CUDA driver to 11.0 should not use this API or provide a fallback.
|
43
Modules
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cudaStreamBeginCapture, cudaStreamGetCaptureInfo,
__host__cudaError_t
cudaStreamUpdateCaptureDependencies_v2 (cudaStream_t
stream, cudaGraphNode_t *dependencies, const
cudaGraphEdgeData *dependencyData, size_t
numDependencies, unsigned int flags)
Update the set of dependencies in a capturing stream (12.3+).
Parameters
stream
- 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
cudaSuccess, cudaErrorInvalidValue, cudaErrorIllegalState
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 cudaStreamAddCaptureDependencies and cudaStreamSetCaptureDependencies. These
control whether the set passed to the API is added to the existing set or replaces it. A flags value of 0
defaults to cudaStreamAddCaptureDependencies.
Nodes that are removed from the dependency set via this API do not result in
cudaErrorStreamCaptureUnjoined if they are unreachable from the stream at cudaStreamEndCapture.
|
44
Modules
Returns cudaErrorIllegalState if the stream is not capturing.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cudaStreamBeginCapture, cudaStreamGetCaptureInfo,
__host____
device__cudaError_t cudaStreamWaitEvent
(cudaStream_t stream, cudaEvent_t event, unsigned int
flags)
Make a compute stream wait on an event.
Parameters
stream
- Stream to wait
event
- Event to wait on
flags
- Parameters for the operation(See above)
Returns
cudaSuccess, cudaErrorInvalidValue, cudaErrorInvalidResourceHandle
Description
Makes all future work submitted to stream wait for all work captured in event. See
cudaEventRecord() for details on what is captured by an event. The synchronization will be performed
efficiently on the device when applicable. event may be from a different device than stream.
flags include:
‣ cudaEventWaitDefault: Default event creation flag.
‣ cudaEventWaitExternal: Event is captured in the graph as an external event node when performing
stream capture.
Note:
‣ This function uses standard default stream semantics.
‣ Note that this function may also return error codes from previous, asynchronous launches.
|
45
Modules
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaStreamCreate, cudaStreamCreateWithFlags, cudaStreamQuery, cudaStreamSynchronize,
cudaStreamAddCallback, cudaStreamDestroy, cuStreamWaitEvent
__host__cudaError_t
cudaThreadExchangeStreamCaptureMode
(cudaStreamCaptureMode *mode)
Swaps the stream capture interaction mode for a thread.
Parameters
mode
- Pointer to mode value to swap with the current mode
Returns
cudaSuccess, cudaErrorInvalidValue
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:
cudaStreamCaptureMode mode = desiredMode;
cudaThreadExchangeStreamCaptureMode(&mode);
cudaThreadExchangeStreamCaptureMode(&mode); // restore previous mode
During stream capture (see cudaStreamBeginCapture), 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
cudaStreamBeginCapture 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
cudaStreamBeginCapture-cudaStreamEndCapture sequence. This behavior can be controlled via this
API and flags to cudaStreamBeginCapture.
A thread's mode is one of the following:
‣ cudaStreamCaptureModeGlobal: This is the default mode. If the local thread has an
ongoing capture sequence that was not initiated with cudaStreamCaptureModeRelaxed at
|
46
Modules
cuStreamBeginCapture, or if any other thread has a concurrent capture sequence initiated
with cudaStreamCaptureModeGlobal, this thread is prohibited from potentially unsafe API
calls.
‣ cudaStreamCaptureModeThreadLocal: If the local thread has an ongoing capture
sequence not initiated with cudaStreamCaptureModeRelaxed, it is prohibited from
potentially unsafe API calls. Concurrent capture sequences in other threads are ignored.
‣ cudaStreamCaptureModeRelaxed: 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 cudaEventQuery 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:
cudaStreamBeginCapture
6.6.
Event Management
This section describes the event management functions of the CUDA runtime application programming
interface.
__host__cudaError_t cudaEventCreate (cudaEvent_t
*event)
Creates an event object.
Parameters
event
- Newly created event
Returns
cudaSuccess, cudaErrorInvalidValue, cudaErrorLaunchFailure, cudaErrorMemoryAllocation
Description
Creates an event object for the current device using cudaEventDefault.
|
47
Modules
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaEventCreate ( C++ API), cudaEventCreateWithFlags, cudaEventRecord, cudaEventQuery,
cudaEventSynchronize, cudaEventDestroy, cudaEventElapsedTime, cudaStreamWaitEvent,
cuEventCreate
__host____
device__cudaError_t
cudaEventCreateWithFlags (cudaEvent_t *event, unsigned
int flags)
Creates an event object with the specified flags.
Parameters
event
- Newly created event
flags
- Flags for new event
Returns
cudaSuccess, cudaErrorInvalidValue, cudaErrorLaunchFailure, cudaErrorMemoryAllocation
Description
Creates an event object for the current device with the specified flags. Valid flags include:
‣ cudaEventDefault: Default event creation flag.
‣ cudaEventBlockingSync: Specifies that event should use blocking synchronization. A host thread
that uses cudaEventSynchronize() to wait on an event created with this flag will block until the
event actually completes.
‣ cudaEventDisableTiming: Specifies that the created event does not need to record timing data.
Events created with this flag specified and the cudaEventBlockingSync flag not specified will
provide the best performance when used with cudaStreamWaitEvent() and cudaEventQuery().
|
48
Modules
‣ cudaEventInterprocess: Specifies that the created event may be used as an interprocess
event by cudaIpcGetEventHandle(). cudaEventInterprocess must be specified along with
cudaEventDisableTiming.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaEventCreate ( C API), cudaEventSynchronize, cudaEventDestroy, cudaEventElapsedTime,
cudaStreamWaitEvent, cuEventCreate
__host____
device__cudaError_t cudaEventDestroy
(cudaEvent_t event)
Destroys an event object.
Parameters
event
- Event to destroy
Returns
cudaSuccess, cudaErrorInvalidValue, cudaErrorInvalidResourceHandle, cudaErrorLaunchFailure
Description
Destroys the event specified by event.
An event may be destroyed before it is complete (i.e., while cudaEventQuery() would return
cudaErrorNotReady). 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.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
|
49
Modules
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
‣ Use of the handle after this call is undefined behavior.
‣ Returns cudaErrorInvalidResourceHandle in the event of being passed NULL as the input event.
See also:
cudaEventCreate ( C API), cudaEventCreateWithFlags, cudaEventQuery, cudaEventSynchronize,
cudaEventRecord, cudaEventElapsedTime, cuEventDestroy
__host__cudaError_t cudaEventElapsedTime (float *ms,
cudaEvent_t start, cudaEvent_t end)
Computes the elapsed time between events.
Parameters
ms
- Time between start and end in ms
start
- Starting event
end
- Ending event
Returns
cudaSuccess, cudaErrorNotReady, cudaErrorInvalidValue, cudaErrorInvalidResourceHandle,
cudaErrorLaunchFailure, cudaErrorUnknown
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 cudaEventRecord()
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 cudaEventRecord() has not been called on either event, then cudaErrorInvalidResourceHandle is
returned. If cudaEventRecord() has been called on both events but one or both of them has not yet been
completed (that is, cudaEventQuery() would return cudaErrorNotReady on at least one of the events),
cudaErrorNotReady is returned. If either event was created with the cudaEventDisableTiming flag,
then this function will return cudaErrorInvalidResourceHandle.
|
50
Modules
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
‣ Returns cudaErrorInvalidResourceHandle in the event of being passed NULL as the input event.
See also:
cudaEventCreate ( C API), cudaEventCreateWithFlags, cudaEventQuery, cudaEventSynchronize,
cudaEventDestroy, cudaEventRecord, cuEventElapsedTime
__host__cudaError_t cudaEventElapsedTime_v2 (float
*ms, cudaEvent_t start, cudaEvent_t end)
Computes the elapsed time between events.
Parameters
ms
- Time between start and end in ms
start
- Starting event
end
- Ending event
Returns
cudaSuccess, cudaErrorNotReady, cudaErrorInvalidValue, cudaErrorInvalidResourceHandle,
cudaErrorLaunchFailure, cudaErrorUnknown
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 a elapsed time calculation only and polling for completion on the events
to be compared should be done with cudaEventQuery 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 cudaEventRecord()
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.
|
51
Modules
If cudaEventRecord() has not been called on either event, then cudaErrorInvalidResourceHandle is
returned. If cudaEventRecord() has been called on both events but one or both of them has not yet been
completed (that is, cudaEventQuery() would return cudaErrorNotReady on at least one of the events),
cudaErrorNotReady is returned. If either event was created with the cudaEventDisableTiming flag,
then this function will return cudaErrorInvalidResourceHandle.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
‣ Returns cudaErrorInvalidResourceHandle in the event of being passed NULL as the input event.
See also:
cudaEventCreate ( C API), cudaEventCreateWithFlags, cudaEventQuery, cudaEventSynchronize,
cudaEventDestroy, cudaEventRecord, cuEventElapsedTime
__host__cudaError_t cudaEventQuery (cudaEvent_t event)
Queries an event's status.
Parameters
event
- Event to query
Returns
cudaSuccess, cudaErrorNotReady, cudaErrorInvalidValue, cudaErrorInvalidResourceHandle,
cudaErrorLaunchFailure
Description
Queries the status of all work currently captured by event. See cudaEventRecord() for details on what
is captured by an event.
Returns cudaSuccess if all captured work has been completed, or cudaErrorNotReady if any captured
work is incomplete.
For the purposes of Unified Memory, a return value of cudaSuccess is equivalent to having called
cudaEventSynchronize().
|
52
Modules
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
‣ Returns cudaErrorInvalidResourceHandle in the event of being passed NULL as the input event.
See also:
cudaEventCreate ( C API), cudaEventCreateWithFlags, cudaEventRecord, cudaEventSynchronize,
cudaEventDestroy, cudaEventElapsedTime, cuEventQuery
__host____
device__cudaError_t cudaEventRecord
(cudaEvent_t event, cudaStream_t stream)
Records an event.
Parameters
event
- Event to record
stream
- Stream in which to record event
Returns
cudaSuccess, cudaErrorInvalidValue, cudaErrorInvalidResourceHandle, cudaErrorLaunchFailure
Description
Captures in event the contents of stream at the time of this call. event and stream must be
on the same CUDA context. Calls such as cudaEventQuery() or cudaStreamWaitEvent() will then
examine or wait for completion of the work that was captured. Uses of stream after this call do not
modify event. See note on default stream behavior for what is captured in the default case.
cudaEventRecord() can be called multiple times on the same event and will overwrite the previously
captured state. Other APIs such as cudaStreamWaitEvent() use the most recently captured state at the
time of the API call, and are not affected by later calls to cudaEventRecord(). Before the first call to
cudaEventRecord(), an event represents an empty set of work, so for example cudaEventQuery() would
return cudaSuccess.
Note:
|
53
Modules
‣ This function uses standard default stream semantics.
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
‣ Returns cudaErrorInvalidResourceHandle in the event of being passed NULL as the input event.
See also:
cudaEventCreate ( C API), cudaEventCreateWithFlags, cudaEventQuery, cudaEventSynchronize,
cudaEventDestroy, cudaEventElapsedTime, cudaStreamWaitEvent, cudaEventRecordWithFlags,
cuEventRecord
__host__cudaError_t cudaEventRecordWithFlags
(cudaEvent_t event, cudaStream_t stream, unsigned int
flags)
Records an event.
Parameters
event
- Event to record
stream
- Stream in which to record event
flags
- Parameters for the operation(See above)
Returns
cudaSuccess, cudaErrorInvalidValue, cudaErrorInvalidResourceHandle, cudaErrorLaunchFailure
Description
Captures in event the contents of stream at the time of this call. event and stream must be
on the same CUDA context. Calls such as cudaEventQuery() or cudaStreamWaitEvent() will then
examine or wait for completion of the work that was captured. Uses of stream after this call do not
modify event. See note on default stream behavior for what is captured in the default case.
cudaEventRecordWithFlags() can be called multiple times on the same event and will overwrite the
previously captured state. Other APIs such as cudaStreamWaitEvent() use the most recently captured
state at the time of the API call, and are not affected by later calls to cudaEventRecordWithFlags().
Before the first call to cudaEventRecordWithFlags(), an event represents an empty set of work, so for
example cudaEventQuery() would return cudaSuccess.
|
54
Modules
flags include:
‣ cudaEventRecordDefault: Default event creation flag.
‣ cudaEventRecordExternal: Event is captured in the graph as an external event node when
performing stream capture.
Note:
‣ This function uses standard default stream semantics.
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
‣ Returns cudaErrorInvalidResourceHandle in the event of being passed NULL as the input event.
See also:
cudaEventCreate ( C API), cudaEventCreateWithFlags, cudaEventQuery, cudaEventSynchronize,
cudaEventDestroy, cudaEventElapsedTime, cudaStreamWaitEvent, cudaEventRecord, cuEventRecord,
__host__cudaError_t cudaEventSynchronize (cudaEvent_t
event)
Waits for an event to complete.
Parameters
event
- Event to wait for
Returns
cudaSuccess, cudaErrorInvalidValue, cudaErrorInvalidResourceHandle, cudaErrorLaunchFailure
Description
Waits until the completion of all work currently captured in event. See cudaEventRecord() for details
on what is captured by an event.
Waiting for an event that was created with the cudaEventBlockingSync flag will cause the calling CPU
thread to block until the event has been completed by the device. If the cudaEventBlockingSync flag
has not been set, then the CPU thread will busy-wait until the event has been completed by the device.
|
55
Modules
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
‣ Returns cudaErrorInvalidResourceHandle in the event of being passed NULL as the input event.
See also:
cudaEventCreate ( C API), cudaEventCreateWithFlags, cudaEventRecord, cudaEventQuery,
cudaEventDestroy, cudaEventElapsedTime, cuEventSynchronize
6.7.
External Resource Interoperability
This section describes the external resource interoperability functions of the CUDA runtime application
programming interface.
__host__cudaError_t cudaDestroyExternalMemory
(cudaExternalMemory_t extMem)
Destroys an external memory object.
Parameters
extMem
- External memory object to be destroyed
Returns
cudaSuccess, cudaErrorInvalidResourceHandle
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 cudaFree and
cudaFreeMipmappedArray respectively.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
|
56
Modules
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
‣ Use of the handle after this call is undefined behavior.
See also:
cudaImportExternalMemory, cudaExternalMemoryGetMappedBuffer,
cudaExternalMemoryGetMappedMipmappedArray
__host__cudaError_t cudaDestroyExternalSemaphore
(cudaExternalSemaphore_t extSem)
Destroys an external semaphore.
Parameters
extSem
- External semaphore to be destroyed
Returns
cudaSuccess, cudaErrorInvalidResourceHandle
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.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
‣ Use of the handle after this call is undefined behavior.
See also:
cudaImportExternalSemaphore, cudaSignalExternalSemaphoresAsync,
cudaWaitExternalSemaphoresAsync
|
57
Modules
__host__cudaError_t
cudaExternalMemoryGetMappedBuffer (void
**devPtr, cudaExternalMemory_t extMem, const
cudaExternalMemoryBufferDesc *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
cudaSuccess, cudaErrorInvalidValue, cudaErrorInvalidResourceHandle
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
cudaExternalMemoryBufferDesc structure is defined as follows:
typedef struct cudaExternalMemoryBufferDesc_st {
unsigned long long offset;
unsigned long long size;
unsigned int flags;
} cudaExternalMemoryBufferDesc;
where cudaExternalMemoryBufferDesc::offset is the offset in the memory object where
the buffer's base address is. cudaExternalMemoryBufferDesc::size is the size of the buffer.
cudaExternalMemoryBufferDesc::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
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 cudaFree.
|
58
Modules
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaImportExternalMemory, cudaDestroyExternalMemory,
cudaExternalMemoryGetMappedMipmappedArray
__host__cudaError_t
cudaExternalMemoryGetMappedMipmappedArray
(cudaMipmappedArray_t *mipmap,
cudaExternalMemory_t extMem, const
cudaExternalMemoryMipmappedArrayDesc
*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
cudaSuccess, cudaErrorInvalidValue, cudaErrorInvalidResourceHandle
Description
Maps a CUDA mipmapped array onto an external object and returns a handle to it in mipmap.
The properties of the CUDA mipmapped array being mapped must be described in mipmapDesc. The
structure cudaExternalMemoryMipmappedArrayDesc is defined as follows:
typedef struct cudaExternalMemoryMipmappedArrayDesc_st {
unsigned long long offset;
cudaChannelFormatDesc formatDesc;
|
59
Modules
cudaExtent extent;
unsigned int flags;
unsigned int numLevels;
} cudaExternalMemoryMipmappedArrayDesc;
where cudaExternalMemoryMipmappedArrayDesc::offset is the offset in the memory object where
the base level of the mipmap chain is. cudaExternalMemoryMipmappedArrayDesc::formatDesc
describes the format of the data. cudaExternalMemoryMipmappedArrayDesc::extent specifies the
dimensions of the base level of the mipmap chain. cudaExternalMemoryMipmappedArrayDesc::flags
are flags associated with CUDA mipmapped arrays. For further details, please refer
to the documentation for cudaMalloc3DArray. Note that if the mipmapped array is
bound as a color target in the graphics API, then the flag cudaArrayColorAttachment
must be specified in cudaExternalMemoryMipmappedArrayDesc::flags.
cudaExternalMemoryMipmappedArrayDesc::numLevels specifies the total number of levels in the
mipmap chain.
The returned CUDA mipmapped array must be freed using cudaFreeMipmappedArray.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaImportExternalMemory, cudaDestroyExternalMemory, cudaExternalMemoryGetMappedBuffer
Note:
If cudaExternalMemoryHandleDesc::type is cudaExternalMemoryHandleTypeNvSciBuf, then
cudaExternalMemoryMipmappedArrayDesc::numLevels must not be greater than 1.
__host__cudaError_t cudaImportExternalMemory
(cudaExternalMemory_t *extMem_out, const
cudaExternalMemoryHandleDesc *memHandleDesc)
Imports an external memory object.
Parameters
extMem_out
- Returned handle to an external memory object
|
60
Modules
memHandleDesc
- Memory import handle descriptor
Returns
cudaSuccess, cudaErrorInvalidValue, cudaErrorInvalidResourceHandle, cudaErrorOperatingSystem
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
cudaExternalMemoryHandleDesc structure is defined as follows:
typedef struct cudaExternalMemoryHandleDesc_st {
cudaExternalMemoryHandleType type;
union {
int fd;
struct {
void *handle;
const void *name;
} win32;
const void *nvSciBufObject;
} handle;
unsigned long long size;
unsigned int flags;
} cudaExternalMemoryHandleDesc;
where cudaExternalMemoryHandleDesc::type specifies the type of handle being imported.
cudaExternalMemoryHandleType is defined as:
typedef enum cudaExternalMemoryHandleType_enum {
cudaExternalMemoryHandleTypeOpaqueFd
= 1,
cudaExternalMemoryHandleTypeOpaqueWin32
= 2,
cudaExternalMemoryHandleTypeOpaqueWin32Kmt
= 3,
cudaExternalMemoryHandleTypeD3D12Heap
= 4,
cudaExternalMemoryHandleTypeD3D12Resource
= 5,
cudaExternalMemoryHandleTypeD3D11Resource
= 6,
cudaExternalMemoryHandleTypeD3D11ResourceKmt = 7,
cudaExternalMemoryHandleTypeNvSciBuf
= 8
} cudaExternalMemoryHandleType;
If cudaExternalMemoryHandleDesc::type is cudaExternalMemoryHandleTypeOpaqueFd, then
cudaExternalMemoryHandleDesc::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 cudaExternalMemoryHandleDesc::type is cudaExternalMemoryHandleTypeOpaqueWin32,
then exactly one of cudaExternalMemoryHandleDesc::handle::win32::handle and
cudaExternalMemoryHandleDesc::handle::win32::name must not be NULL. If
cudaExternalMemoryHandleDesc::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 cudaExternalMemoryHandleDesc::handle::win32::name is not NULL, then it must point
to a NULL-terminated array of UTF-16 characters that refers to a memory object.
|
61
Modules
If cudaExternalMemoryHandleDesc::type is cudaExternalMemoryHandleTypeOpaqueWin32Kmt,
then cudaExternalMemoryHandleDesc::handle::win32::handle must be non-NULL and
cudaExternalMemoryHandleDesc::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 cudaExternalMemoryHandleDesc::type is cudaExternalMemoryHandleTypeD3D12Heap,
then exactly one of cudaExternalMemoryHandleDesc::handle::win32::handle and
cudaExternalMemoryHandleDesc::handle::win32::name must not be NULL. If
cudaExternalMemoryHandleDesc::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
cudaExternalMemoryHandleDesc::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 cudaExternalMemoryHandleDesc::type is cudaExternalMemoryHandleTypeD3D12Resource,
then exactly one of cudaExternalMemoryHandleDesc::handle::win32::handle and
cudaExternalMemoryHandleDesc::handle::win32::name must not be NULL. If
cudaExternalMemoryHandleDesc::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
cudaExternalMemoryHandleDesc::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 cudaExternalMemoryHandleDesc::type is cudaExternalMemoryHandleTypeD3D11Resource,then
exactly one of cudaExternalMemoryHandleDesc::handle::win32::handle and
cudaExternalMemoryHandleDesc::handle::win32::name must not be NULL. If
cudaExternalMemoryHandleDesc::handle::win32::handle is not NULL, then it must represent a valid
shared NT handle that is returned by IDXGIResource1::CreateSharedHandle when referring to a
ID3D11Resource object. If cudaExternalMemoryHandleDesc::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 cudaExternalMemoryHandleDesc::type is cudaExternalMemoryHandleTypeD3D11ResourceKmt,
then cudaExternalMemoryHandleDesc::handle::win32::handle must be non-NULL and
cudaExternalMemoryHandleDesc::handle::win32::name must be NULL. The handle specified must be
a valid shared KMT handle that is returned by IDXGIResource::GetSharedHandle when referring to a
ID3D11Resource object.
If cudaExternalMemoryHandleDesc::type is cudaExternalMemoryHandleTypeNvSciBuf,
then cudaExternalMemoryHandleDesc::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 cudaWaitExternalSemaphoresAsync or
cudaSignalExternalSemaphoresAsync as approprriate barriers to maintain coherence between
CUDA and the other drivers. See cudaExternalSemaphoreWaitSkipNvSciBufMemSync and
cudaExternalSemaphoreSignalSkipNvSciBufMemSync for memory synchronization.
|
62
Modules
The size of the memory object must be specified in cudaExternalMemoryHandleDesc::size.
Specifying the flag cudaExternalMemoryDedicated in cudaExternalMemoryHandleDesc::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 cudaExternalMemoryHandleDesc::type is one of the following:
cudaExternalMemoryHandleTypeD3D12Resource cudaExternalMemoryHandleTypeD3D11Resource
cudaExternalMemoryHandleTypeD3D11ResourceKmt
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
‣ 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:
cudaDestroyExternalMemory, cudaExternalMemoryGetMappedBuffer,
cudaExternalMemoryGetMappedMipmappedArray
__host__cudaError_t cudaImportExternalSemaphore
(cudaExternalSemaphore_t *extSem_out, const
cudaExternalSemaphoreHandleDesc *semHandleDesc)
Imports an external semaphore.
Parameters
extSem_out
- Returned handle to an external semaphore
semHandleDesc
- Semaphore import handle descriptor
Returns
cudaSuccess, cudaErrorInvalidResourceHandle, cudaErrorOperatingSystem
|
63
Modules
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
cudaExternalSemaphoreHandleDesc is defined as follows:
typedef struct cudaExternalSemaphoreHandleDesc_st {
cudaExternalSemaphoreHandleType type;
union {
int fd;
struct {
void *handle;
const void *name;
} win32;
const void* NvSciSyncObj;
} handle;
unsigned int flags;
} cudaExternalSemaphoreHandleDesc;
where cudaExternalSemaphoreHandleDesc::type specifies the type of handle being imported.
cudaExternalSemaphoreHandleType is defined as:
typedef enum cudaExternalSemaphoreHandleType_enum {
cudaExternalSemaphoreHandleTypeOpaqueFd
= 1,
cudaExternalSemaphoreHandleTypeOpaqueWin32
= 2,
cudaExternalSemaphoreHandleTypeOpaqueWin32Kmt
= 3,
cudaExternalSemaphoreHandleTypeD3D12Fence
= 4,
cudaExternalSemaphoreHandleTypeD3D11Fence
= 5,
cudaExternalSemaphoreHandleTypeNvSciSync
= 6,
cudaExternalSemaphoreHandleTypeKeyedMutex
= 7,
cudaExternalSemaphoreHandleTypeKeyedMutexKmt
= 8,
cudaExternalSemaphoreHandleTypeTimelineSemaphoreFd
= 9,
cudaExternalSemaphoreHandleTypeTimelineSemaphoreWin32
= 10
} cudaExternalSemaphoreHandleType;
If cudaExternalSemaphoreHandleDesc::type is cudaExternalSemaphoreHandleTypeOpaqueFd,
then cudaExternalSemaphoreHandleDesc::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 cudaExternalSemaphoreHandleDesc::type is cudaExternalSemaphoreHandleTypeOpaqueWin32,
then exactly one of cudaExternalSemaphoreHandleDesc::handle::win32::handle and
cudaExternalSemaphoreHandleDesc::handle::win32::name must not be NULL. If
cudaExternalSemaphoreHandleDesc::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 cudaExternalSemaphoreHandleDesc::handle::win32::name is not NULL,
then it must name a valid synchronization object.
If cudaExternalSemaphoreHandleDesc::type is
cudaExternalSemaphoreHandleTypeOpaqueWin32Kmt, then
cudaExternalSemaphoreHandleDesc::handle::win32::handle must be non-NULL and
cudaExternalSemaphoreHandleDesc::handle::win32::name must be NULL. The handle specified must
|
64
Modules
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 cudaExternalSemaphoreHandleDesc::type is cudaExternalSemaphoreHandleTypeD3D12Fence,
then exactly one of cudaExternalSemaphoreHandleDesc::handle::win32::handle and
cudaExternalSemaphoreHandleDesc::handle::win32::name must not be NULL. If
cudaExternalSemaphoreHandleDesc::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
cudaExternalSemaphoreHandleDesc::handle::win32::name is not NULL, then it must name a valid
synchronization object that refers to a valid ID3D12Fence object.
If cudaExternalSemaphoreHandleDesc::type is cudaExternalSemaphoreHandleTypeD3D11Fence,
then exactly one of cudaExternalSemaphoreHandleDesc::handle::win32::handle and
cudaExternalSemaphoreHandleDesc::handle::win32::name must not be NULL. If
cudaExternalSemaphoreHandleDesc::handle::win32::handle is not NULL, then it must
represent a valid shared NT handle that is returned by ID3D11Fence::CreateSharedHandle. If
cudaExternalSemaphoreHandleDesc::handle::win32::name is not NULL, then it must name a valid
synchronization object that refers to a valid ID3D11Fence object.
If cudaExternalSemaphoreHandleDesc::type is cudaExternalSemaphoreHandleTypeNvSciSync, then
cudaExternalSemaphoreHandleDesc::handle::nvSciSyncObj represents a valid NvSciSyncObj.
cudaExternalSemaphoreHandleTypeKeyedMutex, then exactly one
of cudaExternalSemaphoreHandleDesc::handle::win32::handle and
cudaExternalSemaphoreHandleDesc::handle::win32::name must not be NULL. If
cudaExternalSemaphoreHandleDesc::handle::win32::handle is not NULL, then it represent a valid
shared NT handle that is returned by IDXGIResource1::CreateSharedHandle when referring to a
IDXGIKeyedMutex object.
If cudaExternalSemaphoreHandleDesc::type is cudaExternalSemaphoreHandleTypeKeyedMutexKmt,
then cudaExternalSemaphoreHandleDesc::handle::win32::handle must be non-NULL and
cudaExternalSemaphoreHandleDesc::handle::win32::name must be NULL. The handle specified must
represent a valid KMT handle that is returned by IDXGIResource::GetSharedHandle when referring to
a IDXGIKeyedMutex object.
If cudaExternalSemaphoreHandleDesc::type is
cudaExternalSemaphoreHandleTypeTimelineSemaphoreFd, then
cudaExternalSemaphoreHandleDesc::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 cudaExternalSemaphoreHandleDesc::type is
cudaExternalSemaphoreHandleTypeTimelineSemaphoreWin32, then exactly
one of cudaExternalSemaphoreHandleDesc::handle::win32::handle and
cudaExternalSemaphoreHandleDesc::handle::win32::name must not be NULL. If
cudaExternalSemaphoreHandleDesc::handle::win32::handle is not NULL, then it must represent
|
65
Modules
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 cudaExternalSemaphoreHandleDesc::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.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaDestroyExternalSemaphore, cudaSignalExternalSemaphoresAsync,
cudaWaitExternalSemaphoresAsync
__host__cudaError_t cudaSignalExternalSemaphoresAsync
(const cudaExternalSemaphore_t *extSemArray, const
cudaExternalSemaphoreSignalParams *paramsArray,
unsigned int numExtSems, cudaStream_t stream)
Signals a set of external semaphore objects.
Parameters
extSemArray
- Set of external semaphores to be signaled
paramsArray
- Array of semaphore parameters
numExtSems
- Number of semaphores to signal
stream
- Stream to enqueue the signal operations in
Returns
cudaSuccess, cudaErrorInvalidResourceHandle
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.
|
66
Modules
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:
cudaExternalSemaphoreHandleTypeOpaqueFd, cudaExternalSemaphoreHandleTypeOpaqueWin32,
cudaExternalSemaphoreHandleTypeOpaqueWin32Kmt then signaling the semaphore will set it to the
signaled state.
If the semaphore object is any one of the following types:
cudaExternalSemaphoreHandleTypeD3D12Fence, cudaExternalSemaphoreHandleTypeD3D11Fence,
cudaExternalSemaphoreHandleTypeTimelineSemaphoreFd,
cudaExternalSemaphoreHandleTypeTimelineSemaphoreWin32 then the semaphore will be set to the
value specified in cudaExternalSemaphoreSignalParams::params::fence::value.
If the semaphore object is of the type cudaExternalSemaphoreHandleTypeNvSciSync this
API sets cudaExternalSemaphoreSignalParams::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
cudaExternalSemaphoreSignalParams::params::nvSciSync::fence. By default, signaling such an
external semaphore object causes appropriate memory synchronization operations to be performed
over all the external memory objects that are imported as cudaExternalMemoryHandleTypeNvSciBuf.
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
cudaExternalSemaphoreSignalSkipNvSciBufMemSync, 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 cudaExternalSemaphoreHandleTypeNvSciSync, if the NvSciSyncAttrList used
to create the NvSciSyncObj had not set the flags in cudaDeviceGetNvSciSyncAttributes to
cudaNvSciSyncAttrSignal, this API will return cudaErrorNotSupported.
cudaExternalSemaphoreSignalParams::params::nvSciSync::fence associated with semaphore
object of the type cudaExternalSemaphoreHandleTypeNvSciSync 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:
cudaExternalSemaphoreHandleTypeKeyedMutex,
|
67
Modules
cudaExternalSemaphoreHandleTypeKeyedMutexKmt, then the keyed mutex will be released with the
key specified in cudaExternalSemaphoreSignalParams::params::keyedmutex::key.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaImportExternalSemaphore, cudaDestroyExternalSemaphore, cudaWaitExternalSemaphoresAsync
__host__cudaError_t cudaWaitExternalSemaphoresAsync
(const cudaExternalSemaphore_t *extSemArray, const
cudaExternalSemaphoreWaitParams *paramsArray,
unsigned int numExtSems, cudaStream_t stream)
Waits on a set of external semaphore objects.
Parameters
extSemArray
- External semaphores to be waited on
paramsArray
- Array of semaphore parameters
numExtSems
- Number of semaphores to wait on
stream
- Stream to enqueue the wait operations in
Returns
cudaSuccess, cudaErrorInvalidResourceHandle cudaErrorTimeout
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.
|
68
Modules
If the semaphore object is any one of the following types:
cudaExternalSemaphoreHandleTypeOpaqueFd, cudaExternalSemaphoreHandleTypeOpaqueWin32,
cudaExternalSemaphoreHandleTypeOpaqueWin32Kmt 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:
cudaExternalSemaphoreHandleTypeD3D12Fence, cudaExternalSemaphoreHandleTypeD3D11Fence,
cudaExternalSemaphoreHandleTypeTimelineSemaphoreFd,
cudaExternalSemaphoreHandleTypeTimelineSemaphoreWin32 then waiting on
the semaphore will wait until the value of the semaphore is greater than or equal to
cudaExternalSemaphoreWaitParams::params::fence::value.
If the semaphore object is of the type cudaExternalSemaphoreHandleTypeNvSciSync then, waiting
on the semaphore will wait until the cudaExternalSemaphoreSignalParams::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
cudaExternalMemoryHandleTypeNvSciBuf. 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 cudaExternalSemaphoreWaitSkipNvSciBufMemSync, 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 cudaExternalSemaphoreHandleTypeNvSciSync, if the NvSciSyncAttrList
used to create the NvSciSyncObj had not set the flags in cudaDeviceGetNvSciSyncAttributes to
cudaNvSciSyncAttrWait, this API will return cudaErrorNotSupported.
If the semaphore object is any one of the following types:
cudaExternalSemaphoreHandleTypeKeyedMutex,
cudaExternalSemaphoreHandleTypeKeyedMutexKmt, then the keyed
mutex will be acquired when it is released with the key specified in
cudaExternalSemaphoreSignalParams::params::keyedmutex::key or until the timeout specified by
cudaExternalSemaphoreSignalParams::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.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
|
69
Modules
See also:
cudaImportExternalSemaphore, cudaDestroyExternalSemaphore,
cudaSignalExternalSemaphoresAsync
6.8.
Execution Control
This section describes the execution control functions of the CUDA runtime application programming
interface.
Some functions have overloaded C++ API template versions documented separately in the C++ API
Routines module.
__host____
device__cudaError_t cudaFuncGetAttributes
(cudaFuncAttributes *attr, const void *func)
Find out attributes for a given function.
Parameters
attr
- Return pointer to function's attributes
func
- Device function symbol
Returns
cudaSuccess, cudaErrorInvalidDeviceFunction
Description
This function obtains the attributes of a function specified via func. func is a device function symbol
and must be declared as a __global__ function. The fetched attributes are placed in attr. If the
specified function does not exist, then it is assumed to be a cudaKernel_t and used as is. For templated
functions, pass the function symbol as follows: func_name<template_arg_0,...,template_arg_N>
Note that some function attributes such as maxThreadsPerBlock may vary based on the device that is
currently being used.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Use of a string naming a function as the func parameter was deprecated in CUDA 4.1 and
removed in CUDA 5.0.
|
70
Modules
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
‣ The API can also be used with a kernel cudaKernel_t by querying the handle using
cudaLibraryGetKernel() or cudaGetKernel and then passing it to the API by casting to void*. The
symbol entryFuncAddr passed to cudaGetKernel should be a symbol that is registered with the
same CUDA Runtime instance.
‣ Passing a symbol that belongs that belongs to a different runtime instance will result in undefined
behavior. The only type that can be reliably passed to a different runtime instance is cudaKernel_t
See also:
cudaFuncSetCacheConfig ( C API), cudaFuncGetAttributes ( C++ API), cudaLaunchKernel ( C API),
cuFuncGetAttribute
__host__cudaError_t cudaFuncGetName (const char
**name, const void *func)
Returns the function name for a device entry function pointer.
Parameters
name
- The returned name of the function
func
- The function pointer to retrieve name for
Returns
cudaSuccess, cudaErrorInvalidValue, cudaErrorInvalidDeviceFunction
Description
Returns in **name the function name associated with the symbol func . The function name is
returned as a null-terminated string. This API may return a mangled name if the function is not
declared as having C linkage. If **name is NULL, cudaErrorInvalidValue is returned. If func is not
a device entry function, then it is assumed to be a cudaKernel_t and used as is.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
|
71
Modules
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
‣ The API can also be used with a kernel cudaKernel_t by querying the handle using
cudaLibraryGetKernel() or cudaGetKernel and then passing it to the API by casting to void*. The
symbol entryFuncAddr passed to cudaGetKernel should be a symbol that is registered with the
same CUDA Runtime instance.
‣ Passing a symbol that belongs that belongs to a different runtime instance will result in undefined
behavior. The only type that can be reliably passed to a different runtime instance is cudaKernel_t
cudaFuncGetName (C++ API)
__host__cudaError_t cudaFuncGetParamInfo (const void
*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
paramOffset
- The offset into the device-side parameter layout at which the parameter resides
paramSize
- 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 in func's list of parameters and returns parameter
information via paramOffset and paramSize. paramOffset returns the offset of the
parameter in the device-side parameter layout. paramSize returns the size in bytes of the
parameter. 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.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
|
72
Modules
‣ The API can also be used with a kernel cudaKernel_t by querying the handle using
cudaLibraryGetKernel() or cudaGetKernel and then passing it to the API by casting to void*. The
symbol entryFuncAddr passed to cudaGetKernel should be a symbol that is registered with the
same CUDA Runtime instance.
‣ Passing a symbol that belongs that belongs to a different runtime instance will result in undefined
behavior. The only type that can be reliably passed to a different runtime instance is cudaKernel_t
__host__cudaError_t cudaFuncSetAttribute (const void
*func, cudaFuncAttribute attr, int value)
Set attributes for a given function.
Parameters
func
- Function to get attributes of
attr
- Attribute to set
value
- Value to set
Returns
cudaSuccess, cudaErrorInvalidDeviceFunction, cudaErrorInvalidValue
Description
This function sets the attributes of a function specified via func. The parameter func must be a
pointer to a function that executes on the device. The parameter specified by func must be declared
as a __global__ function. The enumeration defined by attr is set to the value defined by value.
If the specified function does not exist, then it is assumed to be a cudaKernel_t and used as is. If
the specified attribute cannot be written, or if the value is incorrect, then cudaErrorInvalidValue is
returned.
Valid values for attr are:
‣ cudaFuncAttributeMaxDynamicSharedMemorySize - The requested maximum
size in bytes of dynamically-allocated shared memory. The sum of this value
and the function attribute sharedSizeBytes cannot exceed the device attribute
cudaDevAttrMaxSharedMemoryPerBlockOptin. The maximal size of requestable dynamic shared
memory may differ by GPU architecture.
‣ cudaFuncAttributePreferredSharedMemoryCarveout - On devices where the L1 cache and shared
memory use the same hardware resources, this sets the shared memory carveout preference, in
percent of the total shared memory. See cudaDevAttrMaxSharedMemoryPerMultiprocessor. This
is only a hint, and the driver can choose a different ratio if required to execute the function.
|
73
Modules
‣ cudaFuncAttributeRequiredClusterWidth: The required cluster width in blocks. The width, height,
and depth values must either all be 0 or all be positive. The validity of the cluster dimensions is
checked at launch time. If the value is set during compile time, it cannot be set at runtime. Setting it
at runtime will return cudaErrorNotPermitted.
‣ cudaFuncAttributeRequiredClusterHeight: The required cluster height in blocks. The width, height,
and depth values must either all be 0 or all be positive. The validity of the cluster dimensions is
checked at launch time. If the value is set during compile time, it cannot be set at runtime. Setting it
at runtime will return cudaErrorNotPermitted.
‣ cudaFuncAttributeRequiredClusterDepth: The required cluster depth in blocks. The width, height,
and depth values must either all be 0 or all be positive. The validity of the cluster dimensions is
checked at launch time. If the value is set during compile time, it cannot be set at runtime. Setting it
at runtime will return cudaErrorNotPermitted.
‣ cudaFuncAttributeNonPortableClusterSizeAllowed: Indicates whether the function can be
launched with non-portable cluster size. 1 is allowed, 0 is disallowed.
‣ cudaFuncAttributeClusterSchedulingPolicyPreference: The block scheduling policy of a function.
The value type is cudaClusterSchedulingPolicy.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
‣ The API can also be used with a kernel cudaKernel_t by querying the handle using
cudaLibraryGetKernel() or cudaGetKernel and then passing it to the API by casting to void*. The
symbol entryFuncAddr passed to cudaGetKernel should be a symbol that is registered with the
same CUDA Runtime instance.
‣ Passing a symbol that belongs that belongs to a different runtime instance will result in undefined
behavior. The only type that can be reliably passed to a different runtime instance is cudaKernel_t
cudaLaunchKernel (C++ API), cudaFuncSetCacheConfig ( C++ API), cudaFuncGetAttributes ( C
API),
|
74
Modules
__host__cudaError_t cudaFuncSetCacheConfig (const void
*func, cudaFuncCache cacheConfig)
Sets the preferred cache configuration for a device function.
Parameters
func
- Device function symbol
cacheConfig
- Requested cache configuration
Returns
cudaSuccess, cudaErrorInvalidDeviceFunction
Description
On devices where the L1 cache and shared memory use the same hardware resources, this sets through
cacheConfig the preferred cache configuration for the function specified via func. This is only
a preference. The runtime will use the requested configuration if possible, but it is free to choose a
different configuration if required to execute func.
func is a device function symbol and must be declared as a __global__ function. If the specified
function does not exist, then cudaErrorInvalidDeviceFunction is returned. For templated functions, pass
the function symbol as follows: func_name<template_arg_0,...,template_arg_N>
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:
‣ cudaFuncCachePreferNone: no preference for shared memory or L1 (default)
‣ cudaFuncCachePreferShared: prefer larger shared memory and smaller L1 cache
‣ cudaFuncCachePreferL1: prefer larger L1 cache and smaller shared memory
‣ cudaFuncCachePreferEqual: prefer equal size L1 cache and shared memory
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Use of a string naming a function as the func parameter was deprecated in CUDA 4.1 and
removed in CUDA 5.0.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
|
75
Modules
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
‣ This API does not accept a cudaKernel_t casted as void*. If cache config modification is
required for a cudaKernel_t (or a __global__ function), it can be replaced with a call to
cudaFuncSetAttributes with the attribute cudaFuncAttributePreferredSharedMemoryCarveout to
specify a more granular L1 cache and shared memory split configuration.
See also:
cudaFuncSetCacheConfig ( C++ API), cudaFuncGetAttributes ( C API), cudaLaunchKernel ( C API),
cuFuncSetCacheConfig
__device__ void *cudaGetParameterBuffer (size_t
alignment, size_t size)
Obtains a parameter buffer.
Parameters
alignment
- Specifies alignment requirement of the parameter buffer
size
- Specifies size requirement in bytes
Returns
Returns pointer to the allocated parameterBuffer
Description
Obtains a parameter buffer which can be filled with parameters for a kernel launch. Parameters passed
to cudaLaunchDevice must be allocated via this function.
This is a low level API and can only be accessed from Parallel Thread Execution (PTX). CUDA user
code should use <<< >>> to launch kernels.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cudaLaunchDevice
|
76
Modules
__device__ void cudaGridDependencySynchronize (void)
Programmatic grid dependency synchronization.
Description
This device function will block the thread until all direct grid dependencies have completed.
This API is intended to use in conjuncture with programmatic / launch event / dependency.
See cudaLaunchAttributeID::cudaLaunchAttributeProgrammaticStreamSerialization and
cudaLaunchAttributeID::cudaLaunchAttributeProgrammaticEvent for more information.
__host__cudaError_t cudaLaunchCooperativeKernel (const
void *func, dim3 gridDim, dim3 blockDim, void **args,
size_t sharedMem, cudaStream_t stream)
Launches a device function where thread blocks can cooperate and synchronize as they execute.
Parameters
func
- Device function symbol
gridDim
- Grid dimentions
blockDim
- Block dimentions
args
- Arguments
sharedMem
- Shared memory
stream
- Stream identifier
Returns
cudaSuccess, cudaErrorInvalidDeviceFunction, cudaErrorInvalidConfiguration,
cudaErrorLaunchFailure, cudaErrorLaunchTimeout, cudaErrorLaunchOutOfResources,
cudaErrorCooperativeLaunchTooLarge, cudaErrorSharedObjectInitFailed
Description
The function invokes kernel func on gridDim (gridDim.x gridDim.y gridDim.z) grid of
blocks. Each block contains blockDim (blockDim.x blockDim.y blockDim.z) threads.
The device on which this kernel is invoked must have a non-zero value for the device attribute
cudaDevAttrCooperativeLaunch.
|
77
Modules
The total number of blocks launched cannot exceed the maximum number of blocks per
multiprocessor as returned by cudaOccupancyMaxActiveBlocksPerMultiprocessor (or
cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags) times the number of multiprocessors as
specified by the device attribute cudaDevAttrMultiProcessorCount.
The kernel cannot make use of CUDA dynamic parallelism.
If the kernel has N parameters the args should point to array of N pointers. Each pointer, from
args[0] to args[N - 1], point to the region of memory from which the actual parameter will be
copied.
For templated functions, pass the function symbol as follows:
func_name<template_arg_0,...,template_arg_N>
sharedMem sets the amount of dynamic shared memory that will be available to each thread block.
stream specifies a stream the invocation is associated to.
Note:
‣ This function uses standard default stream semantics.
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
‣ The API can also be used with a kernel cudaKernel_t by querying the handle using
cudaLibraryGetKernel() or cudaGetKernel and then passing it to the API by casting to void*. The
symbol entryFuncAddr passed to cudaGetKernel should be a symbol that is registered with the
same CUDA Runtime instance.
‣ Passing a symbol that belongs that belongs to a different runtime instance will result in undefined
behavior. The only type that can be reliably passed to a different runtime instance is cudaKernel_t
See also:
cudaLaunchCooperativeKernel (C++ API), cudaLaunchCooperativeKernelMultiDevice,
cuLaunchCooperativeKernel
|
78
Modules
__host__cudaError_t
cudaLaunchCooperativeKernelMultiDevice
(cudaLaunchParams *launchParamsList, unsigned int
numDevices, unsigned int flags)
Launches device 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
cudaSuccess, cudaErrorInvalidDeviceFunction, cudaErrorInvalidConfiguration,
cudaErrorLaunchFailure, cudaErrorLaunchTimeout, cudaErrorLaunchOutOfResources,
cudaErrorCooperativeLaunchTooLarge, cudaErrorSharedObjectInitFailed
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
cudaDevAttrCooperativeMultiDeviceLaunch.
The same kernel must be launched on all devices. Note that any __device__ or __constant__ variables
are independently instantiated on every device. It is the application's responsiblity 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 cudaStreamCreate
or cudaStreamCreateWithPriority or cudaStreamCreateWithPriority. The NULL stream or
cudaStreamLegacy or cudaStreamPerThread cannot be used.
The total number of blocks launched per kernel cannot exceed the maximum number of blocks
per multiprocessor as returned by cudaOccupancyMaxActiveBlocksPerMultiprocessor (or
|
79
Modules
cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags) times the number of multiprocessors
as specified by the device attribute cudaDevAttrMultiProcessorCount. 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 kernel cannot make use of CUDA dynamic parallelism.
The cudaLaunchParams structure is defined as:
struct cudaLaunchParams
{
void *func;
dim3 gridDim;
dim3 blockDim;
void **args;
size_t sharedMem;
cudaStream_t
stream;
};
where:
‣ cudaLaunchParams::func specifies the kernel to be launched. This same functions must
be launched on all devices. For templated functions, pass the function symbol as follows:
func_name<template_arg_0,...,template_arg_N>
‣ cudaLaunchParams::gridDim specifies the width, height and depth of the grid in blocks. This must
match across all kernels launched.
‣ cudaLaunchParams::blockDim is the width, height and depth of each thread block. This must
match across all kernels launched.
‣ cudaLaunchParams::args specifies the arguments to the kernel. If the kernel has N parameters
then cudaLaunchParams::args should point to array of N pointers. Each pointer, from
cudaLaunchParams::args[0] to cudaLaunchParams::args[N - 1], point to the region
of memory from which the actual parameter will be copied.
‣ cudaLaunchParams::sharedMem is the dynamic shared-memory size per thread block in bytes.
This must match across all kernels launched.
‣ cudaLaunchParams::stream is the handle to the stream to perform the launch in. This cannot be the
NULL stream or cudaStreamLegacy or cudaStreamPerThread.
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
cudaCooperativeLaunchMultiDeviceNoPreSync. 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 cudaCooperativeLaunchMultiDeviceNoPostSync. 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.
|
80
Modules
Note:
‣ This function uses standard default stream semantics.
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaLaunchCooperativeKernel (C++ API), cudaLaunchCooperativeKernel,
cuLaunchCooperativeKernelMultiDevice
__device__cudaError_t cudaLaunchDevice (void
*func, void *parameterBuffer, dim3 gridDimension,
dim3 blockDimension, unsigned int sharedMemSize,
cudaStream_t stream)
Launches a specified kernel.
Parameters
func
- Pointer to the kernel to be launched
parameterBuffer
- Holds the parameters to the launched kernel. parameterBuffer can be NULL. (Optional)
gridDimension
- Specifies grid dimensions
blockDimension
- Specifies block dimensions
sharedMemSize
- Specifies size of shared memory
stream
- Specifies the stream to be used
Returns
cudaSuccess, cudaErrorInvalidDevice, cudaErrorLaunchMaxDepthExceeded,
cudaErrorInvalidConfiguration, cudaErrorStartupFailure, cudaErrorLaunchPendingCountExceeded,
cudaErrorLaunchOutOfResources
|
81
Modules
Description
Launches a specified kernel with the specified parameter buffer. A parameter buffer can be obtained by
calling cudaGetParameterBuffer().
This is a low level API and can only be accessed from Parallel Thread Execution (PTX). CUDA user
code should use <<< >>> to launch the kernels.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
Please refer to Execution Configuration and Parameter Buffer Layout from the CUDA Programming
Guide for the detailed descriptions of launch configuration and parameter layout respectively.
See also:
cudaGetParameterBuffer
__host__cudaError_t cudaLaunchHostFunc (cudaStream_t
stream, cudaHostFn_t fn, void *userData)
Enqueues a host function call in a stream.
Parameters
stream
fn
- The function to call once preceding stream operations are complete
userData
- User-specified data to be passed to the function
Returns
cudaSuccess, cudaErrorInvalidResourceHandle, cudaErrorInvalidValue, cudaErrorNotSupported
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 cudaErrorNotPermitted, 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.
|
82
Modules
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.
‣ 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 constrast 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.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaStreamCreate, cudaStreamQuery, cudaStreamSynchronize, cudaStreamWaitEvent,
cudaStreamDestroy, cudaMallocManaged, cudaStreamAttachMemAsync, cudaStreamAddCallback,
cuLaunchHostFunc
|
83
Modules
__host__cudaError_t cudaLaunchKernel (const void
*func, dim3 gridDim, dim3 blockDim, void **args, size_t
sharedMem, cudaStream_t stream)
Launches a device function.
Parameters
func
- Device function symbol
gridDim
- Grid dimentions
blockDim
- Block dimentions
args
- Arguments
sharedMem
- Shared memory
stream
- Stream identifier
Returns
cudaSuccess, cudaErrorInvalidDeviceFunction, cudaErrorInvalidConfiguration,
cudaErrorLaunchFailure, cudaErrorLaunchTimeout, cudaErrorLaunchOutOfResources,
cudaErrorSharedObjectInitFailed, cudaErrorInvalidPtx, cudaErrorUnsupportedPtxVersion,
cudaErrorNoKernelImageForDevice, cudaErrorJitCompilerNotFound,
cudaErrorJitCompilationDisabled
Description
The function invokes kernel func on gridDim (gridDim.x gridDim.y gridDim.z) grid of
blocks. Each block contains blockDim (blockDim.x blockDim.y blockDim.z) threads.
If the kernel has N parameters the args should point to array of N pointers. Each pointer, from
args[0] to args[N - 1], point to the region of memory from which the actual parameter will be
copied.
For templated functions, pass the function symbol as follows:
func_name<template_arg_0,...,template_arg_N>
sharedMem sets the amount of dynamic shared memory that will be available to each thread block.
stream specifies a stream the invocation is associated to.
|
84
Modules
Note:
‣ This function uses standard default stream semantics.
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
‣ The API can also be used with a kernel cudaKernel_t by querying the handle using
cudaLibraryGetKernel() or cudaGetKernel and then passing it to the API by casting to void*. The
symbol entryFuncAddr passed to cudaGetKernel should be a symbol that is registered with the
same CUDA Runtime instance.
‣ Passing a symbol that belongs that belongs to a different runtime instance will result in undefined
behavior. The only type that can be reliably passed to a different runtime instance is cudaKernel_t
See also:
cudaLaunchKernel (C++ API), cuLaunchKernel
__host__cudaError_t cudaLaunchKernelExC (const
cudaLaunchConfig_t *config, const void *func, void
**args)
Launches a CUDA function with launch-time configuration.
Parameters
config
- Launch configuration
func
- Kernel to launch
args
- Array of pointers to kernel parameters
Returns
cudaSuccess, cudaErrorInvalidDeviceFunction, cudaErrorInvalidConfiguration,
cudaErrorLaunchFailure, cudaErrorLaunchTimeout, cudaErrorLaunchOutOfResources,
cudaErrorSharedObjectInitFailed, cudaErrorInvalidPtx, cudaErrorUnsupportedPtxVersion,
cudaErrorNoKernelImageForDevice, cudaErrorJitCompilerNotFound,
cudaErrorJitCompilationDisabled
|
85
Modules
Description
Note that the functionally equivalent variadic template cudaLaunchKernelEx is available for C++11
and newer.
Invokes the kernel func on config->gridDim (config->gridDim.x config-
>gridDim.y config->gridDim.z) grid of blocks. Each block contains config->blockDim
(config->blockDim.x config->blockDim.y config->blockDim.z) threads.
config->dynamicSmemBytes sets the amount of dynamic shared memory that will be available
to each thread block.
config->stream specifies a stream the invocation is associated to.
Configuration beyond grid and block dimensions, dynamic shared memory size, and stream can be
provided with the following two fields of config:
config->attrs is an array of config->numAttrs contiguous cudaLaunchAttribute elements.
The value of this pointer is not considered if config->numAttrs is zero. However, in that case,
it is recommended to set the pointer to NULL. config->numAttrs is the number of attributes
populating the first config->numAttrs positions of the config->attrs array.
If the kernel has N parameters the args should point to array of N pointers. Each pointer, from
args[0] to args[N - 1], point to the region of memory from which the actual parameter will be
copied.
N.B. This function is so named to avoid unintentionally invoking the templated version,
cudaLaunchKernelEx, for kernels taking a single void** or void* parameter.
Note:
‣ This function uses standard default stream semantics.
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
‣ The API can also be used with a kernel cudaKernel_t by querying the handle using
cudaLibraryGetKernel() or cudaGetKernel and then passing it to the API by casting to void*. The
symbol entryFuncAddr passed to cudaGetKernel should be a symbol that is registered with the
same CUDA Runtime instance.
‣ Passing a symbol that belongs that belongs to a different runtime instance will result in undefined
behavior. The only type that can be reliably passed to a different runtime instance is cudaKernel_t
See also:
|
86
Modules
cudaLaunchKernelEx(const cudaLaunchConfig_t *config, void (*kernel)(ExpTypes...), ActTypes
&&... args) "cudaLaunchKernelEx (C++ API)", cuLaunchKernelEx
__host__cudaError_t cudaSetDoubleForDevice (double *d)
Converts a double argument to be executed on a device.
Parameters
d
- Double to convert
Returns
cudaSuccess
Description
Deprecated This function is deprecated as of CUDA 7.5
Converts the double value of d to an internal float representation if the device does not support double
arithmetic. If the device does natively support doubles, then this function does nothing.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaFuncSetCacheConfig ( C API), cudaFuncGetAttributes ( C API), cudaSetDoubleForHost
__host__cudaError_t cudaSetDoubleForHost (double *d)
Converts a double argument after execution on a device.
Parameters
d
- Double to convert
Returns
cudaSuccess
|
87
Modules
Description
Deprecated This function is deprecated as of CUDA 7.5
Converts the double value of d from a potentially internal float representation if the device does not
support double arithmetic. If the device does natively support doubles, then this function does nothing.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaFuncSetCacheConfig ( C API), cudaFuncGetAttributes ( C API), cudaSetDoubleForDevice
__device__ void
cudaTriggerProgrammaticLaunchCompletion (void)
Programmatic dependency trigger.
Description
This device function ensures the programmatic launch completion edges / events are fulfilled.
See cudaLaunchAttributeID::cudaLaunchAttributeProgrammaticStreamSerialization and
cudaLaunchAttributeID::cudaLaunchAttributeProgrammaticEvent for more information. The event /
edge kick off only happens when every CTAs in the grid has either exited or called this function at
least once, otherwise the kick off happens automatically after all warps finishes execution but before
the grid completes. The kick off only enables scheduling of the secondary kernel. It provides no
memory visibility guarantee itself. The user could enforce memory visibility by inserting a memory
fence of the correct scope.
6.9.
Execution Control [DEPRECATED]
This section describes the deprecated execution control functions of the CUDA runtime application
programming interface.
Some functions have overloaded C++ API template versions documented separately in the C++ API
Routines module.
|
88
Modules
__host__cudaError_t cudaFuncSetSharedMemConfig
(const void *func, cudaSharedMemConfig config)
Sets the shared memory configuration for a device function.
Parameters
func
- Device function symbol
config
- Requested shared memory configuration
Returns
cudaSuccess, cudaErrorInvalidDeviceFunction, cudaErrorInvalidValue,
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 cudaFuncSetSharedMemConfig will
override the device wide setting set by cudaDeviceSetSharedMemConfig.
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.
For templated functions, pass the function symbol as follows:
func_name<template_arg_0,...,template_arg_N>
The supported bank configurations are:
‣ cudaSharedMemBankSizeDefault: use the device's shared memory configuration when launching
this function.
‣ cudaSharedMemBankSizeFourByte: set shared memory bank width to be four bytes natively when
launching this function.
‣ cudaSharedMemBankSizeEightByte: set shared memory bank width to be eight bytes natively
when launching this function.
|
89
Modules
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Use of a string naming a function as the func parameter was deprecated in CUDA 4.1 and
removed in CUDA 5.0.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaDeviceSetSharedMemConfig, cudaDeviceGetSharedMemConfig, cudaDeviceSetCacheConfig,
cudaDeviceGetCacheConfig, cudaFuncSetCacheConfig, cuFuncSetSharedMemConfig
6.10. Occupancy
This section describes the occupancy calculation functions of the CUDA runtime application
programming interface.
Besides the occupancy calculator functions (cudaOccupancyMaxActiveBlocksPerMultiprocessor and
cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags), there are also C++ only occupancy-
based launch configuration functions documented in C++ API Routines module.
See cudaOccupancyMaxPotentialBlockSize ( C++ API), cudaOccupancyMaxPotentialBlockSize
( C++ API), cudaOccupancyMaxPotentialBlockSizeVariableSMem ( C+
+ API), cudaOccupancyMaxPotentialBlockSizeVariableSMem ( C++ API)
cudaOccupancyAvailableDynamicSMemPerBlock (C++ API),
__host__cudaError_t
cudaOccupancyAvailableDynamicSMemPerBlock (size_t
*dynamicSmemSize, const void *func, int numBlocks, int
blockSize)
Returns dynamic shared memory available per block when launching numBlocks blocks on SM.
Parameters
dynamicSmemSize
- Returned maximum dynamic shared memory
func
- Kernel function for which occupancy is calculated
|
90
Modules
numBlocks
- Number of blocks to fit on SM
blockSize
- Size of the block
Returns
cudaSuccess, cudaErrorInvalidDevice, cudaErrorInvalidDeviceFunction, cudaErrorInvalidValue,
cudaErrorUnknown,
Description
Returns in *dynamicSmemSize the maximum size of dynamic shared memory to allow
numBlocks blocks per SM.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
‣ The API can also be used with a kernel cudaKernel_t by querying the handle using
cudaLibraryGetKernel() or cudaGetKernel and then passing it to the API by casting to void*. The
symbol entryFuncAddr passed to cudaGetKernel should be a symbol that is registered with the
same CUDA Runtime instance.
‣ Passing a symbol that belongs that belongs to a different runtime instance will result in undefined
behavior. The only type that can be reliably passed to a different runtime instance is cudaKernel_t
See also:
cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags,
cudaOccupancyMaxPotentialBlockSize ( C++ API), cudaOccupancyMaxPotentialBlockSizeWithFlags
( C++ API), cudaOccupancyMaxPotentialBlockSizeVariableSMem ( C++ API),
cudaOccupancyMaxPotentialBlockSizeVariableSMemWithFlags ( C++ API),
cudaOccupancyAvailableDynamicSMemPerBlock
|
91
Modules
__host____
device__cudaError_t
cudaOccupancyMaxActiveBlocksPerMultiprocessor
(int *numBlocks, const void *func, int blockSize, size_t
dynamicSMemSize)
Returns occupancy for a device function.
Parameters
numBlocks
- Returned occupancy
func
- Kernel function for which occupancy is calculated
blockSize
- Block size the kernel is intended to be launched with
dynamicSMemSize
- Per-block dynamic shared memory usage intended, in bytes
Returns
cudaSuccess, cudaErrorInvalidDevice, cudaErrorInvalidDeviceFunction, cudaErrorInvalidValue,
cudaErrorUnknown,
Description
Returns in *numBlocks the maximum number of active blocks per streaming multiprocessor for the
device function.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
‣ The API can also be used with a kernel cudaKernel_t by querying the handle using
cudaLibraryGetKernel() or cudaGetKernel and then passing it to the API by casting to void*. The
symbol entryFuncAddr passed to cudaGetKernel should be a symbol that is registered with the
same CUDA Runtime instance.
‣ Passing a symbol that belongs that belongs to a different runtime instance will result in undefined
behavior. The only type that can be reliably passed to a different runtime instance is cudaKernel_t
|
92
Modules
See also:
cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags,
cudaOccupancyMaxPotentialBlockSize ( C++ API), cudaOccupancyMaxPotentialBlockSizeWithFlags
( C++ API), cudaOccupancyMaxPotentialBlockSizeVariableSMem ( C+
+ API), cudaOccupancyMaxPotentialBlockSizeVariableSMemWithFlags
( C++ API), cudaOccupancyAvailableDynamicSMemPerBlock (C++ API),
cuOccupancyMaxActiveBlocksPerMultiprocessor
__host__cudaError_t
cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags
(int *numBlocks, const void *func, int blockSize, size_t
dynamicSMemSize, unsigned int flags)
Returns occupancy for a device function with the specified flags.
Parameters
numBlocks
- Returned occupancy
func
- Kernel function for which occupancy is calculated
blockSize
- Block size the kernel is intended to be launched with
dynamicSMemSize
- Per-block dynamic shared memory usage intended, in bytes
flags
- Requested behavior for the occupancy calculator
Returns
cudaSuccess, cudaErrorInvalidDevice, cudaErrorInvalidDeviceFunction, cudaErrorInvalidValue,
cudaErrorUnknown,
Description
Returns in *numBlocks the maximum number of active blocks per streaming multiprocessor for the
device function.
The flags parameter controls how special cases are handled. Valid flags include:
‣ cudaOccupancyDefault: keeps the default behavior as
cudaOccupancyMaxActiveBlocksPerMultiprocessor
‣ cudaOccupancyDisableCachingOverride: This flag suppresses the default behavior on platform
where global caching affects occupancy. On such platforms, if caching is enabled, but per-block
|
93
Modules
SM resource usage would result in zero occupancy, the occupancy calculator will calculate the
occupancy as if caching is disabled. Setting this flag makes the occupancy calculator to return 0 in
such cases. More information can be found about this feature in the "Unified L1/Texture Cache"
section of the Maxwell tuning guide.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
‣ The API can also be used with a kernel cudaKernel_t by querying the handle using
cudaLibraryGetKernel() or cudaGetKernel and then passing it to the API by casting to void*. The
symbol entryFuncAddr passed to cudaGetKernel should be a symbol that is registered with the
same CUDA Runtime instance.
‣ Passing a symbol that belongs that belongs to a different runtime instance will result in undefined
behavior. The only type that can be reliably passed to a different runtime instance is cudaKernel_t
See also:
cudaOccupancyMaxActiveBlocksPerMultiprocessor, cudaOccupancyMaxPotentialBlockSize
( C++ API), cudaOccupancyMaxPotentialBlockSizeWithFlags ( C++
API), cudaOccupancyMaxPotentialBlockSizeVariableSMem ( C++ API),
cudaOccupancyMaxPotentialBlockSizeVariableSMemWithFlags ( C+
+ API), cudaOccupancyAvailableDynamicSMemPerBlock (C++ API),
cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags
__host__cudaError_t cudaOccupancyMaxActiveClusters
(int *numClusters, const void *func, const
cudaLaunchConfig_t *launchConfig)
Given the kernel function (func) and launch configuration (config), return the maximum number of
clusters that could co-exist on the target device in *numClusters.
Parameters
numClusters
- Returned maximum number of clusters that could co-exist on the target device
func
- Kernel function for which maximum number of clusters are calculated
|
94
Modules
launchConfig
Returns
cudaSuccess, cudaErrorInvalidDeviceFunction, cudaErrorInvalidValue, cudaErrorInvalidClusterSize,
cudaErrorUnknown,
Description
If the function has required cluster size already set (see cudaFuncGetAttributes), the cluster size from
config must either be unspecified or match the required size. Without required sizes, the cluster size
must be specified in config, else the function will return an error.
Note that various attributes of the kernel function may affect occupancy calculation. Runtime
environment may affect how the hardware schedules the clusters, so the calculated occupancy is not
guaranteed to be achievable.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
‣ The API can also be used with a kernel cudaKernel_t by querying the handle using
cudaLibraryGetKernel() or cudaGetKernel and then passing it to the API by casting to void*. The
symbol entryFuncAddr passed to cudaGetKernel should be a symbol that is registered with the
same CUDA Runtime instance.
‣ Passing a symbol that belongs that belongs to a different runtime instance will result in undefined
behavior. The only type that can be reliably passed to a different runtime instance is cudaKernel_t
See also:
cudaFuncGetAttributes cudaOccupancyMaxActiveClusters (C++ API),
cuOccupancyMaxActiveClusters
|
95
Modules
__host__cudaError_t
cudaOccupancyMaxPotentialClusterSize (int
*clusterSize, const void *func, const cudaLaunchConfig_t
*launchConfig)
Given the kernel function (func) and launch configuration (config), return the maximum cluster
size in *clusterSize.
Parameters
clusterSize
- Returned maximum cluster size that can be launched for the given kernel function and launch
configuration
func
- Kernel function for which maximum cluster size is calculated
launchConfig
Returns
cudaSuccess, cudaErrorInvalidDeviceFunction, cudaErrorInvalidValue, cudaErrorUnknown,
Description
The cluster dimensions in config are ignored. If func has a required cluster size set (see
cudaFuncGetAttributes),*clusterSize will reflect the required cluster size.
By default this function will always return a value that's portable on future hardware. A higher value
may be returned if the kernel function allows non-portable cluster sizes.
This function will respect the compile time launch bounds.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
‣ The API can also be used with a kernel cudaKernel_t by querying the handle using
cudaLibraryGetKernel() or cudaGetKernel and then passing it to the API by casting to void*. The
symbol entryFuncAddr passed to cudaGetKernel should be a symbol that is registered with the
same CUDA Runtime instance.
|
96
Modules
‣ Passing a symbol that belongs that belongs to a different runtime instance will result in undefined
behavior. The only type that can be reliably passed to a different runtime instance is cudaKernel_t
See also:
cudaFuncGetAttributes cudaOccupancyMaxPotentialClusterSize (C++ API),
cuOccupancyMaxPotentialClusterSize
6.11. Memory Management
This section describes the memory management functions of the CUDA runtime application
programming interface.
Some functions have overloaded C++ API template versions documented separately in the C++ API
Routines module.
__host__cudaError_t cudaArrayGetInfo
(cudaChannelFormatDesc *desc, cudaExtent *extent,
unsigned int *flags, cudaArray_t array)
Gets info about the specified cudaArray.
Parameters
desc
- Returned array type
extent
- Returned array shape. 2D arrays will have depth of zero
flags
- Returned array flags
array
- The cudaArray to get info for
Returns
cudaSuccess, cudaErrorInvalidValue
Description
Returns in *desc, *extent and *flags respectively, the type, shape and flags of array.
Any of *desc, *extent and *flags may be specified as NULL.
Note:
|
97
Modules
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cuArrayGetDescriptor, cuArray3DGetDescriptor
__host__cudaError_t cudaArrayGetMemoryRequirements
(cudaArrayMemoryRequirements *memoryRequirements,
cudaArray_t array, int device)
Returns the memory requirements of a CUDA array.
Parameters
memoryRequirements
- Pointer to cudaArrayMemoryRequirements
array
- CUDA array to get the memory requirements of
device
- Device to get the memory requirements for
Returns
cudaSuccess cudaErrorInvalidValue
Description
Returns the memory requirements of a CUDA array in memoryRequirements If the CUDA array
is not allocated with flag cudaArrayDeferredMapping cudaErrorInvalidValue will be returned.
The returned value in cudaArrayMemoryRequirements::size represents the total size of the CUDA
array. The returned value in cudaArrayMemoryRequirements::alignment represents the alignment
necessary for mapping the CUDA array.
See also:
cudaMipmappedArrayGetMemoryRequirements
|
98
Modules
__host__cudaError_t cudaArrayGetPlane (cudaArray_t
*pPlaneArray, cudaArray_t hArray, unsigned int planeIdx)
Gets a CUDA array plane from a CUDA array.
Parameters
pPlaneArray
- Returned CUDA array referenced by the planeIdx
hArray
- CUDA array
planeIdx
- Plane index
Returns
cudaSuccess, cudaErrorInvalidValue cudaErrorInvalidResourceHandle
Description
Returns in pPlaneArray a CUDA array that represents a single format plane of the CUDA array
hArray.
If planeIdx is greater than the maximum number of planes in this array or if the array does not have
a multi-planar format e.g: cudaChannelFormatKindNV12, then cudaErrorInvalidValue is returned.
Note that if the hArray has format cudaChannelFormatKindNV12, then passing in 0 for
planeIdx returns a CUDA array of the same size as hArray but with one 8-bit channel and
cudaChannelFormatKindUnsigned as its format kind. If 1 is passed for planeIdx, then the
returned CUDA array has half the height and width of hArray with two 8-bit channels and
cudaChannelFormatKindUnsigned as its format kind.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuArrayGetPlane
|
99
Modules
__host__cudaError_t cudaArrayGetSparseProperties
(cudaArraySparseProperties *sparseProperties,
cudaArray_t array)
Returns the layout properties of a sparse CUDA array.
Parameters
sparseProperties
- Pointer to return the cudaArraySparseProperties
array
- The CUDA array to get the sparse properties of
Returns
cudaSuccess cudaErrorInvalidValue
Description
Returns the layout properties of a sparse CUDA array in sparseProperties. If the CUDA array is
not allocated with flag cudaArraySparse cudaErrorInvalidValue will be returned.
If the returned value in cudaArraySparseProperties::flags contains
cudaArraySparsePropertiesSingleMipTail, then cudaArraySparseProperties::miptailSize
represents the total size of the array. Otherwise, it will be zero. Also, the returned value in
cudaArraySparseProperties::miptailFirstLevel is always zero. Note that the array must have
been allocated using cudaMallocArray or cudaMalloc3DArray. For CUDA arrays obtained
using cudaMipmappedArrayGetLevel, cudaErrorInvalidValue will be returned. Instead,
cudaMipmappedArrayGetSparseProperties must be used to obtain the sparse properties of the entire
CUDA mipmapped array to which array belongs to.
See also:
cudaMipmappedArrayGetSparseProperties, cuMemMapArrayAsync
__host____
device__cudaError_t cudaFree (void *devPtr)
Frees memory on the device.
Parameters
devPtr
- Device pointer to memory to free
Returns
cudaSuccess, cudaErrorInvalidValue
|
100
Modules
Description
Frees the memory space pointed to by devPtr, which must have been returned by a previous
call to one of the following memory allocation APIs - cudaMalloc(), cudaMallocPitch(),
cudaMallocManaged(), cudaMallocAsync(), cudaMallocFromPoolAsync().
Note - This API will not perform any implicit synchronization when the pointer was allocated with
cudaMallocAsync or cudaMallocFromPoolAsync. Callers must ensure that all accesses to these pointer
have completed before invoking cudaFree. For best performance and memory reuse, users should
use cudaFreeAsync to free memory allocated via the stream ordered memory allocator. For all other
pointers, this API may perform implicit synchronization.
If cudaFree(devPtr) has already been called before, an error is returned. If devPtr is 0, no operation
is performed. cudaFree() returns cudaErrorValue in case of failure.
The device version of cudaFree cannot be used with a *devPtr allocated using the host API, and vice
versa.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaMalloc, cudaMallocPitch, cudaMallocManaged, cudaMallocArray, cudaFreeArray,
cudaMallocAsync, cudaMallocFromPoolAsync cudaMallocHost ( C API), cudaFreeHost,
cudaMalloc3D, cudaMalloc3DArray, cudaFreeAsync cudaHostAlloc, cuMemFree
__host__cudaError_t cudaFreeArray (cudaArray_t array)
Frees an array on the device.
Parameters
array
- Pointer to array to free
Returns
cudaSuccess, cudaErrorInvalidValue
|
101
Modules
Description
Frees the CUDA array array, which must have been returned by a previous call to
cudaMallocArray(). If devPtr is 0, no operation is performed.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaMalloc, cudaMallocPitch, cudaFree, cudaMallocArray, cudaMallocHost ( C API), cudaFreeHost,
cudaHostAlloc, cuArrayDestroy
__host__cudaError_t cudaFreeHost (void *ptr)
Frees page-locked memory.
Parameters
ptr
- Pointer to memory to free
Returns
cudaSuccess, cudaErrorInvalidValue
Description
Frees the memory space pointed to by hostPtr, which must have been returned by a previous call to
cudaMallocHost() or cudaHostAlloc().
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
|
102
Modules
See also:
cudaMalloc, cudaMallocPitch, cudaFree, cudaMallocArray, cudaFreeArray, cudaMallocHost ( C API),
cudaMalloc3D, cudaMalloc3DArray, cudaHostAlloc, cuMemFreeHost
__host__cudaError_t cudaFreeMipmappedArray
(cudaMipmappedArray_t mipmappedArray)
Frees a mipmapped array on the device.
Parameters
mipmappedArray
- Pointer to mipmapped array to free
Returns
cudaSuccess, cudaErrorInvalidValue
Description
Frees the CUDA mipmapped array mipmappedArray, which must have been returned by a previous
call to cudaMallocMipmappedArray(). If devPtr is 0, no operation is performed.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaMalloc, cudaMallocPitch, cudaFree, cudaMallocArray, cudaMallocHost ( C API), cudaFreeHost,
cudaHostAlloc, cuMipmappedArrayDestroy
|
103
Modules
__host__cudaError_t cudaGetMipmappedArrayLevel
(cudaArray_t *levelArray, cudaMipmappedArray_const_t
mipmappedArray, unsigned int level)
Gets a mipmap level of a CUDA mipmapped array.
Parameters
levelArray
- Returned mipmap level CUDA array
mipmappedArray
- CUDA mipmapped array
level
- Mipmap level
Returns
cudaSuccess, cudaErrorInvalidValue cudaErrorInvalidResourceHandle
Description
Returns in *levelArray a CUDA array that represents a single mipmap level of the CUDA
mipmapped array mipmappedArray.
If level is greater than the maximum number of levels in this mipmapped array,
cudaErrorInvalidValue is returned.
If mipmappedArray is NULL, cudaErrorInvalidResourceHandle is returned.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaMalloc3D, cudaMalloc, cudaMallocPitch, cudaFree, cudaFreeArray, cudaMallocHost ( C API),
cudaFreeHost, cudaHostAlloc, make_cudaExtent, cuMipmappedArrayGetLevel
|
104
Modules
__host__cudaError_t cudaGetSymbolAddress (void
**devPtr, const void *symbol)
Finds the address associated with a CUDA symbol.
Parameters
devPtr
- Return device pointer associated with symbol
symbol
- Device symbol address
Returns
cudaSuccess, cudaErrorInvalidSymbol, cudaErrorNoKernelImageForDevice
Description
Returns in *devPtr the address of symbol symbol on the device. symbol is a variable that resides
in global or constant memory space. If symbol cannot be found, or if symbol is not declared in the
global or constant memory space, *devPtr is unchanged and the error cudaErrorInvalidSymbol is
returned.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Use of a string naming a variable as the symbol parameter was deprecated in CUDA 4.1 and
removed in CUDA 5.0.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaGetSymbolAddress ( C++ API), cudaGetSymbolSize ( C API), cuModuleGetGlobal
|
105
Modules
__host__cudaError_t cudaGetSymbolSize (size_t *size,
const void *symbol)
Finds the size of the object associated with a CUDA symbol.
Parameters
size
- Size of object associated with symbol
symbol
- Device symbol address
Returns
cudaSuccess, cudaErrorInvalidSymbol, cudaErrorNoKernelImageForDevice
Description
Returns in *size the size of symbol symbol. symbol is a variable that resides in global or constant
memory space. If symbol cannot be found, or if symbol is not declared in global or constant
memory space, *size is unchanged and the error cudaErrorInvalidSymbol is returned.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Use of a string naming a variable as the symbol parameter was deprecated in CUDA 4.1 and
removed in CUDA 5.0.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaGetSymbolAddress ( C API), cudaGetSymbolSize ( C++ API), cuModuleGetGlobal
|
106
Modules
__host__cudaError_t cudaHostAlloc (void **pHost, size_t
size, unsigned int flags)
Allocates page-locked memory on the host.
Parameters
pHost
- Device pointer to allocated memory
size
- Requested allocation size in bytes
flags
- Requested properties of allocated memory
Returns
cudaSuccess, cudaErrorInvalidValue, cudaErrorMemoryAllocation
Description
Allocates size bytes of host memory that is page-locked and accessible to the device. The driver
tracks the virtual memory ranges allocated with this function and automatically accelerates calls to
functions such as cudaMemcpy(). Since the memory can be accessed directly by the device, it can be
read or written with much higher bandwidth than pageable memory obtained with functions such as
malloc(). Allocating excessive amounts of pinned memory may degrade system performance, since it
reduces the amount of memory available to the system for paging. As a result, this function is best used
sparingly to allocate staging areas for data exchange between host and device.
The flags parameter enables different options to be specified that affect the allocation, as follows.
‣ cudaHostAllocDefault: This flag's value is defined to be 0 and causes cudaHostAlloc() to emulate
cudaMallocHost().
‣ cudaHostAllocPortable: The memory returned by this call will be considered as pinned memory by
all CUDA contexts, not just the one that performed the allocation.
‣ cudaHostAllocMapped: Maps the allocation into the CUDA address space. The device pointer to
the memory may be obtained by calling cudaHostGetDevicePointer().
‣ cudaHostAllocWriteCombined: Allocates the memory as write-combined (WC). WC memory can
be transferred across the PCI Express bus more quickly on some system configurations, but cannot
be read efficiently by most CPUs. WC memory is a good option for buffers that will be written by
the CPU and read by the device via mapped pinned memory or host->device transfers.
All of these flags are orthogonal to one another: a developer may allocate memory that is portable,
mapped and/or write-combined with no restrictions.
|
107
Modules
In order for the cudaHostAllocMapped flag to have any effect, the CUDA context must support the
cudaDeviceMapHost flag, which can be checked via cudaGetDeviceFlags(). The cudaDeviceMapHost
flag is implicitly set for contexts created via the runtime API.
The cudaHostAllocMapped flag may be specified on CUDA contexts for devices that do not support
mapped pinned memory. The failure is deferred to cudaHostGetDevicePointer() because the memory
may be mapped into other CUDA contexts via the cudaHostAllocPortable flag.
Memory allocated by this function must be freed with cudaFreeHost().
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaSetDeviceFlags, cudaMallocHost ( C API), cudaFreeHost, cudaGetDeviceFlags, cuMemHostAlloc
__host__cudaError_t cudaHostGetDevicePointer (void
**pDevice, void *pHost, unsigned int flags)
Passes back device pointer of mapped host memory allocated by cudaHostAlloc or registered by
cudaHostRegister.
Parameters
pDevice
- Returned device pointer for mapped memory
pHost
- Requested host pointer mapping
flags
- Flags for extensions (must be 0 for now)
Returns
cudaSuccess, cudaErrorInvalidValue, cudaErrorMemoryAllocation
Description
Passes back the device pointer corresponding to the mapped, pinned host buffer allocated by
cudaHostAlloc() or registered by cudaHostRegister().
|
108
Modules
cudaHostGetDevicePointer() will fail if the cudaDeviceMapHost flag was not specified before deferred
context creation occurred, or if called on a device that does not support mapped, pinned memory.
For devices that have a non-zero value for the device attribute
cudaDevAttrCanUseHostPointerForRegisteredMem, the memory can also be accessed from the
device using the host pointer pHost. The device pointer returned by cudaHostGetDevicePointer()
may or may not match the original host pointer pHost and depends on the devices visible to the
application. If all devices visible to the application have a non-zero value for the device attribute, the
device pointer returned by cudaHostGetDevicePointer() will match the original pointer pHost. If any
device visible to the application has a zero value for the device attribute, the device pointer returned by
cudaHostGetDevicePointer() will not match the original host pointer pHost, but it will be suitable for
use on all devices provided Unified Virtual Addressing is enabled. In such systems, it is valid to access
the memory using either pointer on devices that have a non-zero value for the device attribute. Note
however that such devices should access the memory using only of the two pointers and not both.
flags provides for future releases. For now, it must be set to 0.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaSetDeviceFlags, cudaHostAlloc, cuMemHostGetDevicePointer
__host__cudaError_t cudaHostGetFlags (unsigned int
*pFlags, void *pHost)
Passes back flags used to allocate pinned host memory allocated by cudaHostAlloc.
Parameters
pFlags
- Returned flags word
pHost
- Host pointer
Returns
cudaSuccess, cudaErrorInvalidValue
|
109
Modules
Description
cudaHostGetFlags() will fail if the input pointer does not reside in an address range allocated by
cudaHostAlloc().
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaHostAlloc, cuMemHostGetFlags
__host__cudaError_t cudaHostRegister (void *ptr, size_t
size, unsigned int flags)
Registers an existing host memory range for use by CUDA.
Parameters
ptr
- Host pointer to memory to page-lock
size
- Size in bytes of the address range to page-lock in bytes
flags
- Flags for allocation request
Returns
cudaSuccess, cudaErrorInvalidValue, cudaErrorMemoryAllocation,
cudaErrorHostMemoryAlreadyRegistered, cudaErrorNotSupported
Description
Page-locks the memory range specified by ptr and size and maps it for the device(s) as specified
by flags. This memory range also is added to the same tracking mechanism as cudaHostAlloc() to
automatically accelerate calls to functions such as cudaMemcpy(). Since the memory can be accessed
directly by the device, it can be read or written with much higher bandwidth than pageable memory that
has not been registered. Page-locking excessive amounts of memory may degrade system performance,
|
110
Modules
since it reduces the amount of memory available to the system for paging. As a result, this function is
best used sparingly to register staging areas for data exchange between host and device.
On systems where pageableMemoryAccessUsesHostPageTables is true, cudaHostRegister will not
page-lock the memory range specified by ptr but only populate unpopulated pages.
cudaHostRegister is supported only on I/O coherent devices that have a non-zero value for the device
attribute cudaDevAttrHostRegisterSupported.
The flags parameter enables different options to be specified that affect the allocation, as follows.
‣ cudaHostRegisterDefault: On a system with unified virtual addressing, the memory will be both
mapped and portable. On a system with no unified virtual addressing, the memory will be neither
mapped nor portable.
‣ cudaHostRegisterPortable: The memory returned by this call will be considered as pinned memory
by all CUDA contexts, not just the one that performed the allocation.
‣ cudaHostRegisterMapped: Maps the allocation into the CUDA address space. The device pointer to
the memory may be obtained by calling cudaHostGetDevicePointer().
‣ cudaHostRegisterIoMemory: The passed memory pointer is treated as pointing to some memory-
mapped I/O space, e.g. belonging to a third-party PCIe device, and it will marked as non cache-
coherent and contiguous.
‣ cudaHostRegisterReadOnly: The passed memory pointer is treated as pointing
to memory that is considered read-only by the device. On platforms without
cudaDevAttrPageableMemoryAccessUsesHostPageTables, this flag is required in order to register
memory mapped to the CPU as read-only. Support for the use of this flag can be queried from
the device attribute cudaDevAttrHostRegisterReadOnlySupported. Using this flag with a current
context associated with a device that does not have this attribute set will cause cudaHostRegister to
error with cudaErrorNotSupported.
All of these flags are orthogonal to one another: a developer may page-lock memory that is portable or
mapped with no restrictions.
The CUDA context must have been created with the cudaMapHost flag in order for the
cudaHostRegisterMapped flag to have any effect.
The cudaHostRegisterMapped flag may be specified on CUDA contexts for devices that do not support
mapped pinned memory. The failure is deferred to cudaHostGetDevicePointer() because the memory
may be mapped into other CUDA contexts via the cudaHostRegisterPortable flag.
For devices that have a non-zero value for the device attribute
cudaDevAttrCanUseHostPointerForRegisteredMem, the memory can also be accessed from the device
using the host pointer ptr. The device pointer returned by cudaHostGetDevicePointer() may or
may not match the original host pointer ptr and depends on the devices visible to the application.
If all devices visible to the application have a non-zero value for the device attribute, the device
pointer returned by cudaHostGetDevicePointer() will match the original pointer ptr. If any device
visible to the application has a zero value for the device attribute, the device pointer returned by
|
111
Modules
cudaHostGetDevicePointer() will not match the original host pointer ptr, but it will be suitable for
use on all devices provided Unified Virtual Addressing is enabled. In such systems, it is valid to access
the memory using either pointer on devices that have a non-zero value for the device attribute. Note
however that such devices should access the memory using only of the two pointers and not both.
The memory page-locked by this function must be unregistered with cudaHostUnregister().
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaHostUnregister, cudaHostGetFlags, cudaHostGetDevicePointer, cuMemHostRegister
__host__cudaError_t cudaHostUnregister (void *ptr)
Unregisters a memory range that was registered with cudaHostRegister.
Parameters
ptr
- Host pointer to memory to unregister
Returns
cudaSuccess, cudaErrorInvalidValue, cudaErrorHostMemoryNotRegistered
Description
Unmaps the memory range whose base address is specified by ptr, and makes it pageable again.
The base address must be the same one specified to cudaHostRegister().
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
|
112
Modules
See also:
cudaHostUnregister, cuMemHostUnregister
__host____
device__cudaError_t cudaMalloc (void
**devPtr, size_t size)
Allocate memory on the device.
Parameters
devPtr
- Pointer to allocated device memory
size
- Requested allocation size in bytes
Returns
cudaSuccess, cudaErrorInvalidValue, cudaErrorMemoryAllocation
Description
Allocates size bytes of linear memory on the device and returns in *devPtr a pointer to the
allocated memory. The allocated memory is suitably aligned for any kind of variable. The memory is
not cleared. cudaMalloc() returns cudaErrorMemoryAllocation in case of failure.
The device version of cudaFree cannot be used with a *devPtr allocated using the host API, and vice
versa.
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaMallocPitch, cudaFree, cudaMallocArray, cudaFreeArray, cudaMalloc3D, cudaMalloc3DArray,
cudaMallocHost ( C API), cudaFreeHost, cudaHostAlloc, cuMemAlloc
|
113
Modules
__host__cudaError_t cudaMalloc3D (cudaPitchedPtr
*pitchedDevPtr, cudaExtent extent)
Allocates logical 1D, 2D, or 3D memory objects on the device.
Parameters
pitchedDevPtr
- Pointer to allocated pitched device memory
extent
- Requested allocation size (width field in bytes)
Returns
cudaSuccess, cudaErrorInvalidValue, cudaErrorMemoryAllocation
Description
Allocates at least width * height * depth bytes of linear memory on the device and returns
a cudaPitchedPtr in which ptr is a pointer to the allocated memory. The function may pad the
allocation to ensure hardware alignment requirements are met. The pitch returned in the pitch field of
pitchedDevPtr is the width in bytes of the allocation.
The returned cudaPitchedPtr contains additional fields xsize and ysize, the logical width and
height of the allocation, which are equivalent to the width and height extent parameters
provided by the programmer during allocation.
For allocations of 2D and 3D objects, it is highly recommended that programmers perform allocations
using cudaMalloc3D() or cudaMallocPitch(). Due to alignment restrictions in the hardware, this
is especially true if the application will be performing memory copies involving 2D or 3D objects
(whether linear memory or CUDA arrays).
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaMallocPitch, cudaFree, cudaMemcpy3D, cudaMemset3D, cudaMalloc3DArray, cudaMallocArray,
cudaFreeArray, cudaMallocHost ( C API), cudaFreeHost, cudaHostAlloc, make_cudaPitchedPtr,
make_cudaExtent, cuMemAllocPitch
|
114
Modules
__host__cudaError_t cudaMalloc3DArray (cudaArray_t
*array, const cudaChannelFormatDesc *desc, cudaExtent
extent, unsigned int flags)
Allocate an array on the device.
Parameters
array
- Pointer to allocated array in device memory
desc
- Requested channel format
extent
- Requested allocation size (width field in elements)
flags
- Flags for extensions
Returns
cudaSuccess, cudaErrorInvalidValue, cudaErrorMemoryAllocation
Description
Allocates a CUDA array according to the cudaChannelFormatDesc structure desc and returns a
handle to the new CUDA array in *array.
The cudaChannelFormatDesc is defined as:
struct cudaChannelFormatDesc {
int x, y, z, w;
enum cudaChannelFormatKind
f;
};
where cudaChannelFormatKind is one of cudaChannelFormatKindSigned,
cudaChannelFormatKindUnsigned, or cudaChannelFormatKindFloat.
cudaMalloc3DArray() can allocate the following:
‣ A 1D array is allocated if the height and depth extents are both zero.
‣ A 2D array is allocated if only the depth extent is zero.
‣ A 3D array is allocated if all three extents are non-zero.
‣ A 1D layered CUDA array is allocated if only the height extent is zero and the cudaArrayLayered
flag is set. Each layer is a 1D array. The number of layers is determined by the depth extent.
‣ A 2D layered CUDA array is allocated if all three extents are non-zero and the cudaArrayLayered
flag is set. Each layer is a 2D array. The number of layers is determined by the depth extent.
‣ A cubemap CUDA array is allocated if all three extents are non-zero and the cudaArrayCubemap
flag is set. Width must be equal to height, and depth must be six. A cubemap is a special type of
|
115
Modules
2D layered CUDA array, where the six layers represent the six faces of a cube. The order of the six
layers in memory is the same as that listed in cudaGraphicsCubeFace.
‣ A cubemap layered CUDA array is allocated if all three extents are non-zero, and both,
cudaArrayCubemap and cudaArrayLayered flags are set. Width must be equal to height, and depth
must be a multiple of six. A cubemap layered CUDA array is a special type of 2D layered CUDA
array that consists of a collection of cubemaps. The first six layers represent the first cubemap, the
next six layers form the second cubemap, and so on.
The flags parameter enables different options to be specified that affect the allocation, as follows.
‣ cudaArrayDefault: This flag's value is defined to be 0 and provides default array allocation
‣ cudaArrayLayered: Allocates a layered CUDA array, with the depth extent indicating the number
of layers
‣ cudaArrayCubemap: Allocates a cubemap CUDA array. Width must be equal to height, and depth
must be six. If the cudaArrayLayered flag is also set, depth must be a multiple of six.
‣ cudaArraySurfaceLoadStore: Allocates a CUDA array that could be read from or written to using a
surface reference.
‣ cudaArrayTextureGather: This flag indicates that texture gather operations will be performed on
the CUDA array. Texture gather can only be performed on 2D CUDA arrays.
‣ cudaArraySparse: Allocates a CUDA array without physical backing memory. The subregions
within this sparse array can later be mapped onto a physical memory allocation by calling
cuMemMapArrayAsync. This flag can only be used for creating 2D, 3D or 2D layered sparse
CUDA arrays. The physical backing memory must be allocated via cuMemCreate.
‣ cudaArrayDeferredMapping: Allocates a CUDA array without physical backing memory.
The entire array can later be mapped onto a physical memory allocation by calling
cuMemMapArrayAsync. The physical backing memory must be allocated via cuMemCreate.
The width, height and depth extents must meet certain size requirements as listed in the following table.
All values are specified in elements.
Note that 2D CUDA arrays have different size requirements if the cudaArrayTextureGather flag
is set. In that case, the valid range for (width, height, depth) is ((1,maxTexture2DGather[0]),
(1,maxTexture2DGather[1]), 0).
Valid extents with
Valid extents that must always be met
cudaArraySurfaceLoadStore set
CUDA
{(width range in elements), (height
{(width range in elements), (height
array type
range), (depth range)}
range), (depth range)}
1D
{ (1,maxTexture1D), 0, 0 }
{ (1,maxSurface1D), 0, 0 }
2D
{ (1,maxTexture2D[0]), (1,maxTexture2D[1]),
{ (1,maxSurface2D[0]), (1,maxSurface2D[1]),
0 }
0 }
3D
{ (1,maxTexture3D[0]), (1,maxTexture3D[1]),
{ (1,maxSurface3D[0]), (1,maxSurface3D[1]),
(1,maxTexture3D[2]) } OR
(1,maxSurface3D[2]) }
{ (1,maxTexture3DAlt[0]),
|
116
Modules
Valid extents with
Valid extents that must always be met
cudaArraySurfaceLoadStore set
CUDA
{(width range in elements), (height
{(width range in elements), (height
array type
range), (depth range)}
range), (depth range)}
(1,maxTexture3DAlt[1]),
(1,maxTexture3DAlt[2]) }
1D Layered
{ (1,maxTexture1DLayered[0]), 0,
{ (1,maxSurface1DLayered[0]), 0,
(1,maxTexture1DLayered[1]) }
(1,maxSurface1DLayered[1]) }
2D Layered
{ (1,maxTexture2DLayered[0]),
{ (1,maxSurface2DLayered[0]),
(1,maxTexture2DLayered[1]),
(1,maxSurface2DLayered[1]),
(1,maxTexture2DLayered[2]) }
(1,maxSurface2DLayered[2]) }
Cubemap
{ (1,maxTextureCubemap),
{ (1,maxSurfaceCubemap),
(1,maxTextureCubemap), 6 }
(1,maxSurfaceCubemap), 6 }
Cubemap
{ (1,maxTextureCubemapLayered[0]),
{ (1,maxSurfaceCubemapLayered[0]),
Layered
(1,maxTextureCubemapLayered[0]),
(1,maxSurfaceCubemapLayered[0]),
(1,maxTextureCubemapLayered[1]) }
(1,maxSurfaceCubemapLayered[1]) }
Note:
‣ Note that this function may also return error codes from previous, asynchronous launches.
‣ Note that this function may also return cudaErrorInitializationError, cudaErrorInsufficientDriver or
cudaErrorNoDevice if this call tries to initialize internal CUDA RT state.
‣ Note that as specified by cudaStreamAddCallback no CUDA function may be called from callback.
cudaErrorNotPermitted may, but is not guaranteed to, be returned as a diagnostic in such case.
See also:
cudaMalloc3D, cudaMalloc, cudaMallocPitch, cudaFree, cudaFreeArray, cudaMallocHost ( C API),
cudaFreeHost, cudaHostAlloc, make_cudaExtent, cuArray3DCreate
__host__cudaError_t cudaMallocArray (cudaArray_t
*array, const cudaChannelFormatDesc *desc, size_t width,
size_t height, unsigned int flags)
Allocate an array on the device.
Parameters
array
- Pointer to allocated array in device memory
|
117
|
||
|
|
|