Thu Mar 7 06:48:19 PM EST 2024

diff --git a/src/.kernel.tmp.1 b/src/.kernel.tmp.1
index 04c0834df7c3007bc6287331ef314cffb3fe37a4..1da95816c44ededf150d07c2cba4676f78d3c39a 100644 (file)
--- a/src/.kernel.tmp.1
+++ b/src/.kernel.tmp.1
@@ -135,12 +135,12 @@ __kernel void kernel_knk_2x2
                     ptrR[(rowR + 0) * (colsR * 2) + ((colR + 1) * 2) + 1] = outer + inner;
                 break;
                 case 2:
-                    ptrR[(rowR + 0) * (colsR * 2) + ((colR + 0) * 2)] = first + lasts;
-                    ptrR[(rowR + 0) * (colsR * 2) + ((colR + 0) * 2) + 1] = outer + inner;
+                    ptrR[(rowR + 1) * (colsR * 2) + ((colR + 0) * 2)] = first + lasts;
+                    ptrR[(rowR + 1) * (colsR * 2) + ((colR + 0) * 2) + 1] = outer + inner;
                 break;
                 case 3:
-                    ptrR[(rowR + 0) * (colsR * 2) + ((colR + 1) * 2)] = first + lasts;
-                    ptrR[(rowR + 0) * (colsR * 2) + ((colR + 1) * 2) + 1] = outer + inner;
+                    ptrR[(rowR + 1) * (colsR * 2) + ((colR + 1) * 2)] = first + lasts;
+                    ptrR[(rowR + 1) * (colsR * 2) + ((colR + 1) * 2) + 1] = outer + inner;
                 break;
             }
 

diff --git a/src/.kernel.tmp.2 b/src/.kernel.tmp.2
index 32bfc81c314627764b9396871127e92642f20c4b..1e791cc72a16245bb1931b4a70dfb487507ea110 100644 (file)
Binary files a/src/.kernel.tmp.2 and b/src/.kernel.tmp.2 differ

diff --git a/src/complex.c b/src/complex.c
index a9ca17c8dc6406e31bee517b8c0157916dae573b..d6cbb4d609803145eeb5f862c141a23ca9b0aee1 100644 (file)
--- a/src/complex.c
+++ b/src/complex.c
@@ -227,7 +227,22 @@ void* cpx_mtx_knk_threads_2x2_run(void *context)
        int colsR = (ctx->colsA) * (ctx->colsB);
     for (int i = 0; i < (ctx->delimeterCount); i++)
     {
-               kernel_knk_2x2(ctx->ptrR, ctx->ptrA, ctx->rowsA, ctx->colsA, ctx->ptrB[0], ctx->ptrB[1], ctx->ptrB[2], ctx->ptrB[3], ctx->ptrB[4], ctx->ptrB[5], ctx->ptrB[6], ctx->ptrB[7], i + (ctx->delimeterStart));
+               kernel_knk_2x2
+               (
+                       ctx->ptrR,
+                       ctx->ptrA,
+                       ctx->rowsA,
+                       ctx->colsA,
+                       ctx->ptrB[0],
+                       ctx->ptrB[1],
+                       ctx->ptrB[2],
+                       ctx->ptrB[3],
+                       ctx->ptrB[4],
+                       ctx->ptrB[5],
+                       ctx->ptrB[6],
+                       ctx->ptrB[7],
+                       i + (ctx->delimeterStart)
+               );
     }
 }
 
@@ -681,7 +696,8 @@ void cpx_mtx_knk_metal_2x2(float* ptrR, float* ptrA, float* ptrB, int rowsA, int
        err = clSetKernelArg(kernel,11, sizeof(float), &gate7); gpuerr(clSetKernelArg);
 
        //Run the program
-       err = clEnqueueNDRangeKernel(cpx_mtx_command_queue, kernel, 1, NULL, (size_t[]){rowsR / 2}, NULL, 0, NULL, NULL);
+       size_t q = 16;
+       err = clEnqueueNDRangeKernel(cpx_mtx_command_queue, kernel, 1, NULL, (size_t[]){rowsR / 2}, &q, 0, NULL, NULL);
        gpuerr(clEnqueueNDRangeKernel);
 
        //Wait for completion

diff --git a/src/kernel_cpu.cl b/src/kernel_cpu.cl
index 7376958621439305809d29e25c5c0e65f0a386a7..f8a12f23f1378239ee431e2b4277cf1db16a9c45 100644 (file)
--- a/src/kernel_cpu.cl
+++ b/src/kernel_cpu.cl
@@ -139,12 +139,12 @@ void kernel_knk_2x2
                     ptrR[(rowR + 0) * (colsR * 2) + ((colR + 1) * 2) + 1] = outer + inner;
                 break;
                 case 2:
-                    ptrR[(rowR + 0) * (colsR * 2) + ((colR + 0) * 2)] = first + lasts;
-                    ptrR[(rowR + 0) * (colsR * 2) + ((colR + 0) * 2) + 1] = outer + inner;
+                    ptrR[(rowR + 1) * (colsR * 2) + ((colR + 0) * 2)] = first + lasts;
+                    ptrR[(rowR + 1) * (colsR * 2) + ((colR + 0) * 2) + 1] = outer + inner;
                 break;
                 case 3:
-                    ptrR[(rowR + 0) * (colsR * 2) + ((colR + 1) * 2)] = first + lasts;
-                    ptrR[(rowR + 0) * (colsR * 2) + ((colR + 1) * 2) + 1] = outer + inner;
+                    ptrR[(rowR + 1) * (colsR * 2) + ((colR + 1) * 2)] = first + lasts;
+                    ptrR[(rowR + 1) * (colsR * 2) + ((colR + 1) * 2) + 1] = outer + inner;
                 break;
             }
 

