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.

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.

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 hDependencies

  • 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 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