Mercurial > hg > Game > Cerium
changeset 1923:e801016bd47c draft
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]; } }