Mon Mar 4 12:47:19 AM EST 2024

diff --git a/src/gpu/gpu.c b/src/gpu/gpu.c
index cf8a318f1951828b6a6c2b34b0a109b201a0f969..1e01b0dd018bb9d6b66f359747f4ec170b62ec48 100644 (file)
--- a/src/gpu/gpu.c
+++ b/src/gpu/gpu.c
@@ -209,8 +209,8 @@ void GPU_knk(float* ptrR, int rowsR, int colsR, float* ptrA, int rowsA, int cols
        err = clSetKernelArg(kernel, 8, sizeof(int), &colsB); gpuerr(clSetKernelArg);
 
        //Run the program
-       size_t work_size[] = {rowsR, colsR};
-       err = clEnqueueNDRangeKernel(GPU_command_queue, kernel, 2, NULL, work_size, NULL, 0, NULL, NULL);
+       size_t work_size[] = {rowsR, colsR, 4};
+       err = clEnqueueNDRangeKernel(GPU_command_queue, kernel, 3, NULL, work_size, NULL, 0, NULL, NULL);
        gpuerr(clEnqueueNDRangeKernel);
 
        //Wait for completion

diff --git a/src/gpu/kernel.cl b/src/gpu/kernel.cl
index da234642b708f7b781b9eaab5b3f3972cbe2f176..3b465d13061a3f31c09210818479f8f152cf17d8 100644 (file)
--- a/src/gpu/kernel.cl
+++ b/src/gpu/kernel.cl
@@ -43,6 +43,7 @@ __kernel void gpu_knk
 {
     int rowR = get_global_id(0);
     int colR = get_global_id(1);
+    int oper = get_global_id(2);
 
     int rowA = rowR / rowsB;
     int colA = colR / colsB;
@@ -61,9 +62,12 @@ __kernel void gpu_knk
     r1 = first + last;
     i1 = outer + inner;
 
-    ptrR[(colR * 2) + (rowR * 2) * (colsR * 2)] = r1;
-    ptrR[((colR * 2) + 1) + (rowR * 2) * (colsR * 2)] = -i1;
-    ptrR[(colR * 2) + ((rowR * 2) + 1) * (colsR * 2)] = i1;
-    ptrR[((colR * 2) + 1) + ((rowR * 2) + 1) * (colsR * 2)] = r1;
+    switch (oper)
+    {
+        case 0: ptrR[(colR * 2) + (rowR * 2) * (colsR * 2)] = r1; break;
+        case 1: ptrR[((colR * 2) + 1) + (rowR * 2) * (colsR * 2)] = -i1; break;
+        case 2: ptrR[(colR * 2) + ((rowR * 2) + 1) * (colsR * 2)] = i1; break;
+        case 3: ptrR[((colR * 2) + 1) + ((rowR * 2) + 1) * (colsR * 2)] = r1; break;
+    }
 }