1 /** 2 * Contains compiler-recognized special symbols for dcompute. 3 * 4 * Copyright: Authors 2017 5 * License: $(LINK2 http://www.boost.org/LICENSE_1_0.txt, Boost License 1.0) 6 * Authors: Nicholas Wilson 7 */ 8 9 module ldc.dcompute; 10 11 enum ReflectTarget : uint 12 { 13 // These numbers MUST match DcomputeTarget::target in LDC. 14 Host = 0, 15 OpenCL = 1, 16 CUDA = 2, 17 } 18 /** 19 * The pseudo conditional compilation function. 20 * returns true if t and _version match the target and it's version. 21 * think __ctfe but per target, codegen time conditional compilation. 22 * as oppsed to compiletime/runtime. 23 * arguments MUST be compiletime constants 24 * valid values of _version are for OpenCL 100 110 120 200 210 25 * and for CUDA are x*100 + y*10 for x any valid values of sm x.y 26 * use 0 as a wildcard to match any version. 27 28 * This is mostly used for selecting the correct intrinsic for the 29 * given target and version, but could also be used to tailor for 30 * performance characteristics. See dcompute.std.index for an example 31 * of how to select intrinsics. 32 */ 33 pure nothrow @nogc 34 extern(C) bool __dcompute_reflect(ReflectTarget t, uint _version = 0); 35 36 ///Readability aliases for compute 37 enum CompileFor : int 38 { 39 deviceOnly = 0, 40 hostAndDevice = 1 41 } 42 43 /++ 44 + When applied to a module, specifies that the module should be compiled for 45 + dcompute (-mdcompute-targets=<...>) using the NVPTX and/or SPIRV backends of 46 + LLVM. 47 + 48 + Examples: 49 + --- 50 + @compute(CompileFor.deviceOnly) module foo; 51 + import ldc.dcompute; 52 + --- 53 +/ 54 struct compute 55 { 56 CompileFor codeProduction = CompileFor.deviceOnly; 57 } 58 59 /++ 60 + Mark a function as a 'kernel', a compute API (CUDA, OpenCL) entry point. 61 + Equivalent to __kernel__ in OpenCL and __global__ in CUDA. 62 + 63 + Examples: 64 + --- 65 + @compute(CompileFor.deviceOnly) module foo; 66 + import ldc.dcompute; 67 + 68 + @kernel void bar() 69 + { 70 + //... 71 + } 72 + --- 73 +/ 74 private struct _kernel {} 75 enum kernel = _kernel(); 76 77 /++ 78 + DCompute has the notion of adress spaces, provide by the magic structs below. 79 + The numbers are for the DCompute virtual addess space and are translated into 80 + the correct address space for each DCompute backend (SPIRV, NVPTX). 81 + The table below shows the equivalent annotation between DCompute OpenCL and CUDA 82 + 83 + DCompute OpenCL Cuda 84 + Global __global __device__ 85 + Shared __local __shared__ 86 * Constant __constant __constant__ 87 + Private __private __local__ 88 + Generic __generic (no qualifier) 89 +/ 90 struct Pointer(AddrSpace as, T) 91 { 92 T* ptr; 93 alias ptr this; 94 } 95 96 struct Variable(AddrSpace as, T) 97 { 98 T val; 99 alias val this; 100 } 101 102 enum AddrSpace : uint 103 { 104 Private = 0, 105 Global = 1, 106 Shared = 2, 107 Constant = 3, 108 Generic = 4, 109 } 110 111 alias PrivatePointer(T) = Pointer!(AddrSpace.Private, T); 112 alias GlobalPointer(T) = Pointer!(AddrSpace.Global, T); 113 alias SharedPointer(T) = Pointer!(AddrSpace.Shared, T); 114 alias ConstantPointer(T) = Pointer!(AddrSpace.Constant, immutable(T)); 115 alias GenericPointer(T) = Pointer!(AddrSpace.Generic, T); 116 117 // N.B private variables are declared on the stack and so cannot be declared 118 // at module scope. 119 // No variables exist in the generic address space, it is purely for generalised 120 // pointers to point to. 121 // 122 // The __gshared below does not work. It is kludged into place in `DtoResolveVariable` 123 124 alias Global(T) = /*__gshared*/ Variable!(AddrSpace.Global, T); 125 alias Shared(T) = shared Variable!(AddrSpace.Shared, T); 126 alias Constant(T) = immutable Variable!(AddrSpace.Constant, T);