LDC CUDA and SPIRV
Contents
About
This page is about the requirements and considerations for getting LDC to target the NVPTX and SPIR backends of LLVM i.e. triples - spir-unknown-unknown - spir64-unknown-unknown - nvptx-unknown-unknown - nvptx64-nvidia-cuda
Compilation process
A D module declaration is tagged @ldc.attributes.compute
@compute gets picked up by gen/uda.cpp:applyFuncDeclUDAs
abort the codegenning of the current host llvm module and instead codegen into the (command line) specified targets (CUDA and/or OpenCL)
Run a "conditional compilation pass" (akin to the nvvm_reflect pass) turning calls to "__dcompute_reflect(target t,uint version)" with constants followed by a DCE pass to remove code not intended for the target.
Run a type replacement pass to replace GlobalPointer et. al with a pointer in the correct address space (these differ between CDUA and OpenCL) and other magic types such as images ( if I can't find a way to get regular buffers to be cached in texture cache and thus roll our own)
Attach the relevant metadata to the generated module
In the case of CUDA Pass the module off to the backend to produce .ptx files, or in the case of OpenCL targeting SPIR-V schedule the SPIR-V pass to generate the .spv file.
Address Spaces
CUDA and OpenCL both have notions of regions of memory:
Private. this is memory used by a given thread of execution and contains its stack and registers Global. Memory that is global to the device Local. Memory that is local to a work group (aka warp wave), a group of threads. Constant. Memory (re)writable only by the host, between execution of the batch of kernels.
This is mapped to the LLVM concept of address spaces. In SPIR and CUDA these map to the above regions. In addition to the above there is a fifth address space that pointers may point to which is generic. These differ between SPIR and CUDA handle them with magic structs that get translated (see above).
Note that pointers have two associated address spaces: the space of residence and the pointer space, e.g. one can have a local pointer to global memory, i.e. the pointer resides in local memory but points to somewhere in global memory.
TODO: how useful is this (other than for private which is required)? how should this interact with D default TLS? Don't allow global variables at all in modules marked @compute, figure out how to do a shared_alloc for allocating shared memory.
The following code
@compute module foo;
import dcompute.attrbutes;//publicly imports ldc.attributes : compute
import dcompute.std.transcendental;
extern(C) // for easier mangling
@kernel void bar(GlobalPointer!(Vector!(float,4)) fp, float b)
{
if(__dcompute_reflect(target.OpenCL))
*fp += float4(b);
else if(__dcompute_reflect(target.CUDA))
*fp = float4(cos(b);
}
when targeting OpenCL should produce the following IR
declare void @bar(<float x 4> addrspace(1)* %fp, float %b)
{
%1 = load <float x 4> addrspace(1)* %fp
%2 = add %b %1
store %2 <float x 4> addrspace(1)* %fp
}
!opencl.kernels = {!1}
!1 = {void(<float x 4> addrspace(1)* %fp, float %b) *@bar, ... }
and when targeting CUDA should produce
declare void @bar(<float x 4> addrspace(1)* %fp, float %b)
{
%1 = load <float x 4> addrspace(1)* %fp
%2 = call float cos(b)
store %2 <float x 4> addrspace(1)* %fp
}
!nvvm.annotations = {!1}
!1 = {void(<float x 4> addrspace(1)* %fp, float %b) *@bar, !"kernel", i32 1}
Restrictions
In the environment of CUDA and OpenCL the nature of execution is more restricted than on CPU. In short, there are no exceptions (what to do about assert?), no function pointers (all template delegate parameters MUST be inlined), no recursion either direct or indirect, there is no I/O, no C or D runtime and no OS. However synchronisation primitives such as fence are still available, as are atomics. The expected way to achieve this is to have a transitive attribute (@kernel) that enforces these restriction similar to @nogc, nothrow and pure. If we were to disallow non-builtin globals and make the builtin ones immutable we may be able to get away with @kernel being equivalent to @nogc pure nothrow.
Ranges
Much of the programming power of D comes from ranges. The paradigm does not transfer perfectly to CUDA and OpenCL but should still be usable.
First it is useful to briefly cover different types of ranges.
Generative ranges: these do not take a range as an input put produce one. The produced range is not random access. Examples include iota, recurrence relations such as Fibonacci.
Transformative ranges: These take ranges as input and return them as output. The output(s) are not necessarily random access (but will generally be if the input is). In the context of GPGPU it is useful to further categorise these as to the relation of the number of elements of input when compared to the number of outputs. Of particular interest are ranges that perform an n:n mapping as these can be chained from within the same kernel. The obvious example here is `map`. Some ranges e.g. `filter` do not preserve this and will have to be dealt with differently e.g. changing elements that do not pass the predicate to a sentinel value (e.g. NaN), or stop the range chain and require new buffers to allocated and different kernels to continue the chain.
Consuming Ranges: these take range(s) as input and either return void or a scalar. Examples include min,max and reduce.
In order to be chained within the same kernel the input range(s) need to have the same number of elements as the output and be random access, as the quintessential range on GPUs are arrays. Also the notion of `.save`ing a range doesn't translate.
Vector types
Currently LDC vector types use adapt __vector which rejects invalid types. this doesn't work for thing like float16 that are too big.
Builtins & Intrinsics
CUDA and OpenCL expose a lot of builtin variables (work size, GlobalLinearId) as well as intrinsic functions various FMA as well as other types like images pipes and events.
Metadata & Special Function attributes
the LLVM IR forms of SPIRV and PTX hold a lot of magic metadata.
The form of the metal data can be found in the test modules for the codegen of clang.
All spirv kernels have the attribute spirv_kernel and all CUDA device function have the attribute ptx_device.
Standard Library
As part of providing SPRI-V and CUDA as backends we will need to provide a standard library of function that meet the restriction criteria imposed by the environment.
A non exhaustive list is
- vector operations and functions. both for fixed length and (run time) variable length.
- work group functions e.g. reduce search sort
- provide the builtin variables.
- provide function to deal with the special objects: images, pipes and events.
Misc
Currently the KHRONOS branch of LLVM that supports SPIR-V only supports OpenCL. It is worth considering making supporting GLSL easy after the fact.
Ideally we should make the interface for CUDA and OpenCL as similar and consistent as possible. The higher level library code should be agnostic. The functions will be able to be introspected and so should be easy.
SPIR-V has the notion of capabilities.These include things like, half and double precision floating support, 64 bit integers, atomics, 64bit atomics, images and pipes.
