SPIR-V Environment: Type of Built-In Variables

[bashbaug]

PR #278 clarifies that all built-in variables must be in the input storage class, but does not describe what the type of each built-in variable must be. We should describe the supported types for each built-in variable.

Some questions to answer:

• For built-in vectors that can be vectors for multi-dimensional ND-ranges, are scalars valid if the ND-range is known to be 1D?

• I'm pretty sure all of the built-in variables are currently integers, but we need to define how big the integers representing each built-in variable may be. Specifically, are there cases where a built-in variable must be 64-bits? Are there cases where a built-in variable may be 16-bits?

[pierremoreau]

Specifically, are there cases where a built-in variable must be 64-bits?

Yes, all the built-ins that are of type size_t such as get_global_id() & co, since (regarding size_t and some other types)

They are 32-bit types if the value of the CL_DEVICE_ADDRESS_BITS device query is 32-bits, and 64-bit types if the value of the query is 64-bits.

Note that get_work_dim() does return an uint, so on 64-bit devices you can end up with built-ins of both 32- and 64-bits sizes.

[bashbaug]

Summary Table:

OpenCL C Built-in	SPIR-V Built-In	OpenCL C Return Type
get_num_groups(dim)	NumWorkgroups	size_t
get_local_size(dim)	WorkgroupSize	size_t
get_group_id(dim)	WorkgroupId	size_t
get_num_groups(dim)	NumWorkgroups	size_t
get_local_size(dim)	WorkgroupSize	size_t
get_group_id(dim)	WorkgroupId	size_t
get_local_id(dim)	LocalInvocationId	size_t
get_global_id(dim)	GlobalInvocationId	size_t
get_local_linear_id	LocalInvocationIndex	size_t
get_work_dim(dim)	WorkDim	uint
get_global_size(dim)	GlobalSize	size_t
get_enqueued_local_size(dim)	EnqueuedWorkgroupSize	size_t
get_global_offset(dim)	GlobalOffset	size_t
get_global_linear_id	GlobalLinearId	size_t
get_sub_group_size	SubgroupSize	uint
get_max_sub_group_size	SubgroupMaxSize	uint
get_num_sub_groups	NumSubgroups	uint
get_enqueued_num_sub_groups	NumEnqueuedSubgroups	uint
get_sub_group_id	SubgroupId	uint
get_sub_group_local_id	SubgroupLocalInvocationId	uint
get_sub_group_eq_mask	SubgroupEqMask	uint4
get_sub_group_ge_mask	SubgroupGeMask	uint4
get_sub_group_gt_mask	SubgroupGtMask	uint4
get_sub_group_le_mask	SubgroupLeMask	uint4
get_sub_group_lt_mask	SubgroupLtMask	uint4

So, at the very least we need to define three possible built-in variable types:

• For the OpenCL C built-ins that return size_t, the corresponding SPIR-V built-in must be:
  • An OpTypeInt with Width equal to 32 if the Addressing Model declared in OpMemoryModel is Physical32,
  • An OpTypeInt with Width equal to 64 if the Addressing Model declared in OpMemoryModel is Physical64,
• For the OpenCL C built-ins that return uint, the corresponding SPIR-V built-in must be an OpTypeInt with Width equal to 32.
• For the OpenCL C built-ins that return uint4, the corresponding SPIR-V built-in must be an OpTypeVector with Component Count equal to 4 and Component Type equal to an OpTypeInt with Width equal to 32.

The OpenCL C built-ins that do not accept a dim arg must be scalar types.

If we want to be consistent with the SPIR-V LLVM Translator, it looks like all of the OpenCL C built-ins that accept a dim argument must be vector types with three components.

The main things we need to decide are:

• Should we require that all OpenCL C built-ins that accept a dim argument must be vector types with three components in SPIR-V? Or, do we want to allow scalar types for known 1D dispatches, or vector types with an arbitrary number of components?

• Do we want to require the exact bit width for the integer types, or do we want to allow smaller integer types also?

For now, I will create a PR documenting the types that the SPIR-V LLVM Translator is generating, then we can consider whether to allow other types also.

[alycm]

https://github.com/KhronosGroup/OpenCL-Docs/pull/526 is merged, it does the above but the sub-group entries are commented out here: https://github.com/KhronosGroup/OpenCL-Docs/commit/cee24879a353834e9afc575813df1fe37a509cbf#diff-2c00fa956c00e5219de5612bcedd4ad4b8eae6e6879069843e6cc10d91182250R483-R502

Either this issue can be closed leaving things as they are, or those lines can be uncommented.

[bashbaug]

Closing as discussed in the August 1st teleconference. We can revisit this issue if needed, but the current SPIR-V environment specification is documenting the current behavior of the SPIR-V LLVM Translator and hence what most SPIR-V consumers are expecting.

Note that the commented out lines mentioned above are described in the cl_khr_subgroup_ballot description, hence no action is required for them.