HOLY CRAP t590 WORKS WITH CUDA GEN

CEED · Dec 12, 2024 · 6ed6b16 · 6ed6b16
1 parent 66c3c7c
commit 6ed6b16
Showing 1 changed file with 1 addition and 1 deletion.
diff --git a/backends/cuda-gen/ceed-cuda-gen-operator-build.cpp b/backends/cuda-gen/ceed-cuda-gen-operator-build.cpp
@@ -580,7 +580,7 @@ static int CeedOperatorBuildKernelQFunction_Cuda_gen(std::ostringstream &code, C
     code << "\n    // Note: Using batches of points\n";
     code << "    const CeedInt point_loop_bound = (blockDim.x * blockDim.y) * ceil(1.0 * max_num_points / (blockDim.x * blockDim.y));\n\n";
     code << "    #pragma unroll\n";
-    code << "    for (CeedInt i = 0; i < point_loop_bound; i++) {\n";
+    code << "    for (CeedInt i = threadIdx.x + threadIdx.y * blockDim.x; i < point_loop_bound; i += blockDim.x * blockDim.y) {\n";
     code << "      const CeedInt p = i % max_num_points;\n\n";
 
     code << "      // -- Coordinates\n";