Functional first kernel tests with both APIs.

include/kernels.h

@@ -3,4 +3,4 @@
 
 __device__ void add_two_device(CUdeviceptr number);
 __global__ void add_three_global(CUdeviceptr number);
-__global__ void add_one_global(void* number);
+__global__ void add_one_global(CUdeviceptr number);

src/kernels.cu

@@ -18,7 +18,7 @@ void add_three_global(CUdeviceptr number)
 }
 
 __global__
-void add_one_global(void* number)
+void add_one_global(CUdeviceptr number)
 {
     (* (int*) number)++;
 }

src/main.cu

@@ -56,8 +56,8 @@ int main()
     uint32_t *buf_ptr = (uint32_t *)((char *)va + offset);
 
     printf("All set\n");
-    /* printf("va: %lu\ndptr: %llu\nGPUInfo.va: %lu\noffset: %d\nbuf_ptr: %d\n", */
-    /* 	   (uint64_t) va,dptr,GPUInfo.va,offset,*buf_ptr); */
+    printf("va: 0x%lx\ndptr: 0x%llx\nGPUInfo.va: 0x%lx\noffset: 0x%x\nbuf_ptr: 0x%lx\n",
+    	   (uint64_t) va,dptr,GPUInfo.va,offset,buf_ptr);
     
     /* At this point the GPU's mem is mapped to a CPU buffer to enable DMA */
 
@@ -68,9 +68,21 @@ int main()
     printf("set = %d\nget = %d\n",set,get);
     assert_cu( cuMemcpyHtoD(dptr,&set,sizeof(set)) );
     add_three_global<<< 1,1 >>>(dptr);
+    cudaDeviceSynchronize();
     assert_cu( cuMemcpyDtoH(&get,dptr,sizeof(get)) );
     printf("set = %d\nget = %d\n",set,get);
 
+    printf("Use the gdr api\n");
+    set = 4242;
+    get = 0;
+    printf("set = %d\nget = %d\nva = %d\n",set,get,*(int*) va);
+    assert_gdr( gdr_copy_to_bar(va,&set,sizeof(set)) );
+    add_one_global<<< 1,1 >>>(dptr);
+    cudaDeviceSynchronize();
+    assert_gdr( gdr_copy_from_bar(&get,va,sizeof(get)) );
+    printf("set = %d\nget = %d\nva = %d\n",set,get,*(int*) va);
+
+    
     /* Close everything */
     assert_gdr( gdr_unmap(g,GPUMemHandle,va,(size_t) GPUProp.sharedMemPerBlock) );
     assert_gdr( gdr_unpin_buffer(g,GPUMemHandle) );