cuda - update shared max matrix sizes

CEED · Jan 3, 2025 · 5d36201 · 5d36201
1 parent 0ff6123
commit 5d36201
Show file tree

Hide file tree

Showing 2 changed files with 7 additions and 3 deletions.
diff --git a/backends/cuda-shared/ceed-cuda-shared-basis.c b/backends/cuda-shared/ceed-cuda-shared-basis.c
@@ -641,6 +641,9 @@ int CeedBasisCreateH1_Cuda_shared(CeedElemTopology topo, CeedInt dim, CeedInt nu
   CeedCallBackend(CeedBasisGetCeed(basis, &ceed));
   CeedCallBackend(CeedCalloc(1, &data));
 
+  // Check max dim
+  CeedCheck(dim <= 3, ceed, CEED_ERROR_BACKEND, "Backend does not implement nontensor bases with dim > 3");
+
   // Copy basis data to GPU
   CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_INTERP, &q_comp_interp));
   CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_GRAD, &q_comp_grad));
@@ -651,6 +654,7 @@ int CeedBasisCreateH1_Cuda_shared(CeedElemTopology topo, CeedInt dim, CeedInt nu
   if (interp) {
     const CeedInt interp_bytes = q_bytes * num_nodes * q_comp_interp;
 
+    CeedCheck(interp_bytes < 45 * 45 * 8, ceed, CEED_ERROR_BACKEND, "Backend does not implement nontensor bases with P * Q this large");
     CeedCallCuda(ceed, cudaMalloc((void **)&data->d_interp_1d, interp_bytes));
     CeedCallCuda(ceed, cudaMemcpy(data->d_interp_1d, interp, interp_bytes, cudaMemcpyHostToDevice));
   }

diff --git a/backends/cuda-shared/kernels/cuda-shared-basis.cu b/backends/cuda-shared/kernels/cuda-shared-basis.cu
@@ -8,9 +8,9 @@
 #include <ceed.h>
 #include <cuda.h>
 
-const int               sizeMax = 16;
-__constant__ CeedScalar c_B[sizeMax * sizeMax];
-__constant__ CeedScalar c_G[sizeMax * sizeMax];
+const int               MAX_SIZE = 45, MAX_DIM = 3;
+__constant__ CeedScalar c_B[MAX_SIZE * MAX_SIZE];
+__constant__ CeedScalar c_G[MAX_SIZE * MAX_SIZE * MAX_DIM];
 
 //------------------------------------------------------------------------------
 // Interp device initialization