Mercurial > hg > GearsTemplate

changeset 451:dcc42f3e7e97

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
Auto choice blockDim
author Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
date Tue, 05 Dec 2017 06:33:40 +0900
parents d3d7a7d6a117
children 1432d924c472 95f58f2b2c0e
files src/parallel_execution/CUDAExecutor.cbc src/parallel_execution/CUDAWorker.cbc src/parallel_execution/context.h src/parallel_execution/cuda.c
diffstat 4 files changed, 21 insertions(+), 11 deletions(-) [+]
line wrap: on
line diff
--- a/src/parallel_execution/CUDAExecutor.cbc	Mon Dec 04 04:24:30 2017 +0900
+++ b/src/parallel_execution/CUDAExecutor.cbc	Tue Dec 05 06:33:40 2017 +0900
@@ -7,9 +7,10 @@
 #include "../helper_cuda.h"
 #include "pthread.h"
 
-Executor* createCUDAExecutor(struct Context* context) {
+Executor* createCUDAExecutor(struct Context* context, CUdevice device) {
     struct Executor* executor = new Executor();
     struct CUDAExecutor* cudaExecutor = new CUDAExecutor();
+    checkCudaErrors(cuDeviceGetAttribute(&cudaExecutor->maxThreadPerBlock, CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK, device));
     executor->executor = (union Data*)cudaExecutor;
     executor->read  = C_readCUDAExecutor;
     executor->exec  = C_execCUDAExecutor;
@@ -35,14 +36,21 @@
     goto meta(context, C_execCUDAExecutor);
 }
 
+int computeblockDim(int count, int maxThreadPerBlock) {
+    return count < maxThreadPerBlock ? count : maxThreadPerBlock;
+}
+
 __code execCUDAExecutor(struct CUDAExecutor* executor, struct Context* task, __code next(...)) {
     // Asynchronous launch kernel
     task->num_exec = 1;
     if (task->iterate) {
         struct MultiDimIterator* iterator = &task->iterator->iterator->MultiDimIterator;
+        int blockDimX = computeblockDim(iterator->x, executor->maxThreadPerBlock);
+        int blockDimY = computeblockDim(iterator->y, executor->maxThreadPerBlock);
+        int blockDimZ = computeblockDim(iterator->z, executor->maxThreadPerBlock);
         checkCudaErrors(cuLaunchKernel(task->function,
-                    iterator->x, iterator->y, iterator->z,
-                    1, 1, 1,
+                    iterator->x/blockDimX, iterator->y/blockDimY, iterator->z/blockDimZ,
+                    blockDimX, blockDimY, blockDimZ,
                     0, NULL, (void**)executor->kernelParams, NULL));
     } else {
         checkCudaErrors(cuLaunchKernel(task->function,
--- a/src/parallel_execution/CUDAWorker.cbc	Mon Dec 04 04:24:30 2017 +0900
+++ b/src/parallel_execution/CUDAWorker.cbc	Tue Dec 05 06:33:40 2017 +0900
@@ -18,10 +18,11 @@
 
 static void startCUDAWorker(Worker* worker) {
     struct CUDAWorker* cudaWorker = &worker->worker->CUDAWorker;
-    cudaInit(cudaWorker, 0);
+    int deviceNum = 0;
+    cudaInit(cudaWorker, deviceNum);
     cudaWorker->context  = NEW(struct Context);
     initContext(cudaWorker->context);
-    cudaWorker->executor = createCUDAExecutor(cudaWorker->context);
+    cudaWorker->executor = createCUDAExecutor(cudaWorker->context, cudaWorker->device);
     Gearef(cudaWorker->context, Worker)->worker = (union Data*)worker;
     goto meta(cudaWorker->context, worker->taskReceive);
 }
--- a/src/parallel_execution/context.h	Mon Dec 04 04:24:30 2017 +0900
+++ b/src/parallel_execution/context.h	Tue Dec 05 06:33:40 2017 +0900
@@ -198,7 +198,7 @@
         struct Queue* tasks;
         int runFlag;
         enum Code next;
-        int num_stream;
+        int numStream;
         struct Executor* executor;
         CUstream *stream;
     } CUDAWorker;
@@ -370,6 +370,7 @@
     struct CUDAExecutor {
         CUdeviceptr** kernelParams;
         struct CUDABuffer* buffer;
+        int maxThreadPerBlock;
     } CUDAExecutor;
     struct CUDABuffer {
         int inputLen;
--- a/src/parallel_execution/cuda.c	Mon Dec 04 04:24:30 2017 +0900
+++ b/src/parallel_execution/cuda.c	Tue Dec 05 06:33:40 2017 +0900
@@ -51,7 +51,7 @@
     struct Queue* tasks;
     int runFlag;
     int next;
-    int num_stream;
+    int numStream;
     CUstream *stream;
 } CUDAWorker;
 
@@ -67,14 +67,14 @@
 } Array;
 */
 
-void cudaInit(struct CUDAWorker *cudaWorker,int phase) {
+void cudaInit(struct CUDAWorker *cudaWorker,int phase, int deviceNum) {
     // initialize and load kernel
-    cudaWorker->num_stream = 1; // number of stream
-    //    cudaWorker->stream = NEWN(cudaWorker->num_stream, CUstream );
+    cudaWorker->numStream = 1; // number of stream
+    //    cudaWorker->stream = NEWN(cudaWorker->numStream, CUstream );
     if (phase==0)
         checkCudaErrors(cuInit(0));
     if (phase==0)
-        checkCudaErrors(cuDeviceGet(&cudaWorker->device, 0));
+        checkCudaErrors(cuDeviceGet(&cudaWorker->device, deviceNum));
     if (phase==0)
         checkCudaErrors(cuCtxCreate(&cudaWorker->cuCtx, CU_CTX_SCHED_SPIN, cudaWorker->device));
     //    if (cudaWorker->num_stream) {