bug fix for MultiDimDeviceArray objects passed as kernel arguments

NVIDIA · Oct 2, 2019 · c6f93e6 · c6f93e6
1 parent 02280d5
commit c6f93e6
Show file tree

Hide file tree

Showing 3 changed files with 47 additions and 2 deletions.
diff --git a/projects/com.nvidia.grcuda.test/src/com/nvidia/grcuda/test/MultiDimArrayTest.java b/projects/com.nvidia.grcuda.test/src/com/nvidia/grcuda/test/MultiDimArrayTest.java
@@ -264,4 +264,43 @@ public void test2DimArrayOutOfBoundsOnWriteAccess() {
             matrix.getArrayElement(0).setArrayElement(53, 42);
         }
     }
+
+    /** CUDA C++ source code of incrementing kernel. */
+    private static final String INC2D_KERNEL_SOURCE = "template <typename T>                         \n" +
+                    "__global__ void inc2d(T *matrix, int num_dim1, int num_dim2) {                  \n" +
+                    "  const int num_elements = num_dim1 * num_dim2;                                 \n" +
+                    "  for (auto idx = blockIdx.x * blockDim.x + threadIdx.x; idx < num_elements;    \n" +
+                    "       idx += gridDim.x * blockDim.x) {                                         \n" +
+                    "    matrix[idx] += (T{} + 1);                                                   \n" +
+                    "  }                                                                             \n" +
+                    "}\n";
+    /** NFI Signature of incrementing kernel. */
+    private static final String INC2D_KERNEL_SIGNATURE = "pointer, sint32, sint32";
+
+    @Test
+    public void test2DimArrayAsKernelArgument() {
+        try (Context context = Context.newBuilder().allowAllAccess(true).build()) {
+            final Value deviceArrayConstructor = context.eval("grcuda", "DeviceArray");
+            final int numDim1 = 19;
+            final int numDim2 = 53;
+            Value matrix = deviceArrayConstructor.execute("int", numDim1, numDim2);
+            assertEquals(numDim1, matrix.getArraySize());
+            assertEquals(numDim2, matrix.getArrayElement(0).getArraySize());
+            for (int i = 0; i < numDim1; i++) {
+                for (int j = 0; j < numDim2; j++) {
+                    matrix.getArrayElement(i).setArrayElement(j, i * numDim2 + j);
+                }
+            }
+            final Value buildKernel = context.eval("grcuda", "buildkernel");
+            final Value kernel = buildKernel.execute(INC2D_KERNEL_SOURCE, "inc2d<int>", INC2D_KERNEL_SIGNATURE);
+            final int blocks = 80;
+            final int threadsPerBlock = 256;
+            kernel.execute(blocks, threadsPerBlock).execute(matrix, numDim1, numDim2);
+            for (int i = 0; i < numDim1; i++) {
+                for (int j = 0; j < numDim2; j++) {
+                    assertEquals(i * numDim2 + j + 1, matrix.getArrayElement(i).getArrayElement(j).asInt());
+                }
+            }
+        }
+    }
 }
diff --git a/projects/com.nvidia.grcuda/src/com/nvidia/grcuda/MultiDimDeviceArray.java b/projects/com.nvidia.grcuda/src/com/nvidia/grcuda/MultiDimDeviceArray.java
@@ -167,10 +167,10 @@ final LittleEndianNativeArrayView getNativeView() {
 
     @Override
     public String toString() {
-        return "DeviceArray(elementType=" + elementType +
+        return "MultiDimDeviceArray(elementType=" + elementType +
                         ", dims=" + Arrays.toString(elementsPerDimension) +
                         ", Elements=" + totalElementCount +
-                        ", size=" + getSizeBytes() + " bytes, " +
+                        ", size=" + getSizeBytes() + " bytes" +
                         ", nativeView=" + nativeView + ')';
     }
 

diff --git a/projects/com.nvidia.grcuda/src/com/nvidia/grcuda/gpu/Kernel.java b/projects/com.nvidia.grcuda/src/com/nvidia/grcuda/gpu/Kernel.java
@@ -33,6 +33,7 @@
 import java.util.ArrayList;
 import com.nvidia.grcuda.DeviceArray;
 import com.nvidia.grcuda.DeviceArray.MemberSet;
+import com.nvidia.grcuda.MultiDimDeviceArray;
 import com.nvidia.grcuda.gpu.UnsafeHelper.MemoryObject;
 import com.oracle.truffle.api.CompilerDirectives;
 import com.oracle.truffle.api.dsl.Fallback;
@@ -129,6 +130,11 @@ KernelArguments createKernelArguments(Object[] args, InteropLibrary int32Access,
                             UnsafeHelper.PointerObject pointer = UnsafeHelper.createPointerObject();
                             pointer.setValueOfPointer(deviceArray.getPointer());
                             kernelArgs.setArgument(argIdx, pointer);
+                        } else if (args[argIdx] instanceof MultiDimDeviceArray) {
+                            MultiDimDeviceArray deviceArray = (MultiDimDeviceArray) args[argIdx];
+                            UnsafeHelper.PointerObject pointer = UnsafeHelper.createPointerObject();
+                            pointer.setValueOfPointer(deviceArray.getPointer());
+                            kernelArgs.setArgument(argIdx, pointer);
                         } else {
                             CompilerDirectives.transferToInterpreter();
                             throw UnsupportedTypeException.create(new Object[]{args[argIdx]}, "expected DeviceArray type");