diff --git a/src/kernel_gpu.cl b/src/kernel_gpu.cl
index 6345e20c53d4a934cd6ae887be8dfb96cf92c094..e5e8787456c614b6b2a193478cbc256cafb82112 100644 (file)
--- a/src/kernel_gpu.cl
+++ b/src/kernel_gpu.cl
@@ -332,14 +332,14 @@ unsigned char kernel_gpu[] = {
   0x20, 0x63, 0x61, 0x73, 0x65, 0x20, 0x32, 0x3a, 0x0a, 0x20, 0x20, 0x20,
   0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
   0x20, 0x20, 0x20, 0x20, 0x20, 0x70, 0x74, 0x72, 0x52, 0x5b, 0x28, 0x72,
-  0x6f, 0x77, 0x52, 0x20, 0x2b, 0x20, 0x30, 0x29, 0x20, 0x2a, 0x20, 0x28,
+  0x6f, 0x77, 0x52, 0x20, 0x2b, 0x20, 0x31, 0x29, 0x20, 0x2a, 0x20, 0x28,
   0x63, 0x6f, 0x6c, 0x73, 0x52, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x20, 0x2b,
   0x20, 0x28, 0x28, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2b, 0x20, 0x30, 0x29,
   0x20, 0x2a, 0x20, 0x32, 0x29, 0x5d, 0x20, 0x3d, 0x20, 0x66, 0x69, 0x72,
   0x73, 0x74, 0x20, 0x2b, 0x20, 0x6c, 0x61, 0x73, 0x74, 0x73, 0x3b, 0x0a,
   0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
   0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x70, 0x74, 0x72, 0x52,
-  0x5b, 0x28, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2b, 0x20, 0x30, 0x29, 0x20,
+  0x5b, 0x28, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2b, 0x20, 0x31, 0x29, 0x20,
   0x2a, 0x20, 0x28, 0x63, 0x6f, 0x6c, 0x73, 0x52, 0x20, 0x2a, 0x20, 0x32,
   0x29, 0x20, 0x2b, 0x20, 0x28, 0x28, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2b,
   0x20, 0x30, 0x29, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x31,
@@ -351,14 +351,14 @@ unsigned char kernel_gpu[] = {
   0x73, 0x65, 0x20, 0x33, 0x3a, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
   0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
   0x20, 0x20, 0x70, 0x74, 0x72, 0x52, 0x5b, 0x28, 0x72, 0x6f, 0x77, 0x52,
-  0x20, 0x2b, 0x20, 0x30, 0x29, 0x20, 0x2a, 0x20, 0x28, 0x63, 0x6f, 0x6c,
+  0x20, 0x2b, 0x20, 0x31, 0x29, 0x20, 0x2a, 0x20, 0x28, 0x63, 0x6f, 0x6c,
   0x73, 0x52, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x28,
   0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2b, 0x20, 0x31, 0x29, 0x20, 0x2a, 0x20,
   0x32, 0x29, 0x5d, 0x20, 0x3d, 0x20, 0x66, 0x69, 0x72, 0x73, 0x74, 0x20,
   0x2b, 0x20, 0x6c, 0x61, 0x73, 0x74, 0x73, 0x3b, 0x0a, 0x20, 0x20, 0x20,
   0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
   0x20, 0x20, 0x20, 0x20, 0x20, 0x70, 0x74, 0x72, 0x52, 0x5b, 0x28, 0x72,
-  0x6f, 0x77, 0x52, 0x20, 0x2b, 0x20, 0x30, 0x29, 0x20, 0x2a, 0x20, 0x28,
+  0x6f, 0x77, 0x52, 0x20, 0x2b, 0x20, 0x31, 0x29, 0x20, 0x2a, 0x20, 0x28,
   0x63, 0x6f, 0x6c, 0x73, 0x52, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x20, 0x2b,
   0x20, 0x28, 0x28, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2b, 0x20, 0x31, 0x29,
   0x20, 0x2a, 0x20, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x31, 0x5d, 0x20, 0x3d,