Complete concrete configuration for one compiled kernel launch. More...

#include <launch_config.h>

Data Fields
loomc_structure_type_t	type
	Structure type.
loomc_host_size_t	structure_size
	Size of this structure in bytes.
void *	next
	Unordered extension chain for additional launch properties.
loomc_dimension3_t	workgroup_count
	Number of workgroups dispatched along each dimension.
loomc_dimension3_t	workgroup_size
	Number of invocations in each workgroup along each dimension.
loomc_dimension3_t	workgroup_cluster_size
	Number of workgroups in each cooperative cluster dimension.
uint32_t	subgroup_size
	Number of invocations in each hardware subgroup.
uint64_t	workgroup_storage_bytes
	Total workgroup-local storage required per workgroup, in bytes.

Detailed Description

Complete concrete configuration for one compiled kernel launch.

Callers zero-initialize this structure, set type to LOOMC_STRUCTURE_TYPE_LAUNCH_CONFIG, set structure_size to sizeof(loomc_launch_config_t), and pass it to loomc_launch_config_program_invoke.

Field Documentation

◆ subgroup_size

uint32_t loomc_launch_config_t::subgroup_size

Number of invocations in each hardware subgroup.

Zero indicates that the compiled target has no fixed subgroup size.

◆ type

loomc_structure_type_t loomc_launch_config_t::type

Structure type.

Must be LOOMC_STRUCTURE_TYPE_LAUNCH_CONFIG when nonzero.

◆ workgroup_cluster_size

loomc_dimension3_t loomc_launch_config_t::workgroup_cluster_size

Number of workgroups in each cooperative cluster dimension.

Kernels that do not use clustered launch report {1, 1, 1}.

◆ workgroup_size

loomc_dimension3_t loomc_launch_config_t::workgroup_size

Number of invocations in each workgroup along each dimension.

Every dimension is nonzero after successful evaluation.

◆ workgroup_storage_bytes

uint64_t loomc_launch_config_t::workgroup_storage_bytes

Total workgroup-local storage required per workgroup, in bytes.

This is the final total capacity required by the compiled kernel, including storage derived from workload facts during compilation. Target adapters map this total requirement to their launch ABI; it is not necessarily an API's additional dynamic-shared-memory argument.

The documentation for this struct was generated from the following file:

launch_config.h