1 module dnv.compiler; 2 3 import std.conv; 4 import std.string; 5 6 import dnv.driver; 7 import dnv.typechecker; 8 import dnv.error; 9 10 11 struct Code { 12 /* 13 FIXME: template support 14 - see: /usr/local/cuda/samples/0_Simple/simpleTemplates_nvrtc 15 */ 16 immutable qualifier = `extern "C" __global__ `; 17 immutable returnType = "void "; 18 const string name; 19 const string args; 20 const string source; 21 22 this(in string nameStr, in string argumentsStr, in string bodyStr) { 23 name = nameStr; 24 args = argumentsStr; 25 source = qualifier ~ returnType ~ nameStr ~ 26 "(" ~ args ~ ")" ~ 27 "{" ~ bodyStr ~ "}"; 28 } 29 } 30 31 unittest { 32 auto saxpy = Code( 33 "saxpy", "float *A, float *B, float *C, int numElements", ` 34 int i = blockDim.x * blockIdx.x + threadIdx.x; 35 if (i < numElements) C[i] = A[i] + B[i]; 36 `); 37 assert(saxpy.source == 38 `extern "C" __global__ void saxpy(float *A, float *B, float *C, int numElements){ 39 int i = blockDim.x * blockIdx.x + threadIdx.x; 40 if (i < numElements) C[i] = A[i] + B[i]; 41 }`); 42 } 43 44 45 struct UnsafeCompiler { 46 static void build(void* vfunc, Code c) { 47 check(compile(vfunc, c.name, c.source)); 48 } 49 void assertArgs(Args...)(Args args) {} 50 } 51 52 struct StaticCompiler(Code c) { 53 immutable Code code = c; 54 immutable cargs = c.args; 55 static void build(void* vfunc, Code c) { 56 check(compile(vfunc, c.name, c.source)); 57 } 58 void assertArgs(Args...)(Args targs) { 59 // FIXME: cannot call 60 staticAssert!(AssignableArgTypes, cargs)(targs); 61 } 62 } 63 64 struct SimpleLauncher { 65 uint[3] grids = [256, 1, 1]; 66 uint[3] blocks; 67 68 void setup(Args...)(Args targs) { 69 uint bx = to!uint((grids[0] + targs[0].length - 1) / grids[0]); 70 blocks = [bx, 1, 1]; 71 } 72 }