|
|
Modules
CUresult cuGraphMemcpyNodeGetParams (CUgraphNode
hNode, CUDA_MEMCPY3D *nodeParams)
Returns a memcpy node's parameters.
Parameters
hNode
- Node to get the parameters for
nodeParams
- Pointer to return the parameters
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_VALUE
Description
Returns the parameters of memcpy node hNode in nodeParams.
Note:
‣ Graph objects are not threadsafe. More here.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuMemcpy3D, cuGraphAddMemcpyNode, cuGraphMemcpyNodeSetParams
CUresult cuGraphMemcpyNodeSetParams (CUgraphNode
hNode, const CUDA_MEMCPY3D *nodeParams)
Sets a memcpy node's parameters.
Parameters
hNode
- Node to set the parameters for
nodeParams
- Parameters to copy
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_VALUE,
|
487
Modules
Description
Sets the parameters of memcpy node hNode to nodeParams.
Note:
‣ Graph objects are not threadsafe. More here.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuGraphNodeSetParams, cuMemcpy3D, cuGraphAddMemcpyNode, cuGraphMemcpyNodeGetParams
CUresult cuGraphMemFreeNodeGetParams
(CUgraphNode hNode, CUdeviceptr *dptr_out)
Returns a memory free node's parameters.
Parameters
hNode
- Node to get the parameters for
dptr_out
- Pointer to return the device address
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_VALUE
Description
Returns the address of a memory free node hNode in dptr_out.
Note:
‣ Graph objects are not threadsafe. More here.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuGraphAddMemFreeNode, cuGraphMemAllocNodeGetParams
|
488
Modules
CUresult cuGraphMemsetNodeGetParams (CUgraphNode
hNode, CUDA_MEMSET_NODE_PARAMS
*nodeParams)
Returns a memset node's parameters.
Parameters
hNode
- Node to get the parameters for
nodeParams
- Pointer to return the parameters
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_VALUE
Description
Returns the parameters of memset node hNode in nodeParams.
Note:
‣ Graph objects are not threadsafe. More here.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuMemsetD2D32, cuGraphAddMemsetNode, cuGraphMemsetNodeSetParams
CUresult cuGraphMemsetNodeSetParams (CUgraphNode
hNode, const CUDA_MEMSET_NODE_PARAMS
*nodeParams)
Sets a memset node's parameters.
Parameters
hNode
- Node to set the parameters for
nodeParams
- Parameters to copy
|
489
Modules
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_VALUE
Description
Sets the parameters of memset node hNode to nodeParams.
Note:
‣ Graph objects are not threadsafe. More here.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuGraphNodeSetParams, cuMemsetD2D32, cuGraphAddMemsetNode,
cuGraphMemsetNodeGetParams
CUresult cuGraphNodeFindInClone (CUgraphNode
*phNode, CUgraphNode hOriginalNode, CUgraph
hClonedGraph)
Finds a cloned version of a node.
Parameters
phNode
- Returns handle to the cloned node
hOriginalNode
- Handle to the original node
hClonedGraph
- Cloned graph to query
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE,
Description
This function returns the node in hClonedGraph corresponding to hOriginalNode in the original
graph.
hClonedGraph must have been cloned from hOriginalGraph via cuGraphClone.
hOriginalNode must have been in hOriginalGraph at the time of the call to cuGraphClone,
|
490
Modules
and the corresponding cloned node in hClonedGraph must not have been removed. The cloned node
is then returned via phClonedNode.
Note:
‣ Graph objects are not threadsafe. More here.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuGraphClone
CUresult cuGraphNodeGetDependencies (CUgraphNode
hNode, CUgraphNode *dependencies, size_t
*numDependencies)
Returns a node's dependencies.
Parameters
hNode
- Node to query
dependencies
- Pointer to return the dependencies
numDependencies
- See description
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_VALUE
Description
Returns a list of node's dependencies. dependencies may be NULL, in which case this function
will return the number of dependencies in numDependencies. Otherwise, numDependencies
entries will be filled in. If numDependencies is higher than the actual number of dependencies, the
remaining entries in dependencies will be set to NULL, and the number of nodes actually obtained
will be returned in numDependencies.
Note:
‣ Graph objects are not threadsafe. More here.
|
491
Modules
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuGraphNodeGetDependentNodes, cuGraphGetNodes, cuGraphGetRootNodes, cuGraphGetEdges,
cuGraphAddDependencies, cuGraphRemoveDependencies
CUresult cuGraphNodeGetDependencies_v2
(CUgraphNode hNode, CUgraphNode *dependencies,
CUgraphEdgeData *edgeData, size_t *numDependencies)
Returns a node's dependencies (12.3+).
Parameters
hNode
- Node to query
dependencies
- Pointer to return the dependencies
edgeData
- Optional array to return edge data for each dependency
numDependencies
- See description
Returns
CUDA_SUCCESS, CUDA_ERROR_LOSSY_QUERY, CUDA_ERROR_DEINITIALIZED,
CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_INVALID_VALUE
Description
Returns a list of node's dependencies. dependencies may be NULL, in which case this function
will return the number of dependencies in numDependencies. Otherwise, numDependencies
entries will be filled in. If numDependencies is higher than the actual number of dependencies, the
remaining entries in dependencies will be set to NULL, and the number of nodes actually obtained
will be returned in numDependencies.
Note that if an edge has non-zero (non-default) edge data and edgeData is NULL, this API will
return CUDA_ERROR_LOSSY_QUERY. If edgeData is non-NULL, then dependencies must
be as well.
Note:
‣ Graph objects are not threadsafe. More here.
‣ Note that this function may also return error codes from previous, asynchronous launches.
|
492
Modules
See also:
cuGraphNodeGetDependentNodes, cuGraphGetNodes, cuGraphGetRootNodes, cuGraphGetEdges,
cuGraphAddDependencies, cuGraphRemoveDependencies
CUresult cuGraphNodeGetDependentNodes
(CUgraphNode hNode, CUgraphNode *dependentNodes,
size_t *numDependentNodes)
Returns a node's dependent nodes.
Parameters
hNode
- Node to query
dependentNodes
- Pointer to return the dependent nodes
numDependentNodes
- See description
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_VALUE
Description
Returns a list of node's dependent nodes. dependentNodes may be NULL, in which case
this function will return the number of dependent nodes in numDependentNodes. Otherwise,
numDependentNodes entries will be filled in. If numDependentNodes is higher than the actual
number of dependent nodes, the remaining entries in dependentNodes will be set to NULL, and the
number of nodes actually obtained will be returned in numDependentNodes.
Note:
‣ Graph objects are not threadsafe. More here.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuGraphNodeGetDependencies, cuGraphGetNodes, cuGraphGetRootNodes, cuGraphGetEdges,
cuGraphAddDependencies, cuGraphRemoveDependencies
|
493
Modules
CUresult cuGraphNodeGetDependentNodes_v2
(CUgraphNode hNode, CUgraphNode
*dependentNodes, CUgraphEdgeData *edgeData, size_t
*numDependentNodes)
Returns a node's dependent nodes (12.3+).
Parameters
hNode
- Node to query
dependentNodes
- Pointer to return the dependent nodes
edgeData
- Optional pointer to return edge data for dependent nodes
numDependentNodes
- See description
Returns
CUDA_SUCCESS, CUDA_ERROR_LOSSY_QUERY, CUDA_ERROR_DEINITIALIZED,
CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_INVALID_VALUE
Description
Returns a list of node's dependent nodes. dependentNodes may be NULL, in which case
this function will return the number of dependent nodes in numDependentNodes. Otherwise,
numDependentNodes entries will be filled in. If numDependentNodes is higher than the actual
number of dependent nodes, the remaining entries in dependentNodes will be set to NULL, and the
number of nodes actually obtained will be returned in numDependentNodes.
Note that if an edge has non-zero (non-default) edge data and edgeData is NULL, this API will
return CUDA_ERROR_LOSSY_QUERY. If edgeData is non-NULL, then dependentNodes
must be as well.
Note:
‣ Graph objects are not threadsafe. More here.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuGraphNodeGetDependencies, cuGraphGetNodes, cuGraphGetRootNodes, cuGraphGetEdges,
cuGraphAddDependencies, cuGraphRemoveDependencies
|
494
Modules
CUresult cuGraphNodeGetEnabled (CUgraphExec
hGraphExec, CUgraphNode hNode, unsigned int
*isEnabled)
Query whether a node in the given graphExec is enabled.
Parameters
hGraphExec
- The executable graph in which to set the specified node
hNode
- Node from the graph from which graphExec was instantiated
isEnabled
- Location to return the enabled status of the node
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE,
Description
Sets isEnabled to 1 if hNode is enabled, or 0 if hNode is disabled.
The node is identified by the corresponding node hNode in the non-executable graph, from which the
executable graph was instantiated.
hNode must not have been removed from the original graph.
Note:
‣ Currently only kernel, memset and memcpy nodes are supported.
‣ This function will not reflect device-side updates for device-updatable kernel nodes.
Note:
‣ Graph objects are not threadsafe. More here.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuGraphNodeSetEnabled, cuGraphExecUpdate, cuGraphInstantiate cuGraphLaunch
|
495
Modules
CUresult cuGraphNodeGetType (CUgraphNode hNode,
CUgraphNodeType *type)
Returns a node's type.
Parameters
hNode
- Node to query
type
- Pointer to return the node type
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_VALUE
Description
Returns the node type of hNode in type.
Note:
‣ Graph objects are not threadsafe. More here.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuGraphGetNodes, cuGraphGetRootNodes, cuGraphChildGraphNodeGetGraph,
cuGraphKernelNodeGetParams, cuGraphKernelNodeSetParams, cuGraphHostNodeGetParams,
cuGraphHostNodeSetParams, cuGraphMemcpyNodeGetParams, cuGraphMemcpyNodeSetParams,
cuGraphMemsetNodeGetParams, cuGraphMemsetNodeSetParams
CUresult cuGraphNodeSetEnabled (CUgraphExec
hGraphExec, CUgraphNode hNode, unsigned int
isEnabled)
Enables or disables the specified node in the given graphExec.
Parameters
hGraphExec
- The executable graph in which to set the specified node
|
496
Modules
hNode
- Node from the graph from which graphExec was instantiated
isEnabled
- Node is enabled if != 0, otherwise the node is disabled
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE,
Description
Sets hNode to be either enabled or disabled. Disabled nodes are functionally equivalent to empty
nodes until they are reenabled. Existing node parameters are not affected by disabling/enabling the
node.
The node is identified by the corresponding node hNode in the non-executable graph, from which the
executable graph was instantiated.
hNode must not have been removed from the original graph.
The modifications only affect future launches of hGraphExec. Already enqueued or running
launches of hGraphExec are not affected by this call. hNode is also not modified by this call.
If hNode is a device-updatable kernel node, the next upload/launch of hGraphExec will overwrite
any previous device-side updates. Additionally, applying host updates to a device-updatable kernel
node while it is being updated from the device will result in undefined behavior.
Note:
Currently only kernel, memset and memcpy nodes are supported.
Note:
‣ Graph objects are not threadsafe. More here.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuGraphNodeGetEnabled, cuGraphExecUpdate, cuGraphInstantiate cuGraphLaunch
|
497
Modules
CUresult cuGraphNodeSetParams (CUgraphNode hNode,
CUgraphNodeParams *nodeParams)
Update's a graph node's parameters.
Parameters
hNode
- Node to set the parameters for
nodeParams
- Parameters to copy
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_NOT_SUPPORTED
Description
Sets the parameters of graph node hNode to nodeParams. The node type specified by
nodeParams->type must match the type of hNode. nodeParams must be fully initialized and all
unused bytes (reserved, padding) zeroed.
Modifying parameters is not supported for node types CU_GRAPH_NODE_TYPE_MEM_ALLOC
and CU_GRAPH_NODE_TYPE_MEM_FREE.
Note:
‣ Graph objects are not threadsafe. More here.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuGraphAddNode, cuGraphExecNodeSetParams
CUresult cuGraphReleaseUserObject (CUgraph graph,
CUuserObject object, unsigned int count)
Release a user object reference from a graph.
Parameters
graph
- The graph that will release the reference
object
- The user object to release a reference for
|
498
Modules
count
- The number of references to release, typically 1. Must be nonzero and not larger than INT_MAX.
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE
Description
Releases user object references owned by a graph.
See CUDA User Objects in the CUDA C++ Programming Guide for more information on user objects.
See also:
cuUserObjectCreate, cuUserObjectRetain, cuUserObjectRelease, cuGraphRetainUserObject,
cuGraphCreate
CUresult cuGraphRemoveDependencies (CUgraph hGraph,
const CUgraphNode *from, const CUgraphNode *to, size_t
numDependencies)
Removes dependency edges from a graph.
Parameters
hGraph
- Graph from which to remove dependencies
from
- Array of nodes that provide the dependencies
to
- Array of dependent nodes
numDependencies
- Number of dependencies to be removed
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE
Description
The number of dependencies to be removed is defined by numDependencies. Elements in
from and to at corresponding indices define a dependency. Each node in from and to must belong
to hGraph.
If numDependencies is 0, elements in from and to will be ignored. Specifying a non-existing
dependency will return an error.
|
499
Modules
Dependencies cannot be removed from graphs which contain allocation or free nodes. Any attempt to
do so will return an error.
Note:
‣ Graph objects are not threadsafe. More here.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuGraphAddDependencies, cuGraphGetEdges, cuGraphNodeGetDependencies,
cuGraphNodeGetDependentNodes
CUresult cuGraphRemoveDependencies_v2 (CUgraph
hGraph, const CUgraphNode *from, const CUgraphNode
*to, const CUgraphEdgeData *edgeData, size_t
numDependencies)
Removes dependency edges from a graph (12.3+).
Parameters
hGraph
- Graph from which to remove dependencies
from
- Array of nodes that provide the dependencies
to
- Array of dependent nodes
edgeData
- Optional array of edge data. If NULL, edge data is assumed to be default (zeroed).
numDependencies
- Number of dependencies to be removed
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE
Description
The number of dependencies to be removed is defined by numDependencies. Elements in
from and to at corresponding indices define a dependency. Each node in from and to must belong
to hGraph.
|
500
Modules
If numDependencies is 0, elements in from and to will be ignored. Specifying an edge that does
not exist in the graph, with data matching edgeData, results in an error. edgeData is nullable,
which is equivalent to passing default (zeroed) data for each edge.
Dependencies cannot be removed from graphs which contain allocation or free nodes. Any attempt to
do so will return an error.
Note:
‣ Graph objects are not threadsafe. More here.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuGraphAddDependencies, cuGraphGetEdges, cuGraphNodeGetDependencies,
cuGraphNodeGetDependentNodes
CUresult cuGraphRetainUserObject (CUgraph graph,
CUuserObject object, unsigned int count, unsigned int
flags)
Retain a reference to a user object from a graph.
Parameters
graph
- The graph to associate the reference with
object
- The user object to retain a reference for
count
- The number of references to add to the graph, typically 1. Must be nonzero and not larger than
INT_MAX.
flags
- The optional flag CU_GRAPH_USER_OBJECT_MOVE transfers references from the calling
thread, rather than create new references. Pass 0 to create new references.
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE
Description
Creates or moves user object references that will be owned by a CUDA graph.
See CUDA User Objects in the CUDA C++ Programming Guide for more information on user objects.
|
501
Modules
See also:
cuUserObjectCreate, cuUserObjectRetain, cuUserObjectRelease, cuGraphReleaseUserObject,
cuGraphCreate
CUresult cuGraphUpload (CUgraphExec hGraphExec,
CUstream hStream)
Uploads an executable graph in a stream.
Parameters
hGraphExec
- Executable graph to upload
hStream
- Stream in which to upload the graph
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_VALUE
Description
Uploads hGraphExec to the device in hStream without executing it. Uploads of the same
hGraphExec will be serialized. Each upload is ordered behind both any previous work in hStream
and any previous launches of hGraphExec. Uses memory cached by stream to back the allocations
owned by hGraphExec.
Note:
‣ Graph objects are not threadsafe. More here.
‣ Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuGraphInstantiate, cuGraphLaunch, cuGraphExecDestroy
|
502
Modules
CUresult cuUserObjectCreate (CUuserObject *object_out,
void *ptr, CUhostFn destroy, unsigned int initialRefcount,
unsigned int flags)
Create a user object.
Parameters
object_out
- Location to return the user object handle
ptr
- The pointer to pass to the destroy function
destroy
- Callback to free the user object when it is no longer in use
initialRefcount
- The initial refcount to create the object with, typically 1. The initial references are owned by the
calling thread.
flags
- Currently it is required to pass CU_USER_OBJECT_NO_DESTRUCTOR_SYNC, which is the
only defined flag. This indicates that the destroy callback cannot be waited on by any CUDA API.
Users requiring synchronization of the callback should signal its completion manually.
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE
Description
Create a user object with the specified destructor callback and initial reference count. The initial
references are owned by the caller.
Destructor callbacks cannot make CUDA API calls and should avoid blocking behavior, as they are
executed by a shared internal thread. Another thread may be signaled to perform such actions, if it does
not block forward progress of tasks scheduled through CUDA.
See CUDA User Objects in the CUDA C++ Programming Guide for more information on user objects.
See also:
cuUserObjectRetain, cuUserObjectRelease, cuGraphRetainUserObject, cuGraphReleaseUserObject,
cuGraphCreate
|
503
Modules
CUresult cuUserObjectRelease (CUuserObject object,
unsigned int count)
Release a reference to a user object.
Parameters
object
- The object to release
count
- The number of references to release, typically 1. Must be nonzero and not larger than INT_MAX.
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE
Description
Releases user object references owned by the caller. The object's destructor is invoked if the reference
count reaches zero.
It is undefined behavior to release references not owned by the caller, or to use a user object handle
after all references are released.
See CUDA User Objects in the CUDA C++ Programming Guide for more information on user objects.
See also:
cuUserObjectCreate, cuUserObjectRetain, cuGraphRetainUserObject, cuGraphReleaseUserObject,
cuGraphCreate
CUresult cuUserObjectRetain (CUuserObject object,
unsigned int count)
Retain a reference to a user object.
Parameters
object
- The object to retain
count
- The number of references to retain, typically 1. Must be nonzero and not larger than INT_MAX.
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE
|
504
Modules
Description
Retains new references to a user object. The new references are owned by the caller.
See CUDA User Objects in the CUDA C++ Programming Guide for more information on user objects.
See also:
cuUserObjectCreate, cuUserObjectRelease, cuGraphRetainUserObject, cuGraphReleaseUserObject,
cuGraphCreate
6.25. Occupancy
This section describes the occupancy calculation functions of the low-level CUDA driver application
programming interface.
CUresult cuOccupancyAvailableDynamicSMemPerBlock
(size_t *dynamicSmemSize, CUfunction 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
numBlocks
- Number of blocks to fit on SM
blockSize
- Size of the blocks
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_UNKNOWN
Description
Returns in *dynamicSmemSize the maximum size of dynamic shared memory to allow
numBlocks blocks per SM.
|
505
Modules
Note that the API can also be used with context-less kernel CUkernel by querying the handle using
cuLibraryGetKernel() and then passing it to the API by casting to CUfunction. Here, the context to use
for calculations will be the current context.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
CUresult cuOccupancyMaxActiveBlocksPerMultiprocessor
(int *numBlocks, CUfunction func, int blockSize, size_t
dynamicSMemSize)
Returns occupancy of a function.
Parameters
numBlocks
- Returned occupancy
func
- Kernel 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
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_UNKNOWN
Description
Returns in *numBlocks the number of the maximum active blocks per streaming multiprocessor.
Note that the API can also be used with context-less kernel CUkernel by querying the handle using
cuLibraryGetKernel() and then passing it to the API by casting to CUfunction. Here, the context to use
for calculations will be the current context.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
|
506
Modules
cudaOccupancyMaxActiveBlocksPerMultiprocessor
CUresult
cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags
(int *numBlocks, CUfunction func, int blockSize, size_t
dynamicSMemSize, unsigned int flags)
Returns occupancy of a function.
Parameters
numBlocks
- Returned occupancy
func
- Kernel 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
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_UNKNOWN
Description
Returns in *numBlocks the number of the maximum active blocks per streaming multiprocessor.
The Flags parameter controls how special cases are handled. The valid flags are:
‣ CU_OCCUPANCY_DEFAULT, which maintains the default behavior as
cuOccupancyMaxActiveBlocksPerMultiprocessor;
‣ CU_OCCUPANCY_DISABLE_CACHING_OVERRIDE, which suppresses the default
behavior on platform where global caching affects occupancy. On such platforms, if
caching is enabled, but per-block SM resource usage would result in zero occupancy,
the occupancy calculator will calculate the occupancy as if caching is disabled. Setting
CU_OCCUPANCY_DISABLE_CACHING_OVERRIDE 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.
|
507
Modules
Note that the API can also be with launch context-less kernel CUkernel by querying the handle using
cuLibraryGetKernel() and then passing it to the API by casting to CUfunction. Here, the context to use
for calculations will be the current context.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags
CUresult cuOccupancyMaxActiveClusters (int
*numClusters, CUfunction func, const CUlaunchConfig
*config)
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
config
- Launch configuration for the given kernel function
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_INVALID_CLUSTER_SIZE, CUDA_ERROR_UNKNOWN
Description
If the function has required cluster size already set (see cudaFuncGetAttributes / cuFuncGetAttribute),
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 that the API can also be used with context-less kernel CUkernel by querying the handle using
cuLibraryGetKernel() and then passing it to the API by casting to CUfunction. Here, the context to
|
508
Modules
use for calculations will either be taken from the specified stream config->hStream or the current
context in case of NULL stream.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cudaFuncGetAttributes, cuFuncGetAttribute
CUresult cuOccupancyMaxPotentialBlockSize (int
*minGridSize, int *blockSize, CUfunction func,
CUoccupancyB2DSize blockSizeToDynamicSMemSize,
size_t dynamicSMemSize, int blockSizeLimit)
Suggest a launch configuration with reasonable occupancy.
Parameters
minGridSize
- Returned minimum grid size needed to achieve the maximum occupancy
blockSize
- Returned maximum block size that can achieve the maximum occupancy
func
- Kernel for which launch configuration is calculated
blockSizeToDynamicSMemSize
- A function that calculates how much per-block dynamic shared memory func uses based on the
block size
dynamicSMemSize
- Dynamic shared memory usage intended, in bytes
blockSizeLimit
- The maximum block size func is designed to handle
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_UNKNOWN
|
509
Modules
Description
Returns in *blockSize a reasonable block size that can achieve the maximum occupancy (or, the
maximum number of active warps with the fewest blocks per multiprocessor), and in *minGridSize
the minimum grid size to achieve the maximum occupancy.
If blockSizeLimit is 0, the configurator will use the maximum block size permitted by the device /
function instead.
If per-block dynamic shared memory allocation is not needed, the user should leave both
blockSizeToDynamicSMemSize and dynamicSMemSize as 0.
If per-block dynamic shared memory allocation is needed, then if the dynamic shared memory size
is constant regardless of block size, the size should be passed through dynamicSMemSize, and
blockSizeToDynamicSMemSize should be NULL.
Otherwise, if the per-block dynamic shared memory size varies with different block sizes, the user
needs to provide a unary function through blockSizeToDynamicSMemSize that computes the
dynamic shared memory needed by func for any given block size. dynamicSMemSize is ignored.
An example signature is:
// Take block size, returns dynamic shared memory needed
size_t blockToSmem(int blockSize);
Note that the API can also be used with context-less kernel CUkernel by querying the handle using
cuLibraryGetKernel() and then passing it to the API by casting to CUfunction. Here, the context to use
for calculations will be the current context.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cudaOccupancyMaxPotentialBlockSize
|
510
Modules
CUresult cuOccupancyMaxPotentialBlockSizeWithFlags
(int *minGridSize, int *blockSize, CUfunction func,
CUoccupancyB2DSize blockSizeToDynamicSMemSize,
size_t dynamicSMemSize, int blockSizeLimit, unsigned int
flags)
Suggest a launch configuration with reasonable occupancy.
Parameters
minGridSize
- Returned minimum grid size needed to achieve the maximum occupancy
blockSize
- Returned maximum block size that can achieve the maximum occupancy
func
- Kernel for which launch configuration is calculated
blockSizeToDynamicSMemSize
- A function that calculates how much per-block dynamic shared memory func uses based on the
block size
dynamicSMemSize
- Dynamic shared memory usage intended, in bytes
blockSizeLimit
- The maximum block size func is designed to handle
flags
- Options
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_UNKNOWN
Description
An extended version of cuOccupancyMaxPotentialBlockSize. In addition to arguments passed to
cuOccupancyMaxPotentialBlockSize, cuOccupancyMaxPotentialBlockSizeWithFlags also takes a
Flags parameter.
The Flags parameter controls how special cases are handled. The valid flags are:
‣ CU_OCCUPANCY_DEFAULT, which maintains the default behavior as
cuOccupancyMaxPotentialBlockSize;
‣ CU_OCCUPANCY_DISABLE_CACHING_OVERRIDE, which suppresses the default
behavior on platform where global caching affects occupancy. On such platforms, the launch
|
511
Modules
configurations that produces maximal occupancy might not support global caching. Setting
CU_OCCUPANCY_DISABLE_CACHING_OVERRIDE guarantees that the the produced launch
configuration is global caching compatible at a potential cost of occupancy. More information can
be found about this feature in the "Unified L1/Texture Cache" section of the Maxwell tuning guide.
Note that the API can also be used with context-less kernel CUkernel by querying the handle using
cuLibraryGetKernel() and then passing it to the API by casting to CUfunction. Here, the context to use
for calculations will be the current context.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cudaOccupancyMaxPotentialBlockSizeWithFlags
CUresult cuOccupancyMaxPotentialClusterSize (int
*clusterSize, CUfunction func, const CUlaunchConfig
*config)
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
config
- Launch configuration for the given kernel function
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE,
CUDA_ERROR_UNKNOWN
Description
The cluster dimensions in config are ignored. If func has a required cluster size set (see
cudaFuncGetAttributes / cuFuncGetAttribute),*clusterSize will reflect the required cluster size.
|
512
Modules
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 that the API can also be used with context-less kernel CUkernel by querying the handle using
cuLibraryGetKernel() and then passing it to the API by casting to CUfunction. Here, the context to
use for calculations will either be taken from the specified stream config->hStream or the current
context in case of NULL stream.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cudaFuncGetAttributes, cuFuncGetAttribute
6.26. Texture Reference Management
[DEPRECATED]
This section describes the deprecated texture reference management functions of the low-level CUDA
driver application programming interface.
CUresult cuTexRefCreate (CUtexref *pTexRef)
Creates a texture reference.
Parameters
pTexRef
- Returned texture reference
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Creates a texture reference and returns its handle in *pTexRef. Once created, the application must
call cuTexRefSetArray() or cuTexRefSetAddress() to associate the reference with allocated memory.
Other texture reference functions are used to specify the format and interpretation (addressing, filtering,
etc.) to be used when the memory is read through this texture reference.
|
513
Modules
See also:
cuTexRefDestroy
CUresult cuTexRefDestroy (CUtexref hTexRef)
Destroys a texture reference.
Parameters
hTexRef
- Texture reference to destroy
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Destroys the texture reference specified by hTexRef.
See also:
cuTexRefCreate
CUresult cuTexRefGetAddress (CUdeviceptr *pdptr,
CUtexref hTexRef)
Gets the address associated with a texture reference.
Parameters
pdptr
- Returned device address
hTexRef
- Texture reference
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
|
514
Modules
Returns in *pdptr the base address bound to the texture reference hTexRef, or returns
CUDA_ERROR_INVALID_VALUE if the texture reference is not bound to any device memory
range.
See also:
cuTexRefSetAddress, cuTexRefSetAddress2D, cuTexRefSetAddressMode, cuTexRefSetArray,
cuTexRefSetFilterMode, cuTexRefSetFlags, cuTexRefSetFormat, cuTexRefGetAddressMode,
cuTexRefGetArray, cuTexRefGetFilterMode, cuTexRefGetFlags, cuTexRefGetFormat
CUresult cuTexRefGetAddressMode (CUaddress_mode
*pam, CUtexref hTexRef, int dim)
Gets the addressing mode used by a texture reference.
Parameters
pam
- Returned addressing mode
hTexRef
- Texture reference
dim
- Dimension
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Returns in *pam the addressing mode corresponding to the dimension dim of the texture reference
hTexRef. Currently, the only valid value for dim are 0 and 1.
See also:
cuTexRefSetAddress, cuTexRefSetAddress2D, cuTexRefSetAddressMode, cuTexRefSetArray,
cuTexRefSetFilterMode, cuTexRefSetFlags, cuTexRefSetFormat, cuTexRefGetAddress,
cuTexRefGetArray, cuTexRefGetFilterMode, cuTexRefGetFlags, cuTexRefGetFormat
|
515
Modules
CUresult cuTexRefGetArray (CUarray *phArray, CUtexref
hTexRef)
Gets the array bound to a texture reference.
Parameters
phArray
- Returned array
hTexRef
- Texture reference
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Returns in *phArray the CUDA array bound to the texture reference hTexRef, or returns
CUDA_ERROR_INVALID_VALUE if the texture reference is not bound to any CUDA array.
See also:
cuTexRefSetAddress, cuTexRefSetAddress2D, cuTexRefSetAddressMode, cuTexRefSetArray,
cuTexRefSetFilterMode, cuTexRefSetFlags, cuTexRefSetFormat, cuTexRefGetAddress,
cuTexRefGetAddressMode, cuTexRefGetFilterMode, cuTexRefGetFlags, cuTexRefGetFormat
CUresult cuTexRefGetBorderColor (float *pBorderColor,
CUtexref hTexRef)
Gets the border color used by a texture reference.
Parameters
pBorderColor
- Returned Type and Value of RGBA color
hTexRef
- Texture reference
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
|
516
Modules
Description
Deprecated
Returns in pBorderColor, values of the RGBA color used by the texture reference hTexRef. The
color value is of type float and holds color components in the following sequence: pBorderColor[0]
holds 'R' component pBorderColor[1] holds 'G' component pBorderColor[2] holds 'B' component
pBorderColor[3] holds 'A' component
See also:
cuTexRefSetAddressMode, cuTexRefSetAddressMode, cuTexRefSetBorderColor
CUresult cuTexRefGetFilterMode (CUfilter_mode *pfm,
CUtexref hTexRef)
Gets the filter-mode used by a texture reference.
Parameters
pfm
- Returned filtering mode
hTexRef
- Texture reference
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Returns in *pfm the filtering mode of the texture reference hTexRef.
See also:
cuTexRefSetAddress, cuTexRefSetAddress2D, cuTexRefSetAddressMode, cuTexRefSetArray,
cuTexRefSetFilterMode, cuTexRefSetFlags, cuTexRefSetFormat, cuTexRefGetAddress,
cuTexRefGetAddressMode, cuTexRefGetArray, cuTexRefGetFlags, cuTexRefGetFormat
|
517
Modules
CUresult cuTexRefGetFlags (unsigned int *pFlags,
CUtexref hTexRef)
Gets the flags used by a texture reference.
Parameters
pFlags
- Returned flags
hTexRef
- Texture reference
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Returns in *pFlags the flags of the texture reference hTexRef.
See also:
cuTexRefSetAddress, cuTexRefSetAddress2D, cuTexRefSetAddressMode, cuTexRefSetArray,
cuTexRefSetFilterMode, cuTexRefSetFlags, cuTexRefSetFormat, cuTexRefGetAddress,
cuTexRefGetAddressMode, cuTexRefGetArray, cuTexRefGetFilterMode, cuTexRefGetFormat
CUresult cuTexRefGetFormat (CUarray_format *pFormat,
int *pNumChannels, CUtexref hTexRef)
Gets the format used by a texture reference.
Parameters
pFormat
- Returned format
pNumChannels
- Returned number of components
hTexRef
- Texture reference
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
|
518
Modules
Description
Deprecated
Returns in *pFormat and *pNumChannels the format and number of components of the CUDA
array bound to the texture reference hTexRef. If pFormat or pNumChannels is NULL, it will be
ignored.
See also:
cuTexRefSetAddress, cuTexRefSetAddress2D, cuTexRefSetAddressMode, cuTexRefSetArray,
cuTexRefSetFilterMode, cuTexRefSetFlags, cuTexRefSetFormat, cuTexRefGetAddress,
cuTexRefGetAddressMode, cuTexRefGetArray, cuTexRefGetFilterMode, cuTexRefGetFlags
CUresult cuTexRefGetMaxAnisotropy (int *pmaxAniso,
CUtexref hTexRef)
Gets the maximum anisotropy for a texture reference.
Parameters
pmaxAniso
- Returned maximum anisotropy
hTexRef
- Texture reference
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Returns the maximum anisotropy in pmaxAniso that's used when reading memory through the texture
reference hTexRef.
See also:
cuTexRefSetAddress, cuTexRefSetAddress2D, cuTexRefSetAddressMode, cuTexRefSetArray,
cuTexRefSetFlags, cuTexRefSetFormat, cuTexRefGetAddress, cuTexRefGetAddressMode,
cuTexRefGetArray, cuTexRefGetFilterMode, cuTexRefGetFlags, cuTexRefGetFormat
|
519
Modules
CUresult cuTexRefGetMipmapFilterMode (CUfilter_mode
*pfm, CUtexref hTexRef)
Gets the mipmap filtering mode for a texture reference.
Parameters
pfm
- Returned mipmap filtering mode
hTexRef
- Texture reference
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Returns the mipmap filtering mode in pfm that's used when reading memory through the texture
reference hTexRef.
See also:
cuTexRefSetAddress, cuTexRefSetAddress2D, cuTexRefSetAddressMode, cuTexRefSetArray,
cuTexRefSetFlags, cuTexRefSetFormat, cuTexRefGetAddress, cuTexRefGetAddressMode,
cuTexRefGetArray, cuTexRefGetFilterMode, cuTexRefGetFlags, cuTexRefGetFormat
CUresult cuTexRefGetMipmapLevelBias (float *pbias,
CUtexref hTexRef)
Gets the mipmap level bias for a texture reference.
Parameters
pbias
- Returned mipmap level bias
hTexRef
- Texture reference
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
|
520
Modules
Description
Deprecated
Returns the mipmap level bias in pBias that's added to the specified mipmap level when reading
memory through the texture reference hTexRef.
See also:
cuTexRefSetAddress, cuTexRefSetAddress2D, cuTexRefSetAddressMode, cuTexRefSetArray,
cuTexRefSetFlags, cuTexRefSetFormat, cuTexRefGetAddress, cuTexRefGetAddressMode,
cuTexRefGetArray, cuTexRefGetFilterMode, cuTexRefGetFlags, cuTexRefGetFormat
CUresult cuTexRefGetMipmapLevelClamp
(float *pminMipmapLevelClamp, float
*pmaxMipmapLevelClamp, CUtexref hTexRef)
Gets the min/max mipmap level clamps for a texture reference.
Parameters
pminMipmapLevelClamp
- Returned mipmap min level clamp
pmaxMipmapLevelClamp
- Returned mipmap max level clamp
hTexRef
- Texture reference
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Returns the min/max mipmap level clamps in pminMipmapLevelClamp and
pmaxMipmapLevelClamp that's used when reading memory through the texture reference
hTexRef.
See also:
cuTexRefSetAddress, cuTexRefSetAddress2D, cuTexRefSetAddressMode, cuTexRefSetArray,
cuTexRefSetFlags, cuTexRefSetFormat, cuTexRefGetAddress, cuTexRefGetAddressMode,
cuTexRefGetArray, cuTexRefGetFilterMode, cuTexRefGetFlags, cuTexRefGetFormat
|
521
Modules
CUresult cuTexRefGetMipmappedArray
(CUmipmappedArray *phMipmappedArray, CUtexref
hTexRef)
Gets the mipmapped array bound to a texture reference.
Parameters
phMipmappedArray
- Returned mipmapped array
hTexRef
- Texture reference
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Returns in *phMipmappedArray the CUDA mipmapped array bound to the texture reference
hTexRef, or returns CUDA_ERROR_INVALID_VALUE if the texture reference is not bound to any
CUDA mipmapped array.
See also:
cuTexRefSetAddress, cuTexRefSetAddress2D, cuTexRefSetAddressMode, cuTexRefSetArray,
cuTexRefSetFilterMode, cuTexRefSetFlags, cuTexRefSetFormat, cuTexRefGetAddress,
cuTexRefGetAddressMode, cuTexRefGetFilterMode, cuTexRefGetFlags, cuTexRefGetFormat
CUresult cuTexRefSetAddress (size_t *ByteOffset,
CUtexref hTexRef, CUdeviceptr dptr, size_t bytes)
Binds an address as a texture reference.
Parameters
ByteOffset
- Returned byte offset
hTexRef
- Texture reference to bind
dptr
- Device pointer to bind
|
522
Modules
bytes
- Size of memory to bind in bytes
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Binds a linear address range to the texture reference hTexRef. Any previous address or CUDA array
state associated with the texture reference is superseded by this function. Any memory previously
bound to hTexRef is unbound.
Since the hardware enforces an alignment requirement on texture base addresses,
cuTexRefSetAddress() passes back a byte offset in *ByteOffset that must be applied to texture
fetches in order to read from the desired memory. This offset must be divided by the texel size and
passed to kernels that read from the texture so they can be applied to the tex1Dfetch() function.
If the device memory pointer was returned from cuMemAlloc(), the offset is guaranteed to be 0 and
NULL may be passed as the ByteOffset parameter.
The total number of elements (or texels) in the linear address range cannot exceed
CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_LINEAR_WIDTH. The number of
elements is computed as (bytes / bytesPerElement), where bytesPerElement is determined from the
data format and number of components set using cuTexRefSetFormat().
See also:
cuTexRefSetAddress2D, cuTexRefSetAddressMode, cuTexRefSetArray, cuTexRefSetFilterMode,
cuTexRefSetFlags, cuTexRefSetFormat, cuTexRefGetAddress, cuTexRefGetAddressMode,
cuTexRefGetArray, cuTexRefGetFilterMode, cuTexRefGetFlags, cuTexRefGetFormat
CUresult cuTexRefSetAddress2D (CUtexref hTexRef,
const CUDA_ARRAY_DESCRIPTOR *desc, CUdeviceptr
dptr, size_t Pitch)
Binds an address as a 2D texture reference.
Parameters
hTexRef
- Texture reference to bind
desc
- Descriptor of CUDA array
|
523
Modules
dptr
- Device pointer to bind
Pitch
- Line pitch in bytes
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Binds a linear address range to the texture reference hTexRef. Any previous address or CUDA array
state associated with the texture reference is superseded by this function. Any memory previously
bound to hTexRef is unbound.
Using a tex2D() function inside a kernel requires a call to either cuTexRefSetArray() to bind the
corresponding texture reference to an array, or cuTexRefSetAddress2D() to bind the texture reference
to linear memory.
Function calls to cuTexRefSetFormat() cannot follow calls to cuTexRefSetAddress2D() for the same
texture reference.
It is required that dptr be aligned to the appropriate hardware-specific texture alignment. You can
query this value using the device attribute CU_DEVICE_ATTRIBUTE_TEXTURE_ALIGNMENT. If
an unaligned dptr is supplied, CUDA_ERROR_INVALID_VALUE is returned.
Pitch has to be aligned to the hardware-specific texture pitch alignment. This value can be queried
using the device attribute CU_DEVICE_ATTRIBUTE_TEXTURE_PITCH_ALIGNMENT. If an
unaligned Pitch is supplied, CUDA_ERROR_INVALID_VALUE is returned.
Width and Height, which are specified in elements (or texels), cannot exceed
CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LINEAR_WIDTH and
CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LINEAR_HEIGHT
respectively. Pitch, which is specified in bytes, cannot exceed
CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LINEAR_PITCH.
See also:
cuTexRefSetAddress, cuTexRefSetAddressMode, cuTexRefSetArray, cuTexRefSetFilterMode,
cuTexRefSetFlags, cuTexRefSetFormat, cuTexRefGetAddress, cuTexRefGetAddressMode,
cuTexRefGetArray, cuTexRefGetFilterMode, cuTexRefGetFlags, cuTexRefGetFormat
|
524
Modules
CUresult cuTexRefSetAddressMode (CUtexref hTexRef,
int dim, CUaddress_mode am)
Sets the addressing mode for a texture reference.
Parameters
hTexRef
- Texture reference
dim
- Dimension
am
- Addressing mode to set
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Specifies the addressing mode am for the given dimension dim of the texture reference hTexRef. If
dim is zero, the addressing mode is applied to the first parameter of the functions used to fetch from
the texture; if dim is 1, the second, and so on. CUaddress_mode is defined as:
typedef enum CUaddress_mode_enum {
CU_TR_ADDRESS_MODE_WRAP = 0,
CU_TR_ADDRESS_MODE_CLAMP = 1,
CU_TR_ADDRESS_MODE_MIRROR = 2,
CU_TR_ADDRESS_MODE_BORDER = 3
} CUaddress_mode;
Note that this call has no effect if hTexRef is bound to linear memory. Also, if the flag,
CU_TRSF_NORMALIZED_COORDINATES, is not set, the only supported address mode is
CU_TR_ADDRESS_MODE_CLAMP.
See also:
cuTexRefSetAddress, cuTexRefSetAddress2D, cuTexRefSetArray, cuTexRefSetFilterMode,
cuTexRefSetFlags, cuTexRefSetFormat, cuTexRefGetAddress, cuTexRefGetAddressMode,
cuTexRefGetArray, cuTexRefGetFilterMode, cuTexRefGetFlags, cuTexRefGetFormat
|
525
Modules
CUresult cuTexRefSetArray (CUtexref hTexRef, CUarray
hArray, unsigned int Flags)
Binds an array as a texture reference.
Parameters
hTexRef
- Texture reference to bind
hArray
- Array to bind
Flags
- Options (must be CU_TRSA_OVERRIDE_FORMAT)
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Binds the CUDA array hArray to the texture reference hTexRef. Any previous address or CUDA
array state associated with the texture reference is superseded by this function. Flags must be set to
CU_TRSA_OVERRIDE_FORMAT. Any CUDA array previously bound to hTexRef is unbound.
See also:
cuTexRefSetAddress, cuTexRefSetAddress2D, cuTexRefSetAddressMode, cuTexRefSetFilterMode,
cuTexRefSetFlags, cuTexRefSetFormat, cuTexRefGetAddress, cuTexRefGetAddressMode,
cuTexRefGetArray, cuTexRefGetFilterMode, cuTexRefGetFlags, cuTexRefGetFormat
CUresult cuTexRefSetBorderColor (CUtexref hTexRef,
float *pBorderColor)
Sets the border color for a texture reference.
Parameters
hTexRef
- Texture reference
pBorderColor
- RGBA color
|
526
Modules
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Specifies the value of the RGBA color via the pBorderColor to the texture reference hTexRef.
The color value supports only float type and holds color components in the following sequence:
pBorderColor[0] holds 'R' component pBorderColor[1] holds 'G' component pBorderColor[2] holds 'B'
component pBorderColor[3] holds 'A' component
Note that the color values can be set only when the Address mode is set to
CU_TR_ADDRESS_MODE_BORDER using cuTexRefSetAddressMode. Applications using integer
border color values have to "reinterpret_cast" their values to float.
See also:
cuTexRefSetAddressMode, cuTexRefGetAddressMode, cuTexRefGetBorderColor
CUresult cuTexRefSetFilterMode (CUtexref hTexRef,
CUfilter_mode fm)
Sets the filtering mode for a texture reference.
Parameters
hTexRef
- Texture reference
fm
- Filtering mode to set
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Specifies the filtering mode fm to be used when reading memory through the texture reference
hTexRef. CUfilter_mode_enum is defined as:
typedef enum CUfilter_mode_enum {
CU_TR_FILTER_MODE_POINT = 0,
CU_TR_FILTER_MODE_LINEAR = 1
} CUfilter_mode;
|
527
Modules
Note that this call has no effect if hTexRef is bound to linear memory.
See also:
cuTexRefSetAddress, cuTexRefSetAddress2D, cuTexRefSetAddressMode, cuTexRefSetArray,
cuTexRefSetFlags, cuTexRefSetFormat, cuTexRefGetAddress, cuTexRefGetAddressMode,
cuTexRefGetArray, cuTexRefGetFilterMode, cuTexRefGetFlags, cuTexRefGetFormat
CUresult cuTexRefSetFlags (CUtexref hTexRef, unsigned
int Flags)
Sets the flags for a texture reference.
Parameters
hTexRef
- Texture reference
Flags
- Optional flags to set
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Specifies optional flags via Flags to specify the behavior of data returned through the texture
reference hTexRef. The valid flags are:
‣ CU_TRSF_READ_AS_INTEGER, which suppresses the default behavior of having the texture
promote integer data to floating point data in the range [0, 1]. Note that texture with 32-bit integer
format would not be promoted, regardless of whether or not this flag is specified;
‣ CU_TRSF_NORMALIZED_COORDINATES, which suppresses the default behavior of having
the texture coordinates range from [0, Dim) where Dim is the width or height of the CUDA array.
Instead, the texture coordinates [0, 1.0) reference the entire breadth of the array dimension;
‣ CU_TRSF_DISABLE_TRILINEAR_OPTIMIZATION, which disables any trilinear filtering
optimizations. Trilinear optimizations improve texture filtering performance by allowing bilinear
filtering on textures in scenarios where it can closely approximate the expected results.
See also:
cuTexRefSetAddress, cuTexRefSetAddress2D, cuTexRefSetAddressMode, cuTexRefSetArray,
cuTexRefSetFilterMode, cuTexRefSetFormat, cuTexRefGetAddress, cuTexRefGetAddressMode,
cuTexRefGetArray, cuTexRefGetFilterMode, cuTexRefGetFlags, cuTexRefGetFormat
|
528
Modules
CUresult cuTexRefSetFormat (CUtexref hTexRef,
CUarray_format fmt, int NumPackedComponents)
Sets the format for a texture reference.
Parameters
hTexRef
- Texture reference
fmt
- Format to set
NumPackedComponents
- Number of components per array element
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Specifies the format of the data to be read by the texture reference hTexRef. fmt and
NumPackedComponents are exactly analogous to the Format and NumChannels members of
the CUDA_ARRAY_DESCRIPTOR structure: They specify the format of each component and the
number of components per array element.
See also:
cuTexRefSetAddress, cuTexRefSetAddress2D, cuTexRefSetAddressMode, cuTexRefSetArray,
cuTexRefSetFilterMode, cuTexRefSetFlags, cuTexRefGetAddress, cuTexRefGetAddressMode,
cuTexRefGetArray, cuTexRefGetFilterMode, cuTexRefGetFlags, cuTexRefGetFormat,
cudaCreateChannelDesc
CUresult cuTexRefSetMaxAnisotropy (CUtexref hTexRef,
unsigned int maxAniso)
Sets the maximum anisotropy for a texture reference.
Parameters
hTexRef
- Texture reference
maxAniso
- Maximum anisotropy
|
529
Modules
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Specifies the maximum anisotropy maxAniso to be used when reading memory through the texture
reference hTexRef.
Note that this call has no effect if hTexRef is bound to linear memory.
See also:
cuTexRefSetAddress, cuTexRefSetAddress2D, cuTexRefSetAddressMode, cuTexRefSetArray,
cuTexRefSetFlags, cuTexRefSetFormat, cuTexRefGetAddress, cuTexRefGetAddressMode,
cuTexRefGetArray, cuTexRefGetFilterMode, cuTexRefGetFlags, cuTexRefGetFormat
CUresult cuTexRefSetMipmapFilterMode (CUtexref
hTexRef, CUfilter_mode fm)
Sets the mipmap filtering mode for a texture reference.
Parameters
hTexRef
- Texture reference
fm
- Filtering mode to set
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Specifies the mipmap filtering mode fm to be used when reading memory through the texture reference
hTexRef. CUfilter_mode_enum is defined as:
typedef enum CUfilter_mode_enum {
CU_TR_FILTER_MODE_POINT = 0,
CU_TR_FILTER_MODE_LINEAR = 1
} CUfilter_mode;
Note that this call has no effect if hTexRef is not bound to a mipmapped array.
|
530
Modules
See also:
cuTexRefSetAddress, cuTexRefSetAddress2D, cuTexRefSetAddressMode, cuTexRefSetArray,
cuTexRefSetFlags, cuTexRefSetFormat, cuTexRefGetAddress, cuTexRefGetAddressMode,
cuTexRefGetArray, cuTexRefGetFilterMode, cuTexRefGetFlags, cuTexRefGetFormat
CUresult cuTexRefSetMipmapLevelBias (CUtexref
hTexRef, float bias)
Sets the mipmap level bias for a texture reference.
Parameters
hTexRef
- Texture reference
bias
- Mipmap level bias
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Specifies the mipmap level bias bias to be added to the specified mipmap level when reading memory
through the texture reference hTexRef.
Note that this call has no effect if hTexRef is not bound to a mipmapped array.
See also:
cuTexRefSetAddress, cuTexRefSetAddress2D, cuTexRefSetAddressMode, cuTexRefSetArray,
cuTexRefSetFlags, cuTexRefSetFormat, cuTexRefGetAddress, cuTexRefGetAddressMode,
cuTexRefGetArray, cuTexRefGetFilterMode, cuTexRefGetFlags, cuTexRefGetFormat
CUresult cuTexRefSetMipmapLevelClamp (CUtexref
hTexRef, float minMipmapLevelClamp, float
maxMipmapLevelClamp)
Sets the mipmap min/max mipmap level clamps for a texture reference.
Parameters
hTexRef
- Texture reference
|
531
Modules
minMipmapLevelClamp
- Mipmap min level clamp
maxMipmapLevelClamp
- Mipmap max level clamp
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Specifies the min/max mipmap level clamps, minMipmapLevelClamp and
maxMipmapLevelClamp respectively, to be used when reading memory through the texture
reference hTexRef.
Note that this call has no effect if hTexRef is not bound to a mipmapped array.
See also:
cuTexRefSetAddress, cuTexRefSetAddress2D, cuTexRefSetAddressMode, cuTexRefSetArray,
cuTexRefSetFlags, cuTexRefSetFormat, cuTexRefGetAddress, cuTexRefGetAddressMode,
cuTexRefGetArray, cuTexRefGetFilterMode, cuTexRefGetFlags, cuTexRefGetFormat
CUresult cuTexRefSetMipmappedArray (CUtexref
hTexRef, CUmipmappedArray hMipmappedArray,
unsigned int Flags)
Binds a mipmapped array to a texture reference.
Parameters
hTexRef
- Texture reference to bind
hMipmappedArray
- Mipmapped array to bind
Flags
- Options (must be CU_TRSA_OVERRIDE_FORMAT)
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
|
532
Modules
Description
Deprecated
Binds the CUDA mipmapped array hMipmappedArray to the texture reference hTexRef. Any
previous address or CUDA array state associated with the texture reference is superseded by this
function. Flags must be set to CU_TRSA_OVERRIDE_FORMAT. Any CUDA array previously
bound to hTexRef is unbound.
See also:
cuTexRefSetAddress, cuTexRefSetAddress2D, cuTexRefSetAddressMode, cuTexRefSetFilterMode,
cuTexRefSetFlags, cuTexRefSetFormat, cuTexRefGetAddress, cuTexRefGetAddressMode,
cuTexRefGetArray, cuTexRefGetFilterMode, cuTexRefGetFlags, cuTexRefGetFormat
6.27. Surface Reference Management
[DEPRECATED]
This section describes the surface reference management functions of the low-level CUDA driver
application programming interface.
CUresult cuSurfRefGetArray (CUarray *phArray,
CUsurfref hSurfRef)
Passes back the CUDA array bound to a surface reference.
Parameters
phArray
- Surface reference handle
hSurfRef
- Surface reference handle
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Returns in *phArray the CUDA array bound to the surface reference hSurfRef, or returns
CUDA_ERROR_INVALID_VALUE if the surface reference is not bound to any CUDA array.
|
533
Modules
See also:
cuModuleGetSurfRef, cuSurfRefSetArray
CUresult cuSurfRefSetArray (CUsurfref hSurfRef,
CUarray hArray, unsigned int Flags)
Sets the CUDA array for a surface reference.
Parameters
hSurfRef
- Surface reference handle
hArray
- CUDA array handle
Flags
- set to 0
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Deprecated
Sets the CUDA array hArray to be read and written by the surface reference hSurfRef. Any
previous CUDA array state associated with the surface reference is superseded by this function. Flags
must be set to 0. The CUDA_ARRAY3D_SURFACE_LDST flag must have been set for the CUDA
array. Any CUDA array previously bound to hSurfRef is unbound.
See also:
cuModuleGetSurfRef, cuSurfRefGetArray
6.28. Texture Object Management
This section describes the texture object management functions of the low-level CUDA driver
application programming interface. The texture object API is only supported on devices of compute
capability 3.0 or higher.
|
534
Modules
CUresult cuTexObjectCreate (CUtexObject *pTexObject,
const CUDA_RESOURCE_DESC *pResDesc,
const CUDA_TEXTURE_DESC *pTexDesc, const
CUDA_RESOURCE_VIEW_DESC *pResViewDesc)
Creates a texture object.
Parameters
pTexObject
- Texture object to create
pResDesc
- Resource descriptor
pTexDesc
- Texture descriptor
pResViewDesc
- Resource view descriptor
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Creates a texture object and returns it in pTexObject. pResDesc describes the data to texture from.
pTexDesc describes how the data should be sampled. pResViewDesc is an optional argument that
specifies an alternate format for the data described by pResDesc, and also describes the subresource
region to restrict access to when texturing. pResViewDesc can only be specified if the type of
resource is a CUDA array or a CUDA mipmapped array not in a block compressed format.
Texture objects are only supported on devices of compute capability 3.0 or higher. Additionally, a
texture object is an opaque value, and, as such, should only be accessed through CUDA API calls.
The CUDA_RESOURCE_DESC structure is defined as:
typedef struct CUDA_RESOURCE_DESC_st
{
CUresourcetype resType;
union {
struct {
CUarray hArray;
} array;
struct {
CUmipmappedArray hMipmappedArray;
} mipmap;
struct {
CUdeviceptr devPtr;
CUarray_format format;
unsigned int numChannels;
|
535
Modules
size_t sizeInBytes;
} linear;
struct {
CUdeviceptr devPtr;
CUarray_format format;
unsigned int numChannels;
size_t width;
size_t height;
size_t pitchInBytes;
} pitch2D;
} res;
unsigned int flags;
} CUDA_RESOURCE_DESC;
where:
‣ CUDA_RESOURCE_DESC::resType specifies the type of resource to texture from.
CUresourceType is defined as:
typedef enum CUresourcetype_enum {
CU_RESOURCE_TYPE_ARRAY
= 0x00,
CU_RESOURCE_TYPE_MIPMAPPED_ARRAY = 0x01,
CU_RESOURCE_TYPE_LINEAR
= 0x02,
CU_RESOURCE_TYPE_PITCH2D
= 0x03
} CUresourcetype;
If CUDA_RESOURCE_DESC::resType is set to CU_RESOURCE_TYPE_ARRAY,
CUDA_RESOURCE_DESC::res::array::hArray must be set to a valid CUDA array handle.
If CUDA_RESOURCE_DESC::resType is set to CU_RESOURCE_TYPE_MIPMAPPED_ARRAY,
CUDA_RESOURCE_DESC::res::mipmap::hMipmappedArray must be set to a valid CUDA
mipmapped array handle.
If CUDA_RESOURCE_DESC::resType is set to CU_RESOURCE_TYPE_LINEAR,
CUDA_RESOURCE_DESC::res::linear::devPtr must be set to a valid device
pointer, that is aligned to CU_DEVICE_ATTRIBUTE_TEXTURE_ALIGNMENT.
CUDA_RESOURCE_DESC::res::linear::format and
CUDA_RESOURCE_DESC::res::linear::numChannels describe the format of each component and
the number of components per array element. CUDA_RESOURCE_DESC::res::linear::sizeInBytes
specifies the size of the array in bytes. The total number of elements in the linear address range cannot
exceed CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_LINEAR_WIDTH. The number of
elements is computed as (sizeInBytes / (sizeof(format) * numChannels)).
If CUDA_RESOURCE_DESC::resType is set to CU_RESOURCE_TYPE_PITCH2D,
CUDA_RESOURCE_DESC::res::pitch2D::devPtr must be set to a valid device
pointer, that is aligned to CU_DEVICE_ATTRIBUTE_TEXTURE_ALIGNMENT.
CUDA_RESOURCE_DESC::res::pitch2D::format and
CUDA_RESOURCE_DESC::res::pitch2D::numChannels describe the format of each component
and the number of components per array element. CUDA_RESOURCE_DESC::res::pitch2D::width
and CUDA_RESOURCE_DESC::res::pitch2D::height specify
the width and height of the array in elements, and cannot exceed
CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LINEAR_WIDTH and
CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LINEAR_HEIGHT respectively.
|
536
Modules
CUDA_RESOURCE_DESC::res::pitch2D::pitchInBytes specifies the pitch between two rows in
bytes and has to be aligned to CU_DEVICE_ATTRIBUTE_TEXTURE_PITCH_ALIGNMENT. Pitch
cannot exceed CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LINEAR_PITCH.
‣ flags must be set to zero.
The CUDA_TEXTURE_DESC struct is defined as
typedef struct CUDA_TEXTURE_DESC_st {
CUaddress_mode addressMode[3];
CUfilter_mode filterMode;
unsigned int flags;
unsigned int maxAnisotropy;
CUfilter_mode mipmapFilterMode;
float mipmapLevelBias;
float minMipmapLevelClamp;
float maxMipmapLevelClamp;
} CUDA_TEXTURE_DESC;
where
‣ CUDA_TEXTURE_DESC::addressMode specifies the addressing mode for each dimension of the
texture data. CUaddress_mode is defined as:
typedef enum CUaddress_mode_enum {
CU_TR_ADDRESS_MODE_WRAP = 0,
CU_TR_ADDRESS_MODE_CLAMP = 1,
CU_TR_ADDRESS_MODE_MIRROR = 2,
CU_TR_ADDRESS_MODE_BORDER = 3
} CUaddress_mode;
This is ignored if CUDA_RESOURCE_DESC::resType is CU_RESOURCE_TYPE_LINEAR.
Also, if the flag, CU_TRSF_NORMALIZED_COORDINATES is not set, the only supported
address mode is CU_TR_ADDRESS_MODE_CLAMP.
‣ CUDA_TEXTURE_DESC::filterMode specifies the filtering mode to be used when fetching from
the texture. CUfilter_mode is defined as:
typedef enum CUfilter_mode_enum {
CU_TR_FILTER_MODE_POINT = 0,
CU_TR_FILTER_MODE_LINEAR = 1
} CUfilter_mode;
This is ignored if CUDA_RESOURCE_DESC::resType is CU_RESOURCE_TYPE_LINEAR.
‣ CUDA_TEXTURE_DESC::flags can be any combination of the following:
‣ CU_TRSF_READ_AS_INTEGER, which suppresses the default behavior of having the
texture promote integer data to floating point data in the range [0, 1]. Note that texture with 32-
bit integer format would not be promoted, regardless of whether or not this flag is specified.
‣ CU_TRSF_NORMALIZED_COORDINATES, which suppresses the default behavior of
having the texture coordinates range from [0, Dim) where Dim is the width or height of the
CUDA array. Instead, the texture coordinates [0, 1.0) reference the entire breadth of the array
dimension; Note that for CUDA mipmapped arrays, this flag has to be set.
‣ CU_TRSF_DISABLE_TRILINEAR_OPTIMIZATION, which disables any trilinear filtering
optimizations. Trilinear optimizations improve texture filtering performance by allowing
bilinear filtering on textures in scenarios where it can closely approximate the expected results.
‣ CU_TRSF_SEAMLESS_CUBEMAP, which enables seamless cube map filtering.
This flag can only be specified if the underlying resource is a CUDA array or a CUDA
|
537
Modules
mipmapped array that was created with the flag CUDA_ARRAY3D_CUBEMAP.
When seamless cube map filtering is enabled, texture address modes specified
by CUDA_TEXTURE_DESC::addressMode are ignored. Instead, if the
CUDA_TEXTURE_DESC::filterMode is set to CU_TR_FILTER_MODE_POINT the
address mode CU_TR_ADDRESS_MODE_CLAMP will be applied for all dimensions. If the
CUDA_TEXTURE_DESC::filterMode is set to CU_TR_FILTER_MODE_LINEAR seamless
cube map filtering will be performed when sampling along the cube face borders.
‣ CUDA_TEXTURE_DESC::maxAnisotropy specifies the maximum anisotropy ratio to be used
when doing anisotropic filtering. This value will be clamped to the range [1,16].
‣ CUDA_TEXTURE_DESC::mipmapFilterMode specifies the filter mode when the calculated
mipmap level lies between two defined mipmap levels.
‣ CUDA_TEXTURE_DESC::mipmapLevelBias specifies the offset to be applied to the calculated
mipmap level.
‣ CUDA_TEXTURE_DESC::minMipmapLevelClamp specifies the lower end of the mipmap level
range to clamp access to.
‣ CUDA_TEXTURE_DESC::maxMipmapLevelClamp specifies the upper end of the mipmap level
range to clamp access to.
The CUDA_RESOURCE_VIEW_DESC struct is defined as
typedef struct CUDA_RESOURCE_VIEW_DESC_st
{
CUresourceViewFormat format;
size_t width;
size_t height;
size_t depth;
unsigned int firstMipmapLevel;
unsigned int lastMipmapLevel;
unsigned int firstLayer;
unsigned int lastLayer;
} CUDA_RESOURCE_VIEW_DESC;
where:
‣ CUDA_RESOURCE_VIEW_DESC::format specifies how the data contained in the CUDA
array or CUDA mipmapped array should be interpreted. Note that this can incur a change
in size of the texture data. If the resource view format is a block compressed format, then
the underlying CUDA array or CUDA mipmapped array has to have a base of format
CU_AD_FORMAT_UNSIGNED_INT32. with 2 or 4 channels, depending on the block
compressed format. For ex., BC1 and BC4 require the underlying CUDA array to have a format
of CU_AD_FORMAT_UNSIGNED_INT32 with 2 channels. The other BC formats require the
underlying resource to have the same base format but with 4 channels.
‣ CUDA_RESOURCE_VIEW_DESC::width specifies the new width of the texture data. If the
resource view format is a block compressed format, this value has to be 4 times the original width
of the resource. For non block compressed formats, this value has to be equal to that of the original
resource.
|
538
Modules
‣ CUDA_RESOURCE_VIEW_DESC::height specifies the new height of the texture data. If the
resource view format is a block compressed format, this value has to be 4 times the original height
of the resource. For non block compressed formats, this value has to be equal to that of the original
resource.
‣ CUDA_RESOURCE_VIEW_DESC::depth specifies the new depth of the texture data. This value
has to be equal to that of the original resource.
‣ CUDA_RESOURCE_VIEW_DESC::firstMipmapLevel specifies the most detailed
mipmap level. This will be the new mipmap level zero. For non-mipmapped resources,
this value has to be zero.CUDA_TEXTURE_DESC::minMipmapLevelClamp and
CUDA_TEXTURE_DESC::maxMipmapLevelClamp will be relative to this value. For ex., if the
firstMipmapLevel is set to 2, and a minMipmapLevelClamp of 1.2 is specified, then the actual
minimum mipmap level clamp will be 3.2.
‣ CUDA_RESOURCE_VIEW_DESC::lastMipmapLevel specifies the least detailed mipmap level.
For non-mipmapped resources, this value has to be zero.
‣ CUDA_RESOURCE_VIEW_DESC::firstLayer specifies the first layer index for layered textures.
This will be the new layer zero. For non-layered resources, this value has to be zero.
‣ CUDA_RESOURCE_VIEW_DESC::lastLayer specifies the last layer index for layered textures.
For non-layered resources, this value has to be zero.
See also:
cuTexObjectDestroy, cudaCreateTextureObject
CUresult cuTexObjectDestroy (CUtexObject texObject)
Destroys a texture object.
Parameters
texObject
- Texture object to destroy
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Destroys the texture object specified by texObject.
See also:
cuTexObjectCreate, cudaDestroyTextureObject
|
539
Modules
CUresult cuTexObjectGetResourceDesc
(CUDA_RESOURCE_DESC *pResDesc, CUtexObject
texObject)
Returns a texture object's resource descriptor.
Parameters
pResDesc
- Resource descriptor
texObject
- Texture object
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Returns the resource descriptor for the texture object specified by texObject.
See also:
cuTexObjectCreate, cudaGetTextureObjectResourceDesc,
CUresult cuTexObjectGetResourceViewDesc
(CUDA_RESOURCE_VIEW_DESC *pResViewDesc,
CUtexObject texObject)
Returns a texture object's resource view descriptor.
Parameters
pResViewDesc
- Resource view descriptor
texObject
- Texture object
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
|
540
Modules
Description
Returns the resource view descriptor for the texture object specified by texObject. If no resource
view was set for texObject, the CUDA_ERROR_INVALID_VALUE is returned.
See also:
cuTexObjectCreate, cudaGetTextureObjectResourceViewDesc
CUresult cuTexObjectGetTextureDesc
(CUDA_TEXTURE_DESC *pTexDesc, CUtexObject
texObject)
Returns a texture object's texture descriptor.
Parameters
pTexDesc
- Texture descriptor
texObject
- Texture object
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Returns the texture descriptor for the texture object specified by texObject.
See also:
cuTexObjectCreate, cudaGetTextureObjectTextureDesc
6.29. Surface Object Management
This section describes the surface object management functions of the low-level CUDA driver
application programming interface. The surface object API is only supported on devices of compute
capability 3.0 or higher.
|
541
Modules
CUresult cuSurfObjectCreate (CUsurfObject *pSurfObject,
const CUDA_RESOURCE_DESC *pResDesc)
Creates a surface object.
Parameters
pSurfObject
- Surface object to create
pResDesc
- Resource descriptor
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Creates a surface object and returns it in pSurfObject. pResDesc describes the
data to perform surface load/stores on. CUDA_RESOURCE_DESC::resType must be
CU_RESOURCE_TYPE_ARRAY and CUDA_RESOURCE_DESC::res::array::hArray must be set to
a valid CUDA array handle. CUDA_RESOURCE_DESC::flags must be set to zero.
Surface objects are only supported on devices of compute capability 3.0 or higher. Additionally, a
surface object is an opaque value, and, as such, should only be accessed through CUDA API calls.
See also:
cuSurfObjectDestroy, cudaCreateSurfaceObject
CUresult cuSurfObjectDestroy (CUsurfObject surfObject)
Destroys a surface object.
Parameters
surfObject
- Surface object to destroy
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Destroys the surface object specified by surfObject.
|
542
Modules
See also:
cuSurfObjectCreate, cudaDestroySurfaceObject
CUresult cuSurfObjectGetResourceDesc
(CUDA_RESOURCE_DESC *pResDesc, CUsurfObject
surfObject)
Returns a surface object's resource descriptor.
Parameters
pResDesc
- Resource descriptor
surfObject
- Surface object
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Returns the resource descriptor for the surface object specified by surfObject.
See also:
cuSurfObjectCreate, cudaGetSurfaceObjectResourceDesc
6.30. Tensor Map Object Managment
This section describes the tensor map object management functions of the low-level CUDA driver
application programming interface. The tensor core API is only supported on devices of compute
capability 9.0 or higher.
|
543
Modules
CUresult cuTensorMapEncodeIm2col (CUtensorMap
*tensorMap, CUtensorMapDataType tensorDataType,
cuuint32_t tensorRank, void *globalAddress,
const cuuint64_t *globalDim, const cuuint64_t
*globalStrides, const int *pixelBoxLowerCorner, const
int *pixelBoxUpperCorner, cuuint32_t channelsPerPixel,
cuuint32_t pixelsPerColumn, const cuuint32_t
*elementStrides, CUtensorMapInterleave interleave,
CUtensorMapSwizzle swizzle, CUtensorMapL2promotion
l2Promotion, CUtensorMapFloatOOBfill oobFill)
Create a tensor map descriptor object representing im2col memory region.
Parameters
tensorMap
- Tensor map object to create
tensorDataType
- Tensor data type
tensorRank
- Dimensionality of tensor; must be at least 3
globalAddress
- Starting address of memory region described by tensor
globalDim
- Array containing tensor size (number of elements) along each of the tensorRank dimensions
globalStrides
- Array containing stride size (in bytes) along each of the tensorRank - 1 dimensions
pixelBoxLowerCorner
- Array containing DHW dimensions of lower box corner
pixelBoxUpperCorner
- Array containing DHW dimensions of upper box corner
channelsPerPixel
- Number of channels per pixel
pixelsPerColumn
- Number of pixels per column
elementStrides
- Array containing traversal stride in each of the tensorRank dimensions
interleave
- Type of interleaved layout the tensor addresses
|
544
Modules
swizzle
- Bank swizzling pattern inside shared memory
l2Promotion
- L2 promotion size
oobFill
- Indicate whether zero or special NaN constant will be used to fill out-of-bound elements
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Creates a descriptor for Tensor Memory Access (TMA) object specified by the parameters describing a
im2col memory layout and returns it in tensorMap.
Tensor map objects are only supported on devices of compute capability 9.0 or higher. Additionally, a
tensor map object is an opaque value, and, as such, should only be accessed through CUDA APIs and
PTX.
The parameters passed are bound to the following requirements:
‣ tensorMap address must be aligned to 64 bytes.
‣ tensorDataType has to be an enum from CUtensorMapDataType which is defined as:
typedef enum CUtensorMapDataType_enum {
CU_TENSOR_MAP_DATA_TYPE_UINT8 = 0,
// 1 byte
CU_TENSOR_MAP_DATA_TYPE_UINT16,
// 2 bytes
CU_TENSOR_MAP_DATA_TYPE_UINT32,
// 4 bytes
CU_TENSOR_MAP_DATA_TYPE_INT32,
// 4 bytes
CU_TENSOR_MAP_DATA_TYPE_UINT64,
// 8 bytes
CU_TENSOR_MAP_DATA_TYPE_INT64,
// 8 bytes
CU_TENSOR_MAP_DATA_TYPE_FLOAT16,
// 2 bytes
CU_TENSOR_MAP_DATA_TYPE_FLOAT32,
// 4 bytes
CU_TENSOR_MAP_DATA_TYPE_FLOAT64,
// 8 bytes
CU_TENSOR_MAP_DATA_TYPE_BFLOAT16,
// 2 bytes
CU_TENSOR_MAP_DATA_TYPE_FLOAT32_FTZ,
// 4 bytes
CU_TENSOR_MAP_DATA_TYPE_TFLOAT32,
// 4 bytes
CU_TENSOR_MAP_DATA_TYPE_TFLOAT32_FTZ
// 4 bytes
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B,
// 4 bits
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B,
// 4 bits
CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B
// 6 bits
} CUtensorMapDataType;
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B copies '16 x U4' packed
values to memory aligned as 8 bytes. There are no gaps between packed values.
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B copies '16 x U4' packed values to
memory aligned as 16 bytes. There are 8 byte gaps between every 8 byte chunk of packed values.
CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B copies '16 x U6' packed values to memory
aligned as 16 bytes. There are 4 byte gaps between every 12 byte chunk of packed values.
‣ tensorRank, which specifies the number of tensor dimensions, must be 3, 4, or 5.
|
545
Modules
‣ globalAddress, which specifies the starting address of the memory region described, must be
16 byte aligned. The following requirements need to also be met:
‣ When interleave is CU_TENSOR_MAP_INTERLEAVE_32B, globalAddress must
be 32 byte aligned.
‣ When tensorDataType is CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B or
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, globalAddress must be 32 byte
aligned.
‣ globalDim array, which specifies tensor size of each of the tensorRank dimensions, must be
non-zero and less than or equal to 2^32. Additionally, the following requirements need to be met
for the packed data types:
‣ When tensorDataType is CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B or
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, globalDim[0] must be a multiple of
128.
‣ When tensorDataType is CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B,
globalDim[0] must be a multiple of 2.
‣ Dimension for the packed data types must reflect the number of individual U# values.
‣ globalStrides array, which specifies tensor stride of each of the lower tensorRank - 1
dimensions in bytes, must be a multiple of 16 and less than 2^40. Additionally, the following
requirements need to be met:
‣ When interleave is CU_TENSOR_MAP_INTERLEAVE_32B, the strides must be a
multiple of 32.
‣ When tensorDataType is CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B or
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, the strides must be a multiple of 32.
Each following dimension specified includes previous dimension stride:
globalStrides[0] = globalDim[0] * elementSizeInBytes(tensorDataType) +
padding[0];
for (i = 1; i < tensorRank - 1; i++)
globalStrides[i] = globalStrides[i - 1] * (globalDim[i] +
padding[i]);
assert(globalStrides[i] >= globalDim[i]);
‣ pixelBoxLowerCorner array specifies the coordinate offsets {D, H, W} of the bounding
box from top/left/front corner. The number of offsets and their precision depend on the tensor
dimensionality:
‣ When tensorRank is 3, one signed offset within range [-32768, 32767] is supported.
‣ When tensorRank is 4, two signed offsets each within range [-128, 127] are supported.
‣ When tensorRank is 5, three offsets each within range [-16, 15] are supported.
‣ pixelBoxUpperCorner array specifies the coordinate offsets {D, H, W} of the bounding box
from bottom/right/back corner. The number of offsets and their precision depend on the tensor
dimensionality:
‣ When tensorRank is 3, one signed offset within range [-32768, 32767] is supported.
|
546
Modules
‣ When tensorRank is 4, two signed offsets each within range [-128, 127] are supported.
‣ When tensorRank is 5, three offsets each within range [-16, 15] are supported. The
bounding box specified by pixelBoxLowerCorner and pixelBoxUpperCorner must
have non-zero area.
‣ channelsPerPixel, which specifies the number of elements which must be
accessed along C dimension, must be less than or equal to 256. Additionally, when
tensorDataType is CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B or
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, channelsPerPixel must be 128.
‣ pixelsPerColumn, which specifies the number of elements that must be accessed along the {N,
D, H, W} dimensions, must be less than or equal to 1024.
‣ elementStrides array, which specifies the iteration step along each of the tensorRank
dimensions, must be non-zero and less than or equal to 8. Note that when interleave is
CU_TENSOR_MAP_INTERLEAVE_NONE, the first element of this array is ignored since TMA
doesn’t support the stride for dimension zero. When all elements of the elementStrides array
are one, boxDim specifies the number of elements to load. However, if elementStrides[i] is
not equal to one for some i, then TMA loads ceil( boxDim[i] / elementStrides[i]) number of
elements along i-th dimension. To load N elements along i-th dimension, boxDim[i] must be set to
N * elementStrides[i].
‣ interleave specifies the interleaved layout of type CUtensorMapInterleave, which is defined
as:
typedef enum CUtensorMapInterleave_enum {
CU_TENSOR_MAP_INTERLEAVE_NONE = 0,
CU_TENSOR_MAP_INTERLEAVE_16B,
CU_TENSOR_MAP_INTERLEAVE_32B
} CUtensorMapInterleave;
TMA supports interleaved layouts like NC/8HWC8 where C8 utilizes 16 bytes in
memory assuming 2 byte per channel or NC/16HWC16 where C16 uses 32 bytes. When
interleave is CU_TENSOR_MAP_INTERLEAVE_NONE and swizzle is not
CU_TENSOR_MAP_SWIZZLE_NONE, the bounding box inner dimension (computed as
channelsPerPixel multiplied by element size in bytes derived from tensorDataType)
must be less than or equal to the swizzle size.
‣ CU_TENSOR_MAP_SWIZZLE_32B requires the bounding box inner dimension to be <= 32.
‣ CU_TENSOR_MAP_SWIZZLE_64B requires the bounding box inner dimension to be <= 64.
‣ CU_TENSOR_MAP_SWIZZLE_128B* require the bounding box
inner dimension to be <= 128. Additionally, tensorDataType of
CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B requires interleave to be
CU_TENSOR_MAP_INTERLEAVE_NONE.
‣ swizzle, which specifies the shared memory bank swizzling pattern, has to be of type
CUtensorMapSwizzle which is defined as:
typedef enum CUtensorMapSwizzle_enum {
CU_TENSOR_MAP_SWIZZLE_NONE = 0,
CU_TENSOR_MAP_SWIZZLE_32B,
// Swizzle 16B chunks
within 32B span
|
547
Modules
CU_TENSOR_MAP_SWIZZLE_64B,
// Swizzle 16B chunks
within 64B span
CU_TENSOR_MAP_SWIZZLE_128B,
// Swizzle 16B chunks
within 128B span
CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B,
// Swizzle 32B chunks
within 128B span
CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B_FLIP_8B, // Swizzle 32B chunks
within 128B span, additionally swap lower 8B with upper 8B within each 16B for
every alternate row
CU_TENSOR_MAP_SWIZZLE_128B_ATOM_64B
// Swizzle 64B chunks
within 128B span
} CUtensorMapSwizzle;
Data are organized in a specific order in global memory; however, this may not match the
order in which the application accesses data in shared memory. This difference in data
organization may cause bank conflicts when shared memory is accessed. In order to avoid
this problem, data can be loaded to shared memory with shuffling across shared memory
banks. When interleave is CU_TENSOR_MAP_INTERLEAVE_32B, swizzle must be
CU_TENSOR_MAP_SWIZZLE_32B. Other interleave modes can have any swizzling pattern.
When the tensorDataType is CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B, only
the following swizzle modes are supported:
‣ CU_TENSOR_MAP_SWIZZLE_NONE (Load & Store)
‣ CU_TENSOR_MAP_SWIZZLE_128B (Load & Store)
‣ CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B (Load & Store)
‣ CU_TENSOR_MAP_SWIZZLE_128B_ATOM_64B (Store only) When the
tensorDataType is CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, only the
following swizzle modes are supported:
‣ CU_TENSOR_MAP_SWIZZLE_NONE (Load only)
‣ CU_TENSOR_MAP_SWIZZLE_128B (Load only)
‣ CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B (Load only)
‣ l2Promotion specifies L2 fetch size which indicates the byte granularity at which L2 requests
are filled from DRAM. It must be of type CUtensorMapL2promotion, which is defined as:
typedef enum CUtensorMapL2promotion_enum {
CU_TENSOR_MAP_L2_PROMOTION_NONE = 0,
CU_TENSOR_MAP_L2_PROMOTION_L2_64B,
CU_TENSOR_MAP_L2_PROMOTION_L2_128B,
CU_TENSOR_MAP_L2_PROMOTION_L2_256B
} CUtensorMapL2promotion;
‣ oobFill, which indicates whether zero or a special NaN constant should be used to fill out-of-
bound elements, must be of type CUtensorMapFloatOOBfill which is defined as:
typedef enum CUtensorMapFloatOOBfill_enum {
CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE = 0,
CU_TENSOR_MAP_FLOAT_OOB_FILL_NAN_REQUEST_ZERO_FMA
} CUtensorMapFloatOOBfill;
Note that CU_TENSOR_MAP_FLOAT_OOB_FILL_NAN_REQUEST_ZERO_FMA
can only be used when tensorDataType represents a floating-point data type, and
when tensorDataType is not CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B,
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, and
CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B.
|
548
Modules
See also:
cuTensorMapEncodeTiled, cuTensorMapEncodeIm2colWide, cuTensorMapReplaceAddress
CUresult cuTensorMapEncodeIm2colWide (CUtensorMap
*tensorMap, CUtensorMapDataType tensorDataType,
cuuint32_t tensorRank, void *globalAddress,
const cuuint64_t *globalDim, const cuuint64_t
*globalStrides, int pixelBoxLowerCornerWidth,
int pixelBoxUpperCornerWidth, cuuint32_t
channelsPerPixel, cuuint32_t pixelsPerColumn, const
cuuint32_t *elementStrides, CUtensorMapInterleave
interleave, CUtensorMapIm2ColWideMode mode,
CUtensorMapSwizzle swizzle, CUtensorMapL2promotion
l2Promotion, CUtensorMapFloatOOBfill oobFill)
Create a tensor map descriptor object representing im2col memory region, but where the elements are
exclusively loaded along the W dimension.
Parameters
tensorMap
- Tensor map object to create
tensorDataType
- Tensor data type
tensorRank
- Dimensionality of tensor; must be at least 3
globalAddress
- Starting address of memory region described by tensor
globalDim
- Array containing tensor size (number of elements) along each of the tensorRank dimensions
globalStrides
- Array containing stride size (in bytes) along each of the tensorRank - 1 dimensions
pixelBoxLowerCornerWidth
- Width offset of left box corner
pixelBoxUpperCornerWidth
- Width offset of right box corner
channelsPerPixel
- Number of channels per pixel
|
549
Modules
pixelsPerColumn
- Number of pixels per column
elementStrides
- Array containing traversal stride in each of the tensorRank dimensions
interleave
- Type of interleaved layout the tensor addresses
mode
- W or W128 mode
swizzle
- Bank swizzling pattern inside shared memory
l2Promotion
- L2 promotion size
oobFill
- Indicate whether zero or special NaN constant will be used to fill out-of-bound elements
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Creates a descriptor for Tensor Memory Access (TMA) object specified by the parameters describing
a im2col memory layout and where the row is always loaded along the W dimensuin and returns it in
tensorMap. This assumes the tensor layout in memory is either NDHWC, NHWC, or NWC.
This API is only supported on devices of compute capability 10.0 or higher. Additionally, a tensor map
object is an opaque value, and, as such, should only be accessed through CUDA APIs and PTX.
The parameters passed are bound to the following requirements:
‣ tensorMap address must be aligned to 64 bytes.
‣ tensorDataType has to be an enum from CUtensorMapDataType which is defined as:
typedef enum CUtensorMapDataType_enum {
CU_TENSOR_MAP_DATA_TYPE_UINT8 = 0,
// 1 byte
CU_TENSOR_MAP_DATA_TYPE_UINT16,
// 2 bytes
CU_TENSOR_MAP_DATA_TYPE_UINT32,
// 4 bytes
CU_TENSOR_MAP_DATA_TYPE_INT32,
// 4 bytes
CU_TENSOR_MAP_DATA_TYPE_UINT64,
// 8 bytes
CU_TENSOR_MAP_DATA_TYPE_INT64,
// 8 bytes
CU_TENSOR_MAP_DATA_TYPE_FLOAT16,
// 2 bytes
CU_TENSOR_MAP_DATA_TYPE_FLOAT32,
// 4 bytes
CU_TENSOR_MAP_DATA_TYPE_FLOAT64,
// 8 bytes
CU_TENSOR_MAP_DATA_TYPE_BFLOAT16,
// 2 bytes
CU_TENSOR_MAP_DATA_TYPE_FLOAT32_FTZ,
// 4 bytes
CU_TENSOR_MAP_DATA_TYPE_TFLOAT32,
// 4 bytes
CU_TENSOR_MAP_DATA_TYPE_TFLOAT32_FTZ
// 4 bytes
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B,
// 4 bits
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B,
// 4 bits
CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B
// 6 bits
} CUtensorMapDataType;
|
550
Modules
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B copies '16 x U4' packed
values to memory aligned as 8 bytes. There are no gaps between packed values.
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B copies '16 x U4' packed values to
memory aligned as 16 bytes. There are 8 byte gaps between every 8 byte chunk of packed values.
CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B copies '16 x U6' packed values to memory
aligned as 16 bytes. There are 4 byte gaps between every 12 byte chunk of packed values.
‣ tensorRank, which specifies the number of tensor dimensions, must be 3, 4, or 5.
‣ globalAddress, which specifies the starting address of the memory region described, must be
16 byte aligned. The following requirements need to also be met:
‣ When interleave is CU_TENSOR_MAP_INTERLEAVE_32B, globalAddress must
be 32 byte aligned.
‣ When tensorDataType is CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B or
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, globalAddress must be 32 byte
aligned.
‣ globalDim array, which specifies tensor size of each of the tensorRank dimensions, must be
non-zero and less than or equal to 2^32. Additionally, the following requirements need to be met
for the packed data types:
‣ When tensorDataType is CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B or
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, globalDim[0] must be a multiple of
128.
‣ When tensorDataType is CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B,
globalDim[0] must be a multiple of 2.
‣ Dimension for the packed data types must reflect the number of individual U# values.
‣ globalStrides array, which specifies tensor stride of each of the lower tensorRank - 1
dimensions in bytes, must be a multiple of 16 and less than 2^40. Additionally, the following
requirements need to be met:
‣ When interleave is CU_TENSOR_MAP_INTERLEAVE_32B, the strides must be a
multiple of 32.
‣ When tensorDataType is CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B or
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, the strides must be a multiple of 32.
Each following dimension specified includes previous dimension stride:
globalStrides[0] = globalDim[0] * elementSizeInBytes(tensorDataType) +
padding[0];
for (i = 1; i < tensorRank - 1; i++)
globalStrides[i] = globalStrides[i - 1] * (globalDim[i] +
padding[i]);
assert(globalStrides[i] >= globalDim[i]);
‣ pixelBoxLowerCornerWidth specifies the coordinate offset W of the bounding box from left
corner. The offset must be within range [-32768, 32767].
|
551
Modules
‣ pixelBoxUpperCornerWidth specifies the coordinate offset W of the bounding box from
right corner. The offset must be within range [-32768, 32767].
The bounding box specified by pixelBoxLowerCornerWidth and
pixelBoxUpperCornerWidth must have non-zero area. Note that the size of the box along D and
H dimensions is always equal to one.
‣ channelsPerPixel, which specifies the number of elements which must be
accessed along C dimension, must be less than or equal to 256. Additionally, when
tensorDataType is CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B or
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, channelsPerPixel must be 128.
‣ pixelsPerColumn, which specifies the number of elements that must be accessed along
the W dimension, must be less than or equal to 1024. This field is ignored when mode is
CU_TENSOR_MAP_IM2COL_WIDE_MODE_W128.
‣ elementStrides array, which specifies the iteration step along each of the tensorRank
dimensions, must be non-zero and less than or equal to 8. Note that when interleave is
CU_TENSOR_MAP_INTERLEAVE_NONE, the first element of this array is ignored since TMA
doesn’t support the stride for dimension zero. When all elements of the elementStrides array
are one, boxDim specifies the number of elements to load. However, if elementStrides[i] is
not equal to one for some i, then TMA loads ceil( boxDim[i] / elementStrides[i]) number of
elements along i-th dimension. To load N elements along i-th dimension, boxDim[i] must be set to
N * elementStrides[i].
‣ interleave specifies the interleaved layout of type CUtensorMapInterleave, which is defined
as:
typedef enum CUtensorMapInterleave_enum {
CU_TENSOR_MAP_INTERLEAVE_NONE = 0,
CU_TENSOR_MAP_INTERLEAVE_16B,
CU_TENSOR_MAP_INTERLEAVE_32B
} CUtensorMapInterleave;
TMA supports interleaved layouts like NC/8HWC8 where C8 utilizes 16 bytes in memory
assuming 2 byte per channel or NC/16HWC16 where C16 uses 32 bytes. When interleave is
CU_TENSOR_MAP_INTERLEAVE_NONE, the bounding box inner dimension (computed as
channelsPerPixel multiplied by element size in bytes derived from tensorDataType)
must be less than or equal to the swizzle size.
‣ CU_TENSOR_MAP_SWIZZLE_64B requires the bounding box inner dimension to be <= 64.
‣ CU_TENSOR_MAP_SWIZZLE_128B* require the bounding box
inner dimension to be <= 128. Additionally, tensorDataType of
CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B requires interleave to be
CU_TENSOR_MAP_INTERLEAVE_NONE.
‣ mode, which describes loading of elements loaded along the W dimension, has to be one of the
following CUtensorMapIm2ColWideMode types:
CU_TENSOR_MAP_IM2COL_WIDE_MODE_W,
CU_TENSOR_MAP_IM2COL_WIDE_MODE_W128
|
552
Modules
CU_TENSOR_MAP_IM2COL_WIDE_MODE_W allows the number of elements loaded along
the W dimension to be specified via the pixelsPerColumn field.
‣ swizzle, which specifies the shared memory bank swizzling pattern, must be one of the
following CUtensorMapSwizzle modes (other swizzle modes are not supported):
typedef enum CUtensorMapSwizzle_enum {
CU_TENSOR_MAP_SWIZZLE_64B,
// Swizzle 16B chunks
within 64B span
CU_TENSOR_MAP_SWIZZLE_128B,
// Swizzle 16B chunks
within 128B span
CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B,
// Swizzle 32B chunks
within 128B span
} CUtensorMapSwizzle;
Data are organized in a specific order in global memory; however, this may not match the order
in which the application accesses data in shared memory. This difference in data organization
may cause bank conflicts when shared memory is accessed. In order to avoid this problem,
data can be loaded to shared memory with shuffling across shared memory banks. When the
tensorDataType is CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B, only the
following swizzle modes are supported:
‣ CU_TENSOR_MAP_SWIZZLE_128B (Load & Store)
‣ CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B (Load & Store) When the
tensorDataType is CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, only the
following swizzle modes are supported:
‣ CU_TENSOR_MAP_SWIZZLE_128B (Load only)
‣ CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B (Load only)
‣ l2Promotion specifies L2 fetch size which indicates the byte granularity at which L2 requests
are filled from DRAM. It must be of type CUtensorMapL2promotion, which is defined as:
typedef enum CUtensorMapL2promotion_enum {
CU_TENSOR_MAP_L2_PROMOTION_NONE = 0,
CU_TENSOR_MAP_L2_PROMOTION_L2_64B,
CU_TENSOR_MAP_L2_PROMOTION_L2_128B,
CU_TENSOR_MAP_L2_PROMOTION_L2_256B
} CUtensorMapL2promotion;
‣ oobFill, which indicates whether zero or a special NaN constant should be used to fill out-of-
bound elements, must be of type CUtensorMapFloatOOBfill which is defined as:
typedef enum CUtensorMapFloatOOBfill_enum {
CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE = 0,
CU_TENSOR_MAP_FLOAT_OOB_FILL_NAN_REQUEST_ZERO_FMA
} CUtensorMapFloatOOBfill;
Note that CU_TENSOR_MAP_FLOAT_OOB_FILL_NAN_REQUEST_ZERO_FMA
can only be used when tensorDataType represents a floating-point data type, and
when tensorDataType is not CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B,
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, and
CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B.
See also:
cuTensorMapEncodeTiled, cuTensorMapEncodeIm2col, cuTensorMapReplaceAddress
|
553
Modules
CUresult cuTensorMapEncodeTiled (CUtensorMap
*tensorMap, CUtensorMapDataType tensorDataType,
cuuint32_t tensorRank, void *globalAddress, const
cuuint64_t *globalDim, const cuuint64_t *globalStrides,
const cuuint32_t *boxDim, const cuuint32_t
*elementStrides, CUtensorMapInterleave interleave,
CUtensorMapSwizzle swizzle, CUtensorMapL2promotion
l2Promotion, CUtensorMapFloatOOBfill oobFill)
Create a tensor map descriptor object representing tiled memory region.
Parameters
tensorMap
- Tensor map object to create
tensorDataType
- Tensor data type
tensorRank
- Dimensionality of tensor
globalAddress
- Starting address of memory region described by tensor
globalDim
- Array containing tensor size (number of elements) along each of the tensorRank dimensions
globalStrides
- Array containing stride size (in bytes) along each of the tensorRank - 1 dimensions
boxDim
- Array containing traversal box size (number of elments) along each of the tensorRank
dimensions. Specifies how many elements to be traversed along each tensor dimension.
elementStrides
- Array containing traversal stride in each of the tensorRank dimensions
interleave
- Type of interleaved layout the tensor addresses
swizzle
- Bank swizzling pattern inside shared memory
l2Promotion
- L2 promotion size
oobFill
- Indicate whether zero or special NaN constant must be used to fill out-of-bound elements
|
554
Modules
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
Description
Creates a descriptor for Tensor Memory Access (TMA) object specified by the parameters describing a
tiled region and returns it in tensorMap.
Tensor map objects are only supported on devices of compute capability 9.0 or higher. Additionally, a
tensor map object is an opaque value, and, as such, should only be accessed through CUDA APIs and
PTX.
The parameters passed are bound to the following requirements:
‣ tensorMap address must be aligned to 64 bytes.
‣ tensorDataType has to be an enum from CUtensorMapDataType which is defined as:
typedef enum CUtensorMapDataType_enum {
CU_TENSOR_MAP_DATA_TYPE_UINT8 = 0,
// 1 byte
CU_TENSOR_MAP_DATA_TYPE_UINT16,
// 2 bytes
CU_TENSOR_MAP_DATA_TYPE_UINT32,
// 4 bytes
CU_TENSOR_MAP_DATA_TYPE_INT32,
// 4 bytes
CU_TENSOR_MAP_DATA_TYPE_UINT64,
// 8 bytes
CU_TENSOR_MAP_DATA_TYPE_INT64,
// 8 bytes
CU_TENSOR_MAP_DATA_TYPE_FLOAT16,
// 2 bytes
CU_TENSOR_MAP_DATA_TYPE_FLOAT32,
// 4 bytes
CU_TENSOR_MAP_DATA_TYPE_FLOAT64,
// 8 bytes
CU_TENSOR_MAP_DATA_TYPE_BFLOAT16,
// 2 bytes
CU_TENSOR_MAP_DATA_TYPE_FLOAT32_FTZ,
// 4 bytes
CU_TENSOR_MAP_DATA_TYPE_TFLOAT32,
// 4 bytes
CU_TENSOR_MAP_DATA_TYPE_TFLOAT32_FTZ,
// 4 bytes
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B,
// 4 bits
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B,
// 4 bits
CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B
// 6 bits
} CUtensorMapDataType;
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B copies '16 x U4' packed
values to memory aligned as 8 bytes. There are no gaps between packed values.
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B copies '16 x U4' packed values to
memory aligned as 16 bytes. There are 8 byte gaps between every 8 byte chunk of packed values.
CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B copies '16 x U6' packed values to memory
aligned as 16 bytes. There are 4 byte gaps between every 12 byte chunk of packed values.
‣ tensorRank must be non-zero and less than or equal to the maximum supported dimensionality
of 5. If interleave is not CU_TENSOR_MAP_INTERLEAVE_NONE, then tensorRank
must additionally be greater than or equal to 3.
‣ globalAddress, which specifies the starting address of the memory region described, must be
16 byte aligned. The following requirements need to also be met:
‣ When interleave is CU_TENSOR_MAP_INTERLEAVE_32B, globalAddress must
be 32 byte aligned.
|
555
Modules
‣ When tensorDataType is CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B or
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, globalAddress must be 32 byte
aligned.
‣ globalDim array, which specifies tensor size of each of the tensorRank dimensions, must be
non-zero and less than or equal to 2^32. Additionally, the following requirements need to be met
for the packed data types:
‣ When tensorDataType is CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B or
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, globalDim[0] must be a multiple of
128.
‣ When tensorDataType is CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B,
globalDim[0] must be a multiple of 2.
‣ Dimension for the packed data types must reflect the number of individual U# values.
‣ globalStrides array, which specifies tensor stride of each of the lower tensorRank - 1
dimensions in bytes, must be a multiple of 16 and less than 2^40. Additionally, the following
requirements need to be met:
‣ When interleave is CU_TENSOR_MAP_INTERLEAVE_32B, the strides must be a
multiple of 32.
‣ When tensorDataType is CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B or
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, the strides must be a multiple of 32.
Each following dimension specified includes previous dimension stride:
globalStrides[0] = globalDim[0] * elementSizeInBytes(tensorDataType) +
padding[0];
for (i = 1; i < tensorRank - 1; i++)
globalStrides[i] = globalStrides[i - 1] * (globalDim[i] +
padding[i]);
assert(globalStrides[i] >= globalDim[i]);
‣ boxDim array, which specifies number of elements to be traversed along each of the
tensorRank dimensions, must be non-zero and less than or equal to 256. Additionally, the
following requirements need to be met:
‣ When interleave is CU_TENSOR_MAP_INTERLEAVE_NONE, { boxDim[0] *
elementSizeInBytes( tensorDataType ) } must be a multiple of 16 bytes.
‣ When tensorDataType is CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B or
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, boxDim[0] must be 128.
‣ elementStrides array, which specifies the iteration step along each of the tensorRank
dimensions, must be non-zero and less than or equal to 8. Note that when interleave is
CU_TENSOR_MAP_INTERLEAVE_NONE, the first element of this array is ignored since TMA
doesn’t support the stride for dimension zero. When all elements of elementStrides array is
one, boxDim specifies the number of elements to load. However, if the elementStrides[i] is
not equal to one, then TMA loads ceil( boxDim[i] / elementStrides[i]) number of elements
along i-th dimension. To load N elements along i-th dimension, boxDim[i] must be set to N *
elementStrides[i].
|
556
Modules
‣ interleave specifies the interleaved layout of type CUtensorMapInterleave, which is defined
as:
typedef enum CUtensorMapInterleave_enum {
CU_TENSOR_MAP_INTERLEAVE_NONE = 0,
CU_TENSOR_MAP_INTERLEAVE_16B,
CU_TENSOR_MAP_INTERLEAVE_32B
} CUtensorMapInterleave;
TMA supports interleaved layouts like NC/8HWC8 where C8 utilizes 16 bytes in
memory assuming 2 byte per channel or NC/16HWC16 where C16 uses 32 bytes. When
interleave is CU_TENSOR_MAP_INTERLEAVE_NONE and swizzle is not
CU_TENSOR_MAP_SWIZZLE_NONE, the bounding box inner dimension (computed as
boxDim[0] multiplied by element size derived from tensorDataType) must be less than or
equal to the swizzle size.
‣ CU_TENSOR_MAP_SWIZZLE_32B requires the bounding box inner dimension to be <= 32.
‣ CU_TENSOR_MAP_SWIZZLE_64B requires the bounding box inner dimension to be <= 64.
‣ CU_TENSOR_MAP_SWIZZLE_128B* require the bounding box
inner dimension to be <= 128. Additionally, tensorDataType of
CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B requires interleave to be
CU_TENSOR_MAP_INTERLEAVE_NONE.
‣ swizzle, which specifies the shared memory bank swizzling pattern, has to be of type
CUtensorMapSwizzle which is defined as:
typedef enum CUtensorMapSwizzle_enum {
CU_TENSOR_MAP_SWIZZLE_NONE = 0,
CU_TENSOR_MAP_SWIZZLE_32B,
// Swizzle 16B chunks
within 32B span
CU_TENSOR_MAP_SWIZZLE_64B,
// Swizzle 16B chunks
within 64B span
CU_TENSOR_MAP_SWIZZLE_128B,
// Swizzle 16B chunks
within 128B span
CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B,
// Swizzle 32B chunks
within 128B span
CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B_FLIP_8B, // Swizzle 32B chunks
within 128B span, additionally swap lower 8B with upper 8B within each 16B for
every alternate row
CU_TENSOR_MAP_SWIZZLE_128B_ATOM_64B
// Swizzle 64B chunks
within 128B span
} CUtensorMapSwizzle;
Data are organized in a specific order in global memory; however, this may not match the
order in which the application accesses data in shared memory. This difference in data
organization may cause bank conflicts when shared memory is accessed. In order to avoid
this problem, data can be loaded to shared memory with shuffling across shared memory
banks. When interleave is CU_TENSOR_MAP_INTERLEAVE_32B, swizzle must be
CU_TENSOR_MAP_SWIZZLE_32B. Other interleave modes can have any swizzling pattern.
When the tensorDataType is CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B, only
the following swizzle modes are supported:
‣ CU_TENSOR_MAP_SWIZZLE_NONE (Load & Store)
‣ CU_TENSOR_MAP_SWIZZLE_128B (Load & Store)
‣ CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B (Load & Store)
|
557
Modules
‣ CU_TENSOR_MAP_SWIZZLE_128B_ATOM_64B (Store only) When the
tensorDataType is CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, only the
following swizzle modes are supported:
‣ CU_TENSOR_MAP_SWIZZLE_NONE (Load only)
‣ CU_TENSOR_MAP_SWIZZLE_128B (Load only)
‣ CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B (Load only)
‣ l2Promotion specifies L2 fetch size which indicates the byte granurality at which L2 requests is
filled from DRAM. It must be of type CUtensorMapL2promotion, which is defined as:
typedef enum CUtensorMapL2promotion_enum {
CU_TENSOR_MAP_L2_PROMOTION_NONE = 0,
CU_TENSOR_MAP_L2_PROMOTION_L2_64B,
CU_TENSOR_MAP_L2_PROMOTION_L2_128B,
CU_TENSOR_MAP_L2_PROMOTION_L2_256B
} CUtensorMapL2promotion;
‣ oobFill, which indicates whether zero or a special NaN constant should be used to fill out-of-
bound elements, must be of type CUtensorMapFloatOOBfill which is defined as:
typedef enum CUtensorMapFloatOOBfill_enum {
CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE = 0,
CU_TENSOR_MAP_FLOAT_OOB_FILL_NAN_REQUEST_ZERO_FMA
} CUtensorMapFloatOOBfill;
Note that CU_TENSOR_MAP_FLOAT_OOB_FILL_NAN_REQUEST_ZERO_FMA
can only be used when tensorDataType represents a floating-point data type, and
when tensorDataType is not CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN8B,
CU_TENSOR_MAP_DATA_TYPE_16U4_ALIGN16B, and
CU_TENSOR_MAP_DATA_TYPE_16U6_ALIGN16B.
See also:
cuTensorMapEncodeIm2col, cuTensorMapEncodeIm2colWide, cuTensorMapReplaceAddress
CUresult cuTensorMapReplaceAddress (CUtensorMap
*tensorMap, void *globalAddress)
Modify an existing tensor map descriptor with an updated global address.
Parameters
tensorMap
- Tensor map object to modify
globalAddress
- Starting address of memory region described by tensor, must follow previous alignment
requirements
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE
|
558
Modules
Description
Modifies the descriptor for Tensor Memory Access (TMA) object passed in tensorMap with an
updated globalAddress.
Tensor map objects are only supported on devices of compute capability 9.0 or higher. Additionally, a
tensor map object is an opaque value, and, as such, should only be accessed through CUDA API calls.
See also:
cuTensorMapEncodeTiled, cuTensorMapEncodeIm2col, cuTensorMapEncodeIm2colWide
6.31. Peer Context Memory Access
This section describes the direct peer context memory access functions of the low-level CUDA driver
application programming interface.
CUresult cuCtxDisablePeerAccess (CUcontext
peerContext)
Disables direct access to memory allocations in a peer context and unregisters any registered
allocations.
Parameters
peerContext
- Peer context to disable direct access to
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_PEER_ACCESS_NOT_ENABLED, CUDA_ERROR_INVALID_CONTEXT,
Description
Returns CUDA_ERROR_PEER_ACCESS_NOT_ENABLED if direct peer access has not yet been
enabled from peerContext to the current context.
Returns CUDA_ERROR_INVALID_CONTEXT if there is no current context, or if peerContext is
not a valid context.
Note:
Note that this function may also return error codes from previous, asynchronous launches.
|
559
Modules
See also:
cuDeviceCanAccessPeer, cuCtxEnablePeerAccess, cudaDeviceDisablePeerAccess
CUresult cuCtxEnablePeerAccess (CUcontext peerContext,
unsigned int Flags)
Enables direct access to memory allocations in a peer context.
Parameters
peerContext
- Peer context to enable direct access to from the current context
Flags
- Reserved for future use and must be set to 0
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED, CUDA_ERROR_TOO_MANY_PEERS,
CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_PEER_ACCESS_UNSUPPORTED,
CUDA_ERROR_INVALID_VALUE
Description
If both the current context and peerContext are on devices which support unified addressing (as
may be queried using CU_DEVICE_ATTRIBUTE_UNIFIED_ADDRESSING) and same major
compute capability, then on success all allocations from peerContext will immediately be
accessible by the current context. See Unified Addressing for additional details.
Note that access granted by this call is unidirectional and that in order to access memory from the
current context in peerContext, a separate symmetric call to cuCtxEnablePeerAccess() is required.
Note that there are both device-wide and system-wide limitations per system configuration, as noted in
the CUDA Programming Guide under the section "Peer-to-Peer Memory Access".
Returns CUDA_ERROR_PEER_ACCESS_UNSUPPORTED if cuDeviceCanAccessPeer() indicates
that the CUdevice of the current context cannot directly access memory from the CUdevice of
peerContext.
Returns CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED if direct access of
peerContext from the current context has already been enabled.
Returns CUDA_ERROR_TOO_MANY_PEERS if direct peer access is not possible because hardware
resources required for peer access have been exhausted.
Returns CUDA_ERROR_INVALID_CONTEXT if there is no current context, peerContext is not
a valid context, or if the current context is peerContext.
Returns CUDA_ERROR_INVALID_VALUE if Flags is not 0.
|
560
Modules
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuDeviceCanAccessPeer, cuCtxDisablePeerAccess, cudaDeviceEnablePeerAccess
CUresult cuDeviceCanAccessPeer (int *canAccessPeer,
CUdevice dev, CUdevice peerDev)
Queries if a device may directly access a peer device's memory.
Parameters
canAccessPeer
- Returned access capability
dev
- Device from which allocations on peerDev are to be directly accessed.
peerDev
- Device on which the allocations to be directly accessed by dev reside.
Returns
CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
CUDA_ERROR_INVALID_DEVICE
Description
Returns in *canAccessPeer a value of 1 if contexts on dev are capable of directly accessing
memory from contexts on peerDev and 0 otherwise. If direct access of peerDev from dev is
possible, then access may be enabled on two specific contexts by calling cuCtxEnablePeerAccess().
Note:
Note that this function may also return error codes from previous, asynchronous launches.
See also:
cuCtxEnablePeerAccess, cuCtxDisablePeerAccess, cudaDeviceCanAccessPeer
|
561
|
||
|
|
|