6.48. subgrid management
A subgrid schedule provides a simple scheme for avoiding a full barrier between task switches on the GPU.
It does so by allowing users to split up a kernel grid of work into subgrids for finer grained dependency management.
By organizing workItems and dependency tracking with subgrids, subsequent kernels can start right away and incrementally wait for work done by prior kernels to be finished, rather than having to wait for the prior kernel to complete all of its work before using any of it.
While PDL can achieve similar schedules by explicitly splitting grids into smaller grids, such a schedule would need to pay the startup and exit cost for CTAs in each split. In a subgrid schedule a CTA can stride across multiple subgrids of data, amortizing CTA startup and completion costs.
The subgrid API builds a schedule hierarchically:
create a schedule;
add one or more worker grids that define the worker launch configuration;
add worksets to each worker grid to group subgrids that share a user-data payload;
add subgrids to each workset and optionally attach dependencies on earlier subgrids.
The current scheduling policy is deterministic and append-ordered. Worker grids are considered in creation order. Within a worker grid, worksets are processed in creation order. Within a workset, subgrids are processed in creation order.
Subgrid schedules must be included in a cuda graph and cannot be modified after graph instantiation. The schedule is run by launching the corresponding instantiated graph.
Functions
- CUresult cuSubgridCreate(CUsubgrid *subgrid_out, CUsubgridWorkset hWorkset, const CUsubgrid *hDependencies, cuuint32_t numDependencies, CUoffset3D *workItemStart, CUoffset3D *workItemEnd)
create a subgrid
- CUresult cuSubgridCreate_v2(CUsubgrid *subgrid_out, CUsubgridWorkset hWorkset, const CUsubgrid *hDependencies, const CUsubgridDependencyType *dependencyTypes, cuuint32_t numDependencies, CUoffset3D *workItemStart, CUoffset3D *workItemEnd)
create a subgrid with dependency type support
- CUresult cuSubgridWorkerGridCreate(CUsubgridWorkerGrid *workerGrid_out, CUsubgridSchedule hSchedule, const CUsubgridWorkerGridDesc *workerGridDesc)
create a subgrid worker grid
- CUresult cuSubgridWorksetCreate(CUsubgridWorkset *workset_out, CUsubgridWorkerGrid hWorkerGrid, const void *userData, size_t userDataSize)
create a subgrid workset
6.48.1. Functions
-
CUresult cuSubgridCreate(CUsubgrid *subgrid_out, CUsubgridWorkset hWorkset, const CUsubgrid *hDependencies, cuuint32_t numDependencies, CUoffset3D *workItemStart, CUoffset3D *workItemEnd)
create a subgrid
A subgrid is the fundametal dependency management object in the subgrid schedule. Subgrid aware kernels will call cudaGetNextWorkItem to get items from the subgrid starting at workItemStart. Subgrid aware kernels will call cudaWaitWorkItem to ensure the associated subgrid’s dependencies are met before reading the related data. Each work item in a subgrid will be fetched and processed by only one worker (CTA/CGA) in the work group.
The application must not add dependencies to subgrids belonging to worker grids created after the worker grid the subgrid is joining. The application must not add dependencies to subgrids belonging to worksets from the owning worker grid created after the owning workset.
The easiest way to meet the dependency ordering requirement is to only add subgrid objects to the last created parent object. That is only add CUsubgridWorksets to the last created CUsubgridWorkerGrid and only add CUsubgrids to the last created CUsubgridWorkset.
See also
- Parameters
*subgrid_out – - returns the created subgrid
hWorkset – - workset to add the subgrid to
*hDependencies – - array of subgrids that this subgrid depends on
numDependencies – - size of the hDependencies array
workItemStart – - the three-dimensional start index of the first work item in the subgrid
workItemEnd – - the three-dimensional end index defining the exclusive range of the work items in the subgrid
- Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_OUT_OF_MEMORY
-
CUresult cuSubgridCreate_v2(CUsubgrid *subgrid_out, CUsubgridWorkset hWorkset, const CUsubgrid *hDependencies, const CUsubgridDependencyType *dependencyTypes, cuuint32_t numDependencies, CUoffset3D *workItemStart, CUoffset3D *workItemEnd)
create a subgrid with dependency type support
A subgrid is the fundametal dependency management object in the subgrid schedule. Subgrid aware kernels will call cudaGetNextWorkItem to get items from the subgrid starting at workItemStart. Subgrid aware kernels will call cudaWaitWorkItem to ensure the associated subgrid’s dependencies are met before reading the related data. Each work item in a subgrid will be fetched and processed by only one worker (CTA/CGA) in the work group.
The application must not add dependencies to subgrids belonging to worker grids created after the worker grid the subgrid is joining. The application must not add dependencies to subgrids belonging to worksets from the owning worker grid created after the owning workset.
The easiest way to meet the dependency ordering requirement is to only add subgrid objects to the last created parent object. That is only add CUsubgridWorksets to the last created CUsubgridWorkerGrid and only add CUsubgrids to the last created CUsubgridWorkset.
See also
- Parameters
*subgrid_out – - returns the created subgrid
hWorkset – - workset to add the subgrid to
*hDependencies – - array of subgrids that this subgrid depends on
*dependencyTypes – - array of subgrid dependency types, one per entry in
hDependenciesnumDependencies – - size of the hDependencies array
workItemStart – - the three-dimensional start index of the first work item in the subgrid
workItemEnd – - the three-dimensional end index defining the exclusive range of the work items in the subgrid
- Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_OUT_OF_MEMORY
-
CUresult cuSubgridWorkerGridCreate(CUsubgridWorkerGrid *workerGrid_out, CUsubgridSchedule hSchedule, const CUsubgridWorkerGridDesc *workerGridDesc)
create a subgrid worker grid
A subgrid worker grid is a worker launch operation. It specifies the grid of workers (CTAs/CGAs) that the subgrid schedule launches, and the kernel all of those workers execute. The schedule currently launches worker grids in creation order. Only subgrid aware kernels with cudaGetNextWorkItem system calls are supported.
Subgrids must be added to the worker grid as part of worksets for the worker grid to have workItems to work on.
Subgrid initialization expects to use at least 1 full warp, so blockDim.x * blockDim.y * blockDim.z >= 32 is required.
/sa ::cuSubgridScheduleCreate cuSubgridWorksetCreate
- Parameters
*workerGrid_out – - returns the created worker grid
hSchedule – - the subgrid schedule that is responsible for launching this worker grid
*workerGridDesc – - describes the worker grid to be created
- Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_OUT_OF_MEMORY
-
CUresult cuSubgridWorksetCreate(CUsubgridWorkset *workset_out, CUsubgridWorkerGrid hWorkerGrid, const void *userData, size_t userDataSize)
create a subgrid workset
Creates a subgrid workset with attached user data and assigns it to a worker grid. Worksets provide a means of elegantly fusing sets of device functions residing in a single module into a single kernel. The user data associated with a given workset can be used to direct what operation and/or base pointers the kernel should do or use when processing the contained subgrids.
A worker grid will start work on worksets in creation order. Subgrid workItems must be added to the workset for it to contain workItems.
The kernel can get a pointer to the user data using a GetWorksetUserData device function. The kernel may then use the read only proxy to read the user data with the ld.global.proxy::readonly ptx instruction.
/sa ::cuSubgridScheduleCreate cuSubgridCreate
- Parameters
*workset_out – - returns the created workset
hWorkerGrid – - the CUsubgridWorkerGrid that is responsible for performing all the work in the workset
userData – - pointer to user data to associate with the workset for fast retrieval (may be NULL if userDataSize is 0)
userDataSize – - size of user data in bytes (must not exceed 16 bytes)
- Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_OUT_OF_MEMORY