--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/paper/src/cuLaunchKernel.cbc	Tue Feb 09 18:44:53 2021 +0900
@@ -0,0 +1,21 @@
+int computeblockDim(int count, int maxThreadPerBlock) {
+    return count < maxThreadPerBlock ? count : maxThreadPerBlock;
+}
+
+__code execCUDAExecutor(struct CUDAExecutor* executor, struct Context* task, __code next(...)) {
+    // check data parallelism task
+    if (task->iterate) {
+        struct MultiDimIterator* iterator = &task->iterator->iterator->MultiDimIterator;
+
+        // compute block thread size 
+        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/blockDimX, iterator->y/blockDimY, iterator->z/blockDimZ,
+                    blockDimX, blockDimY, blockDimZ,
+                    0, NULL, (void**)executor->kernelParams, NULL));
+    }
+    ...
+}