Cuda testing (#28)

Author: thewilsonator

dub.json

@@ -1,13 +1,26 @@
 {
 	"name": "dcompute",
 	"description": "Native Heterogeneous Computing for D",
-	"copyright": "Copyright © 2016, Nicholas Wilson",
+	"copyright": "Copyright © 2017, Nicholas Wilson",
 	"authors": ["Nicholas Wilson"],
     "license": "BSL-1.0",
 	"dependencies": {
         "derelict-cl"  : "~>2.0.0",
         "derelict-cuda": "~>2.0.1",
         "taggedalgebraic": "~>0.10.7"
 	},
-    "dflags" : ["-mdcompute-targets=cuda-350" , "-oq", "-betterC"]
+    "dflags" : ["-mdcompute-targets=cuda-210" ,"-oq", "-betterC"],
+	"configurations": [
+		{
+			"name": "library",
+			"targetType": "library",
+			"excludedSourceFiles": ["./source/tests/*"],
+		},
+		{
+			"name": "unittest",
+			"targetType": "executable",
+			"versions": ["DComputeTesting"],
+
+		},
+	]
 }

source/dcompute/driver/cuda/buffer.d

@@ -6,11 +6,40 @@ struct Buffer(T)
 {
     size_t raw;
 
+	// Host memory associated with this buffer
+	T[] hostMemory;
+
     this(size_t elems)
     {
         status = cast(Status)cuMemAlloc(&raw,elems * T.sizeof);
         checkErrors();
+		hostMemory = null;
     }
+
+	this(T[] arr)
+	{
+		status = cast(Status)cuMemAlloc(&raw,arr.length * T.sizeof);
+		checkErrors();
+		hostMemory = arr;
+	}
+	void copy(Copy c)()
+	{
+		static if (c == Copy.hostToDevice)
+		{
+			cuMemcpyHtoD(raw, hostMemory.ptr,hostMemory.length * T.sizeof);
+		}
+		else static if  (c == Copy.deviceToHost)
+		{
+			cuMemcpyDtoH(hostMemory.ptr,raw,hostMemory.length * T.sizeof);
+		}
+	}
+	alias hostArgOf(U : GlobalPointer!T) = raw; 
+	void release()
+	{
+		cuMemFree(raw);
+		raw = 0;
+		hostMemory = null;
+	}
 }
 
 alias bf = Buffer!float;

source/dcompute/driver/cuda/context.d

@@ -5,7 +5,7 @@ import dcompute.driver.cuda;
 struct Context
 {
     void* raw;
-    this(Device dev, uint flags)
+    this(Device dev, uint flags = 0)
     {
         status = cast(Status)cuCtxCreate(&raw, flags,dev.raw);
         checkErrors();

source/dcompute/driver/cuda/package.d

@@ -1,5 +1,6 @@
 module dcompute.driver.cuda;
 
+public import ldc.dcompute;
 public import derelict.cuda.driverapi;
 
 public import dcompute.driver.error;
@@ -14,10 +15,10 @@ public import dcompute.driver.cuda.platform;
 public import dcompute.driver.cuda.program;
 public import dcompute.driver.cuda.queue;
 
-enum CopySource
+enum Copy
 {
-    host,
-    device,
+    hostToDevice,
+    deviceToHost,
     array,
 }
 
@@ -27,3 +28,16 @@ enum MemoryBankConfig : int
     fourBytes,
     eightBytes,
 }
+template HostArgsOf(F) {
+    import std.meta, std.traits;
+    alias HostArgsOf = staticMap!(ReplaceTemplate!(Pointer, Buffer), Parameters!F);
+}
+private template ReplaceTemplate(alias needle, alias replacement) {
+    template ReplaceTemplate(T) {
+        static if (is(T : needle!Args, Args...)) {
+            alias ReplaceTemplate = replacement!(Args[1]);
+        } else {
+            alias ReplaceTemplate = T;
+        }
+    }
+}

source/dcompute/driver/cuda/platform.d

@@ -8,18 +8,21 @@ struct Platform
 {
     static void initialise(uint flags =0)
     {
+        DerelictCUDADriver.load();
         status = cast(Status)cuInit(flags);
         checkErrors();
     }
 
-    Device[] devices(A)(A a)
+    static Device[] getDevices(A)(A a)
     {
         int len;
         TypedAllocator!(A) allocator;
         status = cast(Status)cuDeviceGetCount(&len);
         checkErrors();
-        
-        Device[] ret = allocator.makeArray!(Device)(len);
+
+        //TODO:
+        //Device[] ret = allocator.makeArray!(Device)(len);
+		    Device[] ret = new Device[len];
         foreach(int i; 0 .. len)
         {
             status = cast(Status)cuDeviceGet(&ret[i].raw,i);

source/dcompute/driver/cuda/program.d

@@ -39,6 +39,7 @@ struct Program
         return ret;
     }
 
+	__gshared static Program globalProgram;
     //cuModuleLoadDataEx
     //cuModuleLoadFatBinary
 
@@ -47,7 +48,7 @@ struct Program
         status = cast(Status)cuModuleUnload(raw);
         checkErrors();
     }
-    static Program globalProgram;
+    
     //TODO: linkstate
 }
 

source/dcompute/driver/cuda/queue.d

@@ -57,14 +57,21 @@ struct Queue
         cuStreamAddCallback
     }*/
 
-    auto enqueue(alias k)(uint[3] _grid, uint[3] _block, uint _sharedMem)
+    auto enqueue(alias k)(uint[3] _grid, uint[3] _block, uint _sharedMem = 0)
     {
         static struct Call
         {
             Queue q;
             uint[3] grid, block;
             uint sharedMem;
 
+			this(Queue _q,uint[3] _grid, uint[3] _block, uint _sharedMem)
+			{
+				q= _q;
+				grid = _grid;
+				block = _block;
+				sharedMem = _sharedMem;
+			}
             //TODO integrate evnts into this.
             void opCall(HostArgsOf!(typeof(k)) args)
             {
@@ -76,7 +83,7 @@ struct Queue
                 }
 
                 status = cast(Status)
-                        cuLaunchKernel(kernel,
+                        cuLaunchKernel(kernel.raw,
                                        grid[0], grid[1], grid[2],
                                        block[0],block[1],block[2],
                                        sharedMem,

source/dcompute/driver/error.d

@@ -133,14 +133,14 @@ version (D_BetterC)
 {
     void delegate (Status) nothrow @nogc onDriverError;
     immutable void delegate (Status) nothrow @nogc defaultOnDriverError;
-	static this()
-	{
-		defaultOnDriverError = (Status _status)
-		{
-			import core.stdc.stdio : stderr;
-			import std.conv : to;
-			fprintf(stderr,"*** DCompute driver error:%s\n",
-					_status.to!(string).toStringz);
+    static this()
+    {
+        defaultOnDriverError = (Status _status)
+        {
+            import core.stdc.stdio : stderr;
+            import std.conv : to;
+            fprintf(stderr,"*** DCompute driver error:%s\n",
+                   _status.to!(string).toStringz);
 		};
 	}
 }
@@ -161,13 +161,13 @@ else
     }
     void delegate(Status) onDriverError;
     immutable void delegate(Status) defaultOnDriverError;
-	static this()
+    static this()
 	{
-		defaultOnDriverError = (Status _status)
-		{
-			throw new DComputeDriverException(_status);
-		};
-	}
+        defaultOnDriverError = (Status _status)
+        {
+            throw new DComputeDriverException(_status);
+        };
+    }
 }
 static this() { onDriverError = (Status _status) { defaultOnDriverError(_status);};}
 // Thread local status

source/dcompute/tests/dummykernels.d

@@ -5,10 +5,14 @@ pragma(LDC_no_moduleinfo);
 import ldc.dcompute;
 import dcompute.std.index;
 
-@kernel void saxpy(GlobalPointer!(float) a, GlobalPointer!(float) b, float c )
+@kernel void saxpy(GlobalPointer!(float) res,
+				   float alpha,GlobalPointer!(float) x,
+				   GlobalPointer!(float) y, 
+				   size_t N)
 {
     auto i = GlobalIndex.x;
-    a[i] = b[i] +c;
+	if (i >= N) return;
+    res[i] = alpha*x[i] + y[i];
 }
 
 alias aagf = AutoIndexed!(GlobalPointer!(float));

source/dcompute/tests/main.d

@@ -1,14 +1,20 @@
+version = DComputeTestCUDA;
+
 import dcompute.tests.test;
 
 import std.stdio;
 import std.traits;
 import std.meta;
-
+import std.exception : enforce;
 import std.experimental.allocator;
+import std.array;
 
 import dcompute.tests.dummykernels : saxpy;
 
-import dcompute.driver.ocl120;
+version(DComputeTestOpenCL)
+	import dcompute.driver.ocl120;
+version(DComputeTestCUDA)
+	import dcompute.driver.cuda;
 
 ubyte[] readBinaryFromDisk()
 {
@@ -18,16 +24,56 @@ ubyte[] readBinaryFromDisk()
 float[] somedata() { return [1,2,3,4,5]; }
 int main(string[] args)
 {
-    auto platform = Platform.getPlatforms(theAllocator)[0];
-    auto devices  = platform.getDevices(theAllocator);
-    auto plist    = propertyList!(Context.Properties)(Context.Properties.platform);
-    auto context  = Context(devices,plist);
-    Program.globalProgram = context.createProgramFromSPIR(theAllocator,devices,readBinaryFromDisk());
-    auto queue    = context.createQueue(devices[0],Queue.Properties.outOfOrderExecution);
-    auto data     = somedata();
-    auto buf      = context.createBuffer!(float)(Memory.Flags.none,data);
-    Event e = queue.enqueue!(saxpy)([5])(buf,buf,42.0);
-    e.wait();
+	version(DComputeTestOpenCL)
+	{
+		auto platform = Platform.getPlatforms(theAllocator)[0];
+		auto devices  = platform.getDevices(theAllocator);
+		auto plist    = propertyList!(Context.Properties)(Context.Properties.platform);
+		auto ctx      = Context(devices,plist);
+		Program.globalProgram = context.createProgramFromSPIR(theAllocator,devices,readBinaryFromDisk());
+		auto queue    = context.createQueue(devices[0],Queue.Properties.outOfOrderExecution);
+		auto data     = somedata();
+		auto buf      = context.createBuffer!(float)(Memory.Flags.none,data);
+		Event e = queue.enqueue!(saxpy)([5])(buf,buf,42.0);
+		e.wait();
+	}
+	version(DComputeTestCUDA)
+	{
+		Platform.initialise();
+		
+		auto devs	= Platform.getDevices(theAllocator);
+		auto dev	= devs[0]; 
+		auto ctx	= Context(dev); scope(exit) ctx.detach();
+
+		// Change the file path to match your GPU.
+		Program.globalProgram = Program.fromFile("./kernels_cuda210_64.ptx");
+		auto q = Queue(false);
+
+		enum size_t N = 128;
+		float alpha = 5.0;
+		float[N] res, x,y;
+		foreach (i; 0 .. N)
+		{
+			x[i] = N - i;
+			y[i] = i * i;
+		}
+		Buffer!(float) b_res, b_x, b_y;
+		b_res	=  Buffer!(float)(res[]); scope(exit) b_res.release();
+		b_x		=  Buffer!(float)(x[]);   scope(exit) b_x.release();
+		b_y		=  Buffer!(float)(y[]);   scope(exit) b_y.release();
+
+		b_x.copy!(Copy.hostToDevice);
+		b_y.copy!(Copy.hostToDevice);
+
+		q.enqueue!(saxpy)
+			([N,1,1],[1,1,1])
+			(b_res,alpha,b_x,b_y, N);
+		b_res.copy!(Copy.deviceToHost);
+		foreach(i; 0 .. N)
+			enforce(res[i] == alpha * x[i] + y[i]);
+		writeln(res[]);
+	}
+
     return 0;
 }