LDC CUDA and SPIRV

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-nvidia-cuda - nvptx64-nvidia-cuda

Needs Work

Merge the CodeGenVisitors and StatementsVisitors as these were separated for ease of development, but have the vast majority of code the same.

Refactor the isFromDCompute functions and the similar functions in ida.cpp

The ABI code should be moved to the rest of the ABI code

Remove IF_LOG s that are left over from me attempting to figure out what is going on. Basically anything with a %p format specifier.

Do the TODOs.

Attribute code adaption from clang/llvm.

Issues

Need to create a CMake option to link with a SPIRV compatible LLVM to allow use with unmodified llvm.

Currently segfaults if -singleobj is passed and the first file on the command line is marked @compute

Desired Features not yet implemented

Suport for OpenCL Images and pipes (type resolution and metadata)

Compilation process

A D module declaration is tagged @dcompute.attributes.compute

@compute gets picked up prior to the start of module code generation and is deferred until after the rest of the host code generation is done.

codegen is then done foreach of the command line specified DCompute targets e.g. -mdcompute-targets=ocl-120,cuda-350 for OpenCL 1.2 and CUDA SM3.5

during statement codegen calls to "if(__dcompute_reflect(target t,uint version))" are selectively codegen'd for the appropriate target. during Type translation all types in dcompute.types are altered to the appropriate type e.g.

GLobalPointer!(float) becomes {float addrspace(1)*}  

Relevant metadata is attached to the generated module. This includes things like specifying which function are kernels (entry points)

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.

The following code

@compute module foo;
import dcompute.attrbutes;
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 IR semantically equivalent to 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}

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.

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.

Resources

KHRONOS's SPRIV capable LLVM

OpenCL C++ standard library

KHRONOS's SPIR-V clang

Compiling CUDA C/C++ with LLVM

CGO 2016 gpucc tutorial

A Journey of OpenCL 2.0 Development in Clang

The Diff

this section details the differences made to LDC

driver/cl_options.h	added mDcomputeTargets command line option
driver/linker.cpp	in linkObjToBinaryGcc: skip files ending with .spv and .ptx
driver/main.cpp	add includes, in hideLLVMOptions: hide spirv options , in registerPredefinedVersions: version = D_DCompute, in cppmain emit for dcompute
driver/toobj.cpp	add includes, in codegenModule: if module is spirv pass it to createSPIRVWriterPass, if is ptx Target->addPassesToEmitFile but always assembly
gen/functions.cpp	disable debuginfo generation for DCompute, no uwtable, alloca point => allocaPoint ( I think this fixes a crash)
gen/llvmhelpers.cpp	disable debuginfo generation for DCompute
gen/optimizer.cpp	dont optimise spirv modules as turning GEPs into extracts causes crashes.
gen/statements.cpp	comments
gen/tollvm.cpp	DtoBitCast
gen/uda.cpp	_kernel and _compute, isDComputeAttibutes, isFromDComputeAttibutes, getDComputeAttributesStruct, hasKernelAttr, hasComputeAttr
ir/irtypestruct.cpp	all
gen/irstate.{h,cpp}	all

---

Pull Requests

this section details the pull requests associated with this project.

Name	Number	Status	Depends On	Description
Dcompute	1786	Superseded	None	the original PR, now obsolete, superseded by the following PRs.
Cmake	1795	Not working yet	None	Adds the #define LDC_WITH_DCOMPUTE_SPIRV . Works. link with libLLVMSPRIVlib.a don't think this works. Equivalent to the changes to CMakeLitst.txt in 1786.
Attributes	1800	Merged	None	adds the magic attributes '@compute(int)' and '@kernel'. REQUIRES associated druntime changes, see the PR. Euqivlent to the changes in gen/uda.* in 1786. CONFLICTS 1877 : severity trivial.
Addrspace	1801	Pending merge	None	Correctly handle pointers with the llvm addrspace qualification, namely don't allow bit casting between pointers of different addrspaces. Necessary due to the way the addrspace qualification is attached. FILES gen/dvalue.cpp and gen/tollvm.* DEFERED optimise StripAddrSpaces when gDComputeTarget is added and add an early return if it is false.
Toobj and optimiser	1879	Merged	None	dont optimise SPIR-V as turning GEPs into extracts causes the back end to assert. Always use llvm::TargetMachine::CGFT_AssemblyFile for ptx, if we don't create target machine will return null. llvm::createSPIRVWriterPass(out)->runOnModule(m) is the way to make output happen for SPIR-V as there is no target machine for SPIR-V. FILES driver/toobj.cpp and gen/optimizer.cpp.
ABI	1878	Merged	1800	Sets the calling conventions for @kernel and non-@kernel function correctly. FILES dcompute/abi* -> gen/abi* . FUTURE put gen/abi* into gen/abi/*
target + codegenmanager	1922	Pending merge	ABI	Targets hold the address space mapping and other information about the backend, set gTargetMachine and add required metadata to the module being codegen'd. CONTENTION location of the magic types (currently 'dcompute.types', 'ldc.dcomputetypes'?) . FILES dcompute/codegen* -> driver/dcomputecodegen* , dcompute/target* -> gen/dcomputetarget* . DEFER update call to Declaration_codegen when its signature changes. FUTURE handle restrict if/when it becomes a thing
irtypestruct	None	Dependencies not met	1922	Attach the addrspace qualification to the magic type above. FILES ir/irtypestruct.cpp
code generation	None	Dependencies not met	1922	Statement and Declaration visitors.
main	None	Dependencies not met	1922	main.cpp
druntime	None	PR Not opened	None	ldc.attributes (@compute + @kernel) + ldc.dcomputetypes (Pointer)
lit tests	None	Dependencies not met	1801 + 1922	Lit based IR tests