6.39. 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 coresponding instantiated graph.

Functions

__host__ cudaError_t cudaSubgridCreate(cudaSubgrid_t *subgrid_out, cudaSubgridWorkset_t hWorkset, const cudaSubgrid_t *hDependencies, unsigned int numDependencies, struct cudaPos workItemStart, dim3 workItemEnd)

Create a subgrid.

__host__ cudaError_t cudaSubgridCreate_v2(cudaSubgrid_t *subgrid_out, cudaSubgridWorkset_t hWorkset, const cudaSubgrid_t *hDependencies, const cudaSubgridDependencyType *dependencyTypes, unsigned int numDependencies, struct cudaPos workItemStart, dim3 workItemEnd)

Create a subgrid with dependency type support.

__host__ cudaError_t cudaSubgridWorkerGridCreate(cudaSubgridWorkerGrid_t *workerGrid_out, cudaSubgridSchedule_t hSchedule, const cudaSubgridWorkerGridDesc_t *workerGridDesc)

Create a subgrid worker grid.

__host__ cudaError_t cudaSubgridWorksetCreate(cudaSubgridWorkset_t *workset_out, cudaSubgridWorkerGrid_t hWorkerGrid, const void *userData, size_t userDataSize)

Create a subgrid workset.

6.39.1. Functions

__host__ cudaError_t cudaSubgridCreate(cudaSubgrid_t *subgrid_out, cudaSubgridWorkset_t hWorkset, const cudaSubgrid_t *hDependencies, unsigned int numDependencies, struct cudaPos workItemStart, dim3 workItemEnd)

Create a subgrid.

Creates a subgrid within a workset. A subgrid defines a range of work items with shared dependency structure. Subgrids can depend on other subgrids to create execution dependencies.

See also

cudaSubgridWorkerGridCreate, cudaSubgridWorksetCreate

Parameters

  • subgrid_out – - Returns the created subgrid handle

  • hWorkset – - The workset to add the subgrid to

  • hDependencies – - Array of subgrids that this subgrid depends on

  • numDependencies – - Number of dependencies in hDependencies array

  • workItemStart – - 3D starting work item index (inclusive). Uses cudaPos which defaults unspecified dimensions to 0, making 1D initialization natural: {N} becomes (N, 0, 0).

  • workItemEnd – - 3D ending work item index (exclusive). Uses ::dim3 which defaults unspecified dimensions to 1, making 1D initialization natural: {N} becomes (N, 1, 1).

Returns

cudaSuccess, cudaErrorInvalidValue, cudaErrorMemoryAllocation

__host__ cudaError_t cudaSubgridCreate_v2(cudaSubgrid_t *subgrid_out, cudaSubgridWorkset_t hWorkset, const cudaSubgrid_t *hDependencies, const cudaSubgridDependencyType *dependencyTypes, unsigned int numDependencies, struct cudaPos workItemStart, dim3 workItemEnd)

Create a subgrid with dependency type support.

Creates a subgrid within a workset. A subgrid defines a range of work items with shared dependency structure. Subgrids can depend on other subgrids to create execution dependencies.

See also

cudaSubgridWorkerGridCreate, cudaSubgridWorksetCreate

Parameters

  • subgrid_out – - Returns the created subgrid handle

  • hWorkset – - The 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 – - Number of dependencies in hDependencies array

  • workItemStart – - 3D starting work item index (inclusive). Uses cudaPos which defaults unspecified dimensions to 0, making 1D initialization natural: {N} becomes (N, 0, 0).

  • workItemEnd – - 3D ending work item index (exclusive). Uses ::dim3 which defaults unspecified dimensions to 1, making 1D initialization natural: {N} becomes (N, 1, 1).

Returns

cudaSuccess, cudaErrorInvalidValue, cudaErrorMemoryAllocation

__host__ cudaError_t cudaSubgridWorkerGridCreate(cudaSubgridWorkerGrid_t *workerGrid_out, cudaSubgridSchedule_t hSchedule, const cudaSubgridWorkerGridDesc_t *workerGridDesc)

Create a subgrid worker grid.

Creates a worker grid within a subgrid schedule. The worker grid defines the kernel and launch configuration that will process work items.

See also

cudaSubgridWorksetCreate, cudaSubgridCreate

Parameters

  • workerGrid_out – - Returns the created worker grid handle

  • hSchedule – - The subgrid schedule to add the worker grid to

  • workerGridDesc – - Descriptor containing kernel and launch parameters

Returns

cudaSuccess, cudaErrorInvalidValue, cudaErrorMemoryAllocation

__host__ cudaError_t cudaSubgridWorksetCreate(cudaSubgridWorkset_t *workset_out, cudaSubgridWorkerGrid_t 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 group subgrids together and can carry user-defined data for processing.

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.

See also

cudaSubgridWorkerGridCreate, cudaSubgridCreate

Parameters

  • workset_out – - Returns the created workset handle

  • hWorkerGrid – - The worker grid that will process this workset

  • userData – - Pointer to user data to associate with the workset (may be NULL if userDataSize is 0)

  • userDataSize – - Size of user data in bytes (must not exceed 16 bytes)

Returns

cudaSuccess, cudaErrorInvalidValue, cudaErrorMemoryAllocation