From 58aaaa91469141799a5f258c5d0a02ce9c2b5c69 Mon Sep 17 00:00:00 2001 From: Zach Atkins Date: Mon, 8 Jun 2026 13:14:21 -0600 Subject: [PATCH] gpu(hip) - fix compilation error in gen at-points diagonal assembly --- backends/hip-gen/ceed-hip-gen-operator-build.cpp | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) diff --git a/backends/hip-gen/ceed-hip-gen-operator-build.cpp b/backends/hip-gen/ceed-hip-gen-operator-build.cpp index a0e4021e1b..a5eb82a765 100644 --- a/backends/hip-gen/ceed-hip-gen-operator-build.cpp +++ b/backends/hip-gen/ceed-hip-gen-operator-build.cpp @@ -2492,8 +2492,15 @@ extern "C" int CeedOperatorBuildKernelLinearAssembleQFunction_Hip_gen(CeedOperat // Loop over all elements code << "\n" << tab << "// Element loop\n"; code << tab << "__syncthreads();\n"; - code << tab << "for (CeedInt elem = blockIdx.x*blockDim.z + threadIdx.z; elem < num_elem; elem += gridDim.x*blockDim.z) {\n"; + code << tab << "#if CEED_HIP_USE_CHIPSTAR\n"; + code << tab << "// Pad out elements so all threads hit syncthreads()\n"; + code << tab << "const CeedInt elem_loop_bound = (gridDim.x*blockDim.z) * ceil(1.0*num_elem/(gridDim.x*blockDim.z));\n"; + code << tab << "#else\n"; + code << tab << "const CeedInt elem_loop_bound = num_elem;\n"; + code << tab << "#endif\n"; + code << tab << "for (CeedInt e = blockIdx.x*blockDim.z + threadIdx.z; e < elem_loop_bound; e += gridDim.x*blockDim.z) {\n"; tab.push(); + code << tab << "const CeedInt elem = e % num_elem;\n\n"; // -- Compute minimum buffer space needed CeedInt max_rstr_buffer_size = 1;