added tests for LaunchAndSync

Author: chewxy

batch_test.go

@@ -1,6 +1,9 @@
 package cu
 
-import "testing"
+import (
+	"testing"
+	"unsafe"
+)
 
 func TestAttributes(t *testing.T) {
 	devices, _ := NumDevices()
@@ -40,3 +43,82 @@ func TestAttributes(t *testing.T) {
 		t.Errorf("Expected ComputeCapabilityMinor to be %v. Got %v instead", min, attrs[2])
 	}
 }
+
+func TestLaunchAndSync(t *testing.T) {
+	devices, _ := NumDevices()
+
+	if devices == 0 {
+		return
+	}
+
+	var err error
+	var ctx Context
+	var mod Module
+	var fn Function
+
+	d := Device(0)
+	if ctx, err = d.MakeContext(SchedAuto); err != nil {
+		t.Fatal(err)
+	}
+
+	a := make([]float32, 1000)
+	b := make([]float32, 1000)
+	for i := range b {
+		a[i] = 1
+		b[i] = 1
+	}
+
+	size := int64(len(a) * 4)
+
+	var memA, memB DevicePtr
+	if memA, err = MemAlloc(size); err != nil {
+		t.Fatalf("Failed to allocate for a: %v", err)
+	}
+	if memB, err = MemAlloc(size); err != nil {
+		t.Fatalf("Failed to allocate for b: %v", err)
+	}
+
+	if err = MemcpyHtoD(memA, unsafe.Pointer(&a[0]), size); err != nil {
+		t.Fatalf("Failed to copy memory from a: %v", err)
+	}
+
+	if err = MemcpyHtoD(memB, unsafe.Pointer(&b[0]), size); err != nil {
+		t.Fatalf("Failed to copy memory from b: %v", err)
+	}
+
+	if mod, err = LoadData(add32PTX); err != nil {
+		t.Fatalf("Cannot load add32: %v", err)
+	}
+
+	if fn, err = mod.Function("add32"); err != nil {
+		t.Fatalf("Cannot get add32(): %v", err)
+	}
+
+	args := []unsafe.Pointer{
+		unsafe.Pointer(&memA),
+		unsafe.Pointer(&memB),
+		unsafe.Pointer(&size),
+	}
+
+	if err = fn.LaunchAndSync(1, 1, 1, len(a), 1, 1, 1, Stream(0), args); err != nil {
+		t.Error("Launch and Sync Failed: %v", err)
+	}
+
+	if err = MemcpyDtoH(unsafe.Pointer(&a[0]), memA, size); err != nil {
+		t.Fatalf("Failed to copy memory to a: %v", err)
+	}
+
+	if err = MemcpyDtoH(unsafe.Pointer(&b[0]), memA, size); err != nil {
+		t.Fatalf("Failed to copy memory to b: %v", err)
+	}
+
+	for _, v := range a {
+		if v != float32(2) {
+			t.Error("Expected all values to be 2.")
+			break
+		}
+	}
+
+	Unload(mod)
+	DestroyContext(&ctx)
+}

test_test.go

@@ -0,0 +1,67 @@
+package cu
+
+const add32PTX = `//
+// Generated by NVIDIA NVVM Compiler
+//
+// Compiler Build ID: CL-21554848
+// Cuda compilation tools, release 8.0, V8.0.61
+// Based on LLVM 3.4svn
+//
+
+.version 5.0
+.target sm_30
+.address_size 64
+
+	// .globl	add32
+
+.visible .entry add32(
+	.param .u64 add32_param_0,
+	.param .u64 add32_param_1,
+	.param .u32 add32_param_2
+)
+{
+	.reg .pred 	%p<2>;
+	.reg .f32 	%f<4>;
+	.reg .b32 	%r<19>;
+	.reg .b64 	%rd<8>;
+
+
+	ld.param.u64 	%rd1, [add32_param_0];
+	ld.param.u64 	%rd2, [add32_param_1];
+	ld.param.u32 	%r2, [add32_param_2];
+	mov.u32 	%r3, %ctaid.x;
+	mov.u32 	%r4, %ctaid.z;
+	mov.u32 	%r5, %nctaid.y;
+	mov.u32 	%r6, %ctaid.y;
+	mad.lo.s32 	%r7, %r4, %r5, %r6;
+	mov.u32 	%r8, %nctaid.x;
+	mad.lo.s32 	%r9, %r7, %r8, %r3;
+	mov.u32 	%r10, %ntid.y;
+	mov.u32 	%r11, %ntid.x;
+	mul.lo.s32 	%r12, %r10, %r11;
+	mov.u32 	%r13, %ntid.z;
+	mov.u32 	%r14, %tid.y;
+	mov.u32 	%r15, %tid.z;
+	mad.lo.s32 	%r16, %r9, %r13, %r15;
+	mov.u32 	%r17, %tid.x;
+	mad.lo.s32 	%r18, %r14, %r11, %r17;
+	mad.lo.s32 	%r1, %r12, %r16, %r18;
+	setp.ge.s32	%p1, %r1, %r2;
+	@%p1 bra 	BB0_2;
+
+	cvta.to.global.u64 	%rd3, %rd1;
+	mul.wide.s32 	%rd4, %r1, 4;
+	add.s64 	%rd5, %rd3, %rd4;
+	cvta.to.global.u64 	%rd6, %rd2;
+	add.s64 	%rd7, %rd6, %rd4;
+	ld.global.f32 	%f1, [%rd7];
+	ld.global.f32 	%f2, [%rd5];
+	add.rn.f32 	%f3, %f2, %f1;
+	st.global.f32 	[%rd5], %f3;
+
+BB0_2:
+	ret;
+}
+
+
+`