Mercurial > hg > Game > Cerium

changeset 1923:e801016bd47c draft

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
fix
author kkb
date Wed, 29 Jan 2014 20:32:24 +0900
parents bbd209709ca1
children 14f9fc88872c
files TaskManager/Cuda/CudaScheduler.cc TaskManager/Cuda/CudaScheduler.h example/Cuda/main.cc example/Cuda/multiply.cu
diffstat 4 files changed, 15 insertions(+), 6 deletions(-) [+]
line wrap: on
line diff
--- a/TaskManager/Cuda/CudaScheduler.cc	Wed Jan 29 18:55:59 2014 +0900
+++ b/TaskManager/Cuda/CudaScheduler.cc	Wed Jan 29 20:32:24 2014 +0900
@@ -47,12 +47,14 @@
     m->allcate_size = 64;
     m->buf = (CUdeviceptr*)malloc(m->allcate_size*sizeof(CUdeviceptr*));
     m->event = (CUevent*)malloc(m->allcate_size*sizeof(CUevent*));
+    m->stream = (CUStream*)malloc(m->allcate_size*sizeof(CUStream*));
 }
 
 void
 CudaScheduler::destroyCudaBuffer(CudaBufferPtr m) {
     free(m->buf);
     free(m->event);
+    free(m->stream);
     m->size = 0;
     m->allcate_size = 0;
     m->buf = 0;
@@ -65,6 +67,7 @@
         m->allcate_size *= 2;
         m->buf = (CUdeviceptr*)realloc(m->buf, m->allcate_size*sizeof(CUdeviceptr*));
         m->event = (CUevent*)remalloc(m->allcate_size*sizeof(CUevent*));
+        m->stream = (CUStream*)remalloc(m->allcate_size*sizeof(CUStream*));
     }
 
     error = cuMemAlloc(&m->buf[i], size);
@@ -250,7 +253,7 @@
                 
                 if (tasklist->dim > 0) {
                     ret = cuLaunchKernel(kernel[cur],
-                                         tasklist->x*tasklist->y*tasklist->z, 0, 0,
+                                         tasklist->x*tasklist->y*tasklist->z, 1, 1,
                                          1, 1, 1,
                                          stream, kernelParams, NULL);
                 } else {
--- a/TaskManager/Cuda/CudaScheduler.h	Wed Jan 29 18:55:59 2014 +0900
+++ b/TaskManager/Cuda/CudaScheduler.h	Wed Jan 29 20:32:24 2014 +0900
@@ -19,6 +19,7 @@
         int size;
         CUdeviceptr* buf;
         CUevent* event;
+        CUStream* stream;
     } CudaBuffer;
     cudabuffer* CudaBufferPtr;
     CudaScheduler();
--- a/example/Cuda/main.cc	Wed Jan 29 18:55:59 2014 +0900
+++ b/example/Cuda/main.cc	Wed Jan 29 20:32:24 2014 +0900
@@ -26,6 +26,7 @@
     CUcontext context;
     CUmodule module;
     CUfunction function;
+    CUStream stream;
 
     cuInit(0);
     cuDeviceGet(&device, 0);
@@ -33,6 +34,9 @@
     cuModuleLoad(&module, "multiply.ptx");
     cuModuleGetFunction(&function, module, "multiply");
     
+    cuStramCreate(&steam,0);
+
+
     float* A = new float[LENGTH];
     float* B = new float[LENGTH];
     float* C = new float[LENGTH];
@@ -48,18 +52,19 @@
     cuMemAlloc(&devB, LENGTH*sizeof(float));
     cuMemAlloc(&devC, LENGTH*sizeof(float));
 
-    cuMemcpyHtoD(devA, A, LENGTH*sizeof(float));
-    cuMemcpyHtoD(devB, B, LENGTH*sizeof(float));
-    cuMemcpyHtoD(devC, C, LENGTH*sizeof(float));
+    cuMemcpyHtoDAsync(devA, A, LENGTH*sizeof(float), stream);
+    cuMemcpyHtoDAsync(devB, B, LENGTH*sizeof(float), stream);
+    cuMemcpyHtoDAsync(devC, C, LENGTH*sizeof(float), stream);
 
     void* args[] = {&devA, &devB, &devC};
     
     cuLaunchKernel(function,
                    LENGTH, 1, 1,
                    1, 1, 1,
-                   0, NULL, args, NULL);
+                   0, stream, args, NULL);
     
     cuMemcpyDtoH(C, devC, LENGTH*sizeof(float));
+    cuStreamWaitEvent(stream, ,0);
 
     //    print_result(C);
     check_data(A, B, C);
@@ -71,6 +76,7 @@
     cuMemFree(devB);
     cuMemFree(devC);
     cuModuleUnload(module);
+    cuStreamDestroy(stream);
     cuCtxDestroy(context);
 
     return 0;
--- a/example/Cuda/multiply.cu	Wed Jan 29 18:55:59 2014 +0900
+++ b/example/Cuda/multiply.cu	Wed Jan 29 20:32:24 2014 +0900
@@ -1,7 +1,6 @@
 extern "C" {
     __global__ void multiply(float* A, float* B, float* C) {
         int index = blockIdx.x * blockDim.x + threadIdx.x;
-        printf("%d\n",index);
         C[index] = A[index] * B[index];
     }
 }