From 9c3dcd83cb90d2f31f94abebf4dc474e564f226d Mon Sep 17 00:00:00 2001 From: Zach Atkins Date: Fri, 17 Jul 2026 11:17:41 -0600 Subject: [PATCH 1/3] cuda: Fix block size determination for non-tensor to use max block size (instead of constant 512) --- backends/cuda-gen/ceed-cuda-gen-operator.c | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/backends/cuda-gen/ceed-cuda-gen-operator.c b/backends/cuda-gen/ceed-cuda-gen-operator.c index f8228fbfbd..ab59cefdf7 100644 --- a/backends/cuda-gen/ceed-cuda-gen-operator.c +++ b/backends/cuda-gen/ceed-cuda-gen-operator.c @@ -218,7 +218,7 @@ static int CeedOperatorApplyAddCore_Cuda_gen(CeedOperator op, CUstream stream, c CeedCallBackend(BlockGridCalculate(num_elem, min_grid_size / cuda_data->device_prop.multiProcessorCount, max_threads_per_block, cuda_data->device_prop.maxThreadsDim[2], cuda_data->device_prop.warpSize, block, &grid)); } else { - CeedInt elems_per_block = CeedIntMin(cuda_data->device_prop.maxThreadsDim[2], CeedIntMax(512 / data->thread_1d, 1)); + CeedInt elems_per_block = CeedIntMin(cuda_data->device_prop.maxThreadsDim[2], CeedIntMax(max_threads_per_block / data->thread_1d, 1)); grid = num_elem / elems_per_block + (num_elem % elems_per_block > 0); block[2] = elems_per_block; @@ -476,7 +476,7 @@ static int CeedOperatorLinearAssembleQFunctionCore_Cuda_gen(CeedOperator op, boo CeedCallBackend(BlockGridCalculate(num_elem, min_grid_size / cuda_data->device_prop.multiProcessorCount, max_threads_per_block, cuda_data->device_prop.maxThreadsDim[2], cuda_data->device_prop.warpSize, block, &grid)); } else { - CeedInt elems_per_block = CeedIntMin(cuda_data->device_prop.maxThreadsDim[2], CeedIntMax(512 / data->thread_1d, 1)); + CeedInt elems_per_block = CeedIntMin(cuda_data->device_prop.maxThreadsDim[2], CeedIntMax(max_threads_per_block / data->thread_1d, 1)); grid = num_elem / elems_per_block + (num_elem % elems_per_block > 0); block[2] = elems_per_block; From aef039d3ae1242e69fe6aa7fe1a0d7a4a069ee4c Mon Sep 17 00:00:00 2001 From: Zach Atkins Date: Fri, 17 Jul 2026 11:24:51 -0600 Subject: [PATCH 2/3] hip(gen) - fix bug in field reuse which caused RTC compile error --- backends/hip-gen/ceed-hip-gen-operator-build.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/backends/hip-gen/ceed-hip-gen-operator-build.cpp b/backends/hip-gen/ceed-hip-gen-operator-build.cpp index a5eb82a765..ff6e49b848 100644 --- a/backends/hip-gen/ceed-hip-gen-operator-build.cpp +++ b/backends/hip-gen/ceed-hip-gen-operator-build.cpp @@ -1479,6 +1479,7 @@ extern "C" int CeedOperatorBuildKernel_Hip_gen(CeedOperator op, bool *is_good_bu CeedCallBackend(CeedQFunctionFieldGetEvalMode(qf_output_fields[i], &eval_mode_i)); CeedCallBackend(CeedOperatorFieldGetBasis(op_output_fields[i], &basis_i)); + CeedCallBackend(CeedBasisIsTensor(basis_i, &is_tensor)); for (CeedInt j = 0; (output_matrix_reuse[i].index == -1) && (j < num_input_fields); j++) { CeedEvalMode eval_mode_j; CeedBasis basis_j; From dd9e5e109dc4e43758332ac4e500b2c4f144628c Mon Sep 17 00:00:00 2001 From: Zach Atkins Date: Fri, 17 Jul 2026 11:25:58 -0600 Subject: [PATCH 3/3] hip(gen) - simplify and optimize block size determination to allow higher order non-tensor elements without fallback to ref --- .../hip-gen/ceed-hip-gen-operator-build.cpp | 70 ++------ .../hip-gen/ceed-hip-gen-operator-build.h | 1 - backends/hip-gen/ceed-hip-gen-operator.c | 154 ++++++++---------- 3 files changed, 81 insertions(+), 144 deletions(-) diff --git a/backends/hip-gen/ceed-hip-gen-operator-build.cpp b/backends/hip-gen/ceed-hip-gen-operator-build.cpp index ff6e49b848..d826879618 100644 --- a/backends/hip-gen/ceed-hip-gen-operator-build.cpp +++ b/backends/hip-gen/ceed-hip-gen-operator-build.cpp @@ -30,34 +30,6 @@ struct FieldReuse_Hip { CeedEvalMode eval_mode; }; -//------------------------------------------------------------------------------ -// Calculate the block size used for launching the operator kernel -//------------------------------------------------------------------------------ -extern "C" int BlockGridCalculate_Hip_gen(const CeedInt dim, const CeedInt num_elem, const CeedInt P_1d, const CeedInt Q_1d, CeedInt *block_sizes) { - const CeedInt thread_1d = CeedIntMax(Q_1d, P_1d); - if (dim == 1) { - CeedInt elems_per_block = 64 * thread_1d > 256 ? 256 / thread_1d : 64; - - elems_per_block = elems_per_block > 0 ? elems_per_block : 1; - block_sizes[0] = thread_1d; - block_sizes[1] = 1; - block_sizes[2] = elems_per_block; - } else if (dim == 2) { - const CeedInt elems_per_block = thread_1d < 4 ? 16 : 2; - - block_sizes[0] = thread_1d; - block_sizes[1] = thread_1d; - block_sizes[2] = elems_per_block; - } else if (dim == 3) { - const CeedInt elems_per_block = thread_1d < 6 ? 4 : (thread_1d < 8 ? 2 : 1); - - block_sizes[0] = thread_1d; - block_sizes[1] = thread_1d; - block_sizes[2] = elems_per_block; - } - return CEED_ERROR_SUCCESS; -} - //------------------------------------------------------------------------------ // Determine type of operator //------------------------------------------------------------------------------ @@ -1386,8 +1358,7 @@ extern "C" int CeedOperatorBuildKernel_Hip_gen(CeedOperator op, bool *is_good_bu code << tab << "// s_B_[in,out]_i: Interpolation matrix, shared memory\n"; code << tab << "// s_G_[in,out]_i: Gradient matrix, shared memory\n"; code << tab << "// -----------------------------------------------------------------------------\n"; - code << tab << "extern \"C\" __launch_bounds__(BLOCK_SIZE)\n"; - code << "__global__ void " << operator_name + code << tab << "extern \"C\" __global__ void " << operator_name << "(CeedInt num_elem, void* ctx, FieldsInt_Hip indices, Fields_Hip fields, Fields_Hip B, Fields_Hip G, CeedScalar* W, Points_Hip points) {\n"; tab.push(); @@ -1680,18 +1651,14 @@ extern "C" int CeedOperatorBuildKernel_Hip_gen(CeedOperator op, bool *is_good_bu code << tab << "}\n"; code << tab << "// -----------------------------------------------------------------------------\n\n"; - CeedInt block_sizes[3] = {0, 0, 0}; - CeedInt num_elem; - // Compile - CeedCallBackend(CeedOperatorGetNumElements(op, &num_elem)); - CeedCallBackend(BlockGridCalculate_Hip_gen(is_all_tensor ? max_dim : 1, num_elem, data->max_P_1d, is_all_tensor ? Q_1d : Q, block_sizes)); { - bool is_compile_good = false; + bool is_compile_good = false; + const CeedInt T_1d = CeedIntMax(is_all_tensor ? Q_1d : Q, data->max_P_1d); - data->thread_1d = block_sizes[0]; + data->thread_1d = T_1d; CeedCallBackend(CeedTryCompile_Hip(ceed, code.str().c_str(), (std::string("operator_") + qfunction_name).c_str(), &is_compile_good, &data->module, - 2, "OP_T_1D", block_sizes[0], "BLOCK_SIZE", block_sizes[0] * block_sizes[1] * block_sizes[2])); + 1, "OP_T_1D", T_1d)); if (is_compile_good) { *is_good_build = true; CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, operator_name.c_str(), &data->op)); @@ -2158,19 +2125,15 @@ static int CeedOperatorBuildKernelAssemblyAtPoints_Hip_gen(CeedOperator op, bool code << tab << "}\n"; code << tab << "// -----------------------------------------------------------------------------\n\n"; - CeedInt block_sizes[3] = {0, 0, 0}; - CeedInt num_elem; - // Compile - CeedCallBackend(CeedOperatorGetNumElements(op, &num_elem)); - CeedCallBackend(BlockGridCalculate_Hip_gen(max_dim, num_elem, data->max_P_1d, Q_1d, block_sizes)); { - bool is_compile_good = false; + bool is_compile_good = false; + const CeedInt T_1d = CeedIntMax(is_all_tensor ? Q_1d : Q, data->max_P_1d); - data->thread_1d = block_sizes[0]; + data->thread_1d = T_1d; CeedCallBackend(CeedTryCompile_Hip(ceed, code.str().c_str(), (std::string("operator_assembly_at_points") + qfunction_name).c_str(), - &is_compile_good, is_full ? &data->module_assemble_full : &data->module_assemble_diagonal, 2, "OP_T_1D", - block_sizes[0], "BLOCK_SIZE", block_sizes[0] * block_sizes[1] * block_sizes[2])); + &is_compile_good, is_full ? &data->module_assemble_full : &data->module_assemble_diagonal, 1, "OP_T_1D", + T_1d)); if (is_compile_good) { *is_good_build = true; CeedCallBackend(CeedGetKernel_Hip(ceed, is_full ? data->module_assemble_full : data->module_assemble_diagonal, operator_name.c_str(), @@ -2751,19 +2714,14 @@ extern "C" int CeedOperatorBuildKernelLinearAssembleQFunction_Hip_gen(CeedOperat code << tab << "}\n"; code << tab << "// -----------------------------------------------------------------------------\n\n"; - CeedInt block_sizes[3] = {0, 0, 0}; - CeedInt num_elem; - // Compile - CeedCallBackend(CeedOperatorGetNumElements(op, &num_elem)); - CeedCallBackend(BlockGridCalculate_Hip_gen(max_dim, num_elem, data->max_P_1d, Q_1d, block_sizes)); { - bool is_compile_good = false; + bool is_compile_good = false; + const CeedInt T_1d = CeedIntMax(is_all_tensor ? Q_1d : Q, data->max_P_1d); - data->thread_1d = block_sizes[0]; + data->thread_1d = T_1d; CeedCallBackend(CeedTryCompile_Hip(ceed, code.str().c_str(), (std::string("operator_assembly") + qfunction_name).c_str(), &is_compile_good, - &data->module_assemble_qfunction, 2, "OP_T_1D", block_sizes[0], "BLOCK_SIZE", - block_sizes[0] * block_sizes[1] * block_sizes[2])); + &data->module_assemble_qfunction, 1, "OP_T_1D", T_1d)); if (is_compile_good) { *is_good_build = true; CeedCallBackend(CeedGetKernel_Hip(ceed, data->module_assemble_qfunction, operator_name.c_str(), &data->assemble_qfunction)); diff --git a/backends/hip-gen/ceed-hip-gen-operator-build.h b/backends/hip-gen/ceed-hip-gen-operator-build.h index 019643e54e..c69396fc7c 100644 --- a/backends/hip-gen/ceed-hip-gen-operator-build.h +++ b/backends/hip-gen/ceed-hip-gen-operator-build.h @@ -9,7 +9,6 @@ #include #include -CEED_INTERN int BlockGridCalculate_Hip_gen(CeedInt dim, CeedInt num_elem, CeedInt P_1d, CeedInt Q_1d, CeedInt *block_sizes); CEED_INTERN int CeedOperatorBuildKernel_Hip_gen(CeedOperator op, bool *is_good_build); CEED_INTERN int CeedOperatorBuildKernelFullAssemblyAtPoints_Hip_gen(CeedOperator op, bool *is_good_build); CEED_INTERN int CeedOperatorBuildKernelDiagonalAssemblyAtPoints_Hip_gen(CeedOperator op, bool *is_good_build); diff --git a/backends/hip-gen/ceed-hip-gen-operator.c b/backends/hip-gen/ceed-hip-gen-operator.c index 87c8fb5a4d..89f084225f 100644 --- a/backends/hip-gen/ceed-hip-gen-operator.c +++ b/backends/hip-gen/ceed-hip-gen-operator.c @@ -16,6 +16,34 @@ #include "ceed-hip-gen-operator-build.h" #include "ceed-hip-gen.h" +//------------------------------------------------------------------------------ +// Calculate the block size used for launching the operator kernel +//------------------------------------------------------------------------------ +static int BlockGridCalculate_Hip_gen(const CeedInt dim, const CeedInt num_elem, const CeedInt P_1d, const CeedInt Q_1d, CeedInt *block_sizes) { + const CeedInt thread_1d = CeedIntMax(Q_1d, P_1d); + if (dim == 1) { + CeedInt elems_per_block = 64 * thread_1d > 256 ? 256 / thread_1d : 64; + + elems_per_block = elems_per_block > 0 ? elems_per_block : 1; + block_sizes[0] = thread_1d; + block_sizes[1] = 1; + block_sizes[2] = elems_per_block; + } else if (dim == 2) { + const CeedInt elems_per_block = thread_1d < 4 ? 16 : 2; + + block_sizes[0] = thread_1d; + block_sizes[1] = thread_1d; + block_sizes[2] = elems_per_block; + } else if (dim == 3) { + const CeedInt elems_per_block = thread_1d < 6 ? 4 : (thread_1d < 8 ? 2 : 1); + + block_sizes[0] = thread_1d; + block_sizes[1] = thread_1d; + block_sizes[2] = elems_per_block; + } + return CEED_ERROR_SUCCESS; +} + //------------------------------------------------------------------------------ // Destroy operator //------------------------------------------------------------------------------ @@ -60,12 +88,14 @@ static int CeedOperatorApplyAddCore_Hip_gen(CeedOperator op, hipStream_t stream, CeedQFunction qf; CeedOperatorField *op_input_fields, *op_output_fields; CeedOperator_Hip_gen *data; + Ceed_Hip *hip_data; // Creation of the operator CeedCallBackend(CeedOperatorBuildKernel_Hip_gen(op, is_run_good)); if (!(*is_run_good)) return CEED_ERROR_SUCCESS; CeedCallBackend(CeedOperatorGetCeed(op, &ceed)); + CeedCallBackend(CeedGetData(ceed, &hip_data)); CeedCallBackend(CeedOperatorGetData(op, &data)); CeedCallBackend(CeedOperatorGetQFunction(op, &qf)); CeedCallBackend(CeedQFunctionGetData(qf, &qf_data)); @@ -152,38 +182,27 @@ static int CeedOperatorApplyAddCore_Hip_gen(CeedOperator op, hipStream_t stream, CeedCallBackend(CeedQFunctionGetInnerContextData(qf, CEED_MEM_DEVICE, &qf_data->d_c)); // Apply operator - void *opargs[] = {(void *)&num_elem, &qf_data->d_c, &data->indices, &data->fields, &data->B, &data->G, &data->W, &data->points}; + void *opargs[] = {(void *)&num_elem, &qf_data->d_c, &data->indices, &data->fields, &data->B, &data->G, &data->W, &data->points}; + int min_grid_size, max_threads_per_block; + CeedInt grid; CeedCallBackend(CeedOperatorHasTensorBases(op, &is_tensor)); + CeedCallHip(ceed, hipModuleOccupancyMaxPotentialBlockSize(&min_grid_size, &max_threads_per_block, data->op, sizeof(CeedScalar), 0x10000)); CeedInt block_sizes[3] = {data->thread_1d, ((!is_tensor || data->dim == 1) ? 1 : data->thread_1d), -1}; if (is_tensor) { CeedCallBackend(BlockGridCalculate_Hip_gen(data->dim, num_elem, data->max_P_1d, data->Q_1d, block_sizes)); + grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num_elem) ? 1 : 0); } else { - CeedInt elems_per_block = 64 * data->thread_1d > 256 ? 256 / data->thread_1d : 64; + CeedInt elems_per_block = CeedIntMin(hip_data->device_prop.maxThreadsDim[2], CeedIntMax(max_threads_per_block / data->thread_1d, 1)); - elems_per_block = elems_per_block > 0 ? elems_per_block : 1; - block_sizes[2] = elems_per_block; - } - if (data->dim == 1 || !is_tensor) { - CeedInt grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num_elem) ? 1 : 0); - CeedInt sharedMem = block_sizes[2] * data->thread_1d * sizeof(CeedScalar); - - CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->op, stream, grid, block_sizes[0], block_sizes[1], block_sizes[2], sharedMem, - is_run_good, opargs)); - } else if (data->dim == 2) { - CeedInt grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num_elem) ? 1 : 0); - CeedInt sharedMem = block_sizes[2] * data->thread_1d * data->thread_1d * sizeof(CeedScalar); - - CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->op, stream, grid, block_sizes[0], block_sizes[1], block_sizes[2], sharedMem, - is_run_good, opargs)); - } else if (data->dim == 3) { - CeedInt grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num_elem) ? 1 : 0); - CeedInt sharedMem = block_sizes[2] * data->thread_1d * data->thread_1d * sizeof(CeedScalar); - - CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->op, stream, grid, block_sizes[0], block_sizes[1], block_sizes[2], sharedMem, - is_run_good, opargs)); + grid = num_elem / elems_per_block + (num_elem % elems_per_block > 0); + block_sizes[2] = elems_per_block; } + CeedInt shared_mem = block_sizes[0] * block_sizes[1] * block_sizes[2] * sizeof(CeedScalar); + + CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->op, stream, grid, block_sizes[0], block_sizes[1], block_sizes[2], shared_mem, is_run_good, + opargs)); // Restore input arrays for (CeedInt i = 0; i < num_input_fields; i++) { @@ -435,37 +454,28 @@ static int CeedOperatorLinearAssembleQFunctionCore_Hip_gen(CeedOperator op, bool // Assemble QFunction bool is_tensor = false; void *opargs[] = {(void *)&num_elem, &qf_data->d_c, &data->indices, &data->fields, &data->B, &data->G, &data->W, &data->points, &assembled_array}; + int min_grid_size, max_threads_per_block; + CeedInt grid; CeedCallBackend(CeedOperatorHasTensorBases(op, &is_tensor)); + CeedCallHip(ceed, hipModuleOccupancyMaxPotentialBlockSize(&min_grid_size, &max_threads_per_block, data->assemble_qfunction, sizeof(CeedScalar), + 0x10000)); CeedInt block_sizes[3] = {data->thread_1d, ((!is_tensor || data->dim == 1) ? 1 : data->thread_1d), -1}; if (is_tensor) { CeedCallBackend(BlockGridCalculate_Hip_gen(data->dim, num_elem, data->max_P_1d, data->Q_1d, block_sizes)); + grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num_elem) ? 1 : 0); } else { - CeedInt elems_per_block = 64 * data->thread_1d > 256 ? 256 / data->thread_1d : 64; + CeedInt elems_per_block = CeedIntMin(hip_data->device_prop.maxThreadsDim[2], CeedIntMax(max_threads_per_block / data->thread_1d, 1)); elems_per_block = elems_per_block > 0 ? elems_per_block : 1; + grid = num_elem / elems_per_block + (num_elem % elems_per_block > 0); block_sizes[2] = elems_per_block; } - if (data->dim == 1 || !is_tensor) { - CeedInt grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num_elem) ? 1 : 0); - CeedInt sharedMem = block_sizes[2] * data->thread_1d * sizeof(CeedScalar); - - CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->assemble_qfunction, NULL, grid, block_sizes[0], block_sizes[1], block_sizes[2], - sharedMem, &is_run_good, opargs)); - } else if (data->dim == 2) { - CeedInt grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num_elem) ? 1 : 0); - CeedInt sharedMem = block_sizes[2] * data->thread_1d * data->thread_1d * sizeof(CeedScalar); - - CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->assemble_qfunction, NULL, grid, block_sizes[0], block_sizes[1], block_sizes[2], - sharedMem, &is_run_good, opargs)); - } else if (data->dim == 3) { - CeedInt grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num_elem) ? 1 : 0); - CeedInt sharedMem = block_sizes[2] * data->thread_1d * data->thread_1d * sizeof(CeedScalar); - - CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->assemble_qfunction, NULL, grid, block_sizes[0], block_sizes[1], block_sizes[2], - sharedMem, &is_run_good, opargs)); - } + CeedInt shared_mem = block_sizes[0] * block_sizes[1] * block_sizes[2] * sizeof(CeedScalar); + + CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->assemble_qfunction, NULL, grid, block_sizes[0], block_sizes[1], block_sizes[2], + shared_mem, &is_run_good, opargs)); // Restore input arrays for (CeedInt i = 0; i < num_input_fields; i++) { @@ -625,30 +635,15 @@ static int CeedOperatorLinearAssembleAddDiagonalAtPoints_Hip_gen(CeedOperator op // Assemble diagonal void *opargs[] = {(void *)&num_elem, &qf_data->d_c, &data->indices, &data->fields, &data->B, &data->G, &data->W, &data->points, &assembled_array}; - CeedInt block_sizes[3] = {data->thread_1d, (data->dim == 1 ? 1 : data->thread_1d), -1}; CeedCallBackend(BlockGridCalculate_Hip_gen(data->dim, num_elem, data->max_P_1d, data->Q_1d, block_sizes)); - block_sizes[2] = 1; - if (data->dim == 1) { - CeedInt grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num_elem) ? 1 : 0); - CeedInt sharedMem = block_sizes[2] * data->thread_1d * sizeof(CeedScalar); - - CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->assemble_diagonal, NULL, grid, block_sizes[0], block_sizes[1], block_sizes[2], - sharedMem, &is_run_good, opargs)); - } else if (data->dim == 2) { - CeedInt grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num_elem) ? 1 : 0); - CeedInt sharedMem = block_sizes[2] * data->thread_1d * data->thread_1d * sizeof(CeedScalar); - - CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->assemble_diagonal, NULL, grid, block_sizes[0], block_sizes[1], block_sizes[2], - sharedMem, &is_run_good, opargs)); - } else if (data->dim == 3) { - CeedInt grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num_elem) ? 1 : 0); - CeedInt sharedMem = block_sizes[2] * data->thread_1d * data->thread_1d * sizeof(CeedScalar); - - CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->assemble_diagonal, NULL, grid, block_sizes[0], block_sizes[1], block_sizes[2], - sharedMem, &is_run_good, opargs)); - } + + CeedInt grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num_elem) ? 1 : 0); + CeedInt shared_mem = block_sizes[0] * block_sizes[1] * block_sizes[2] * sizeof(CeedScalar); + + CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->assemble_diagonal, NULL, grid, block_sizes[0], block_sizes[1], block_sizes[2], + shared_mem, &is_run_good, opargs)); CeedCallHip(ceed, hipDeviceSynchronize()); // Restore input arrays @@ -808,32 +803,17 @@ static int CeedOperatorAssembleSingleAtPoints_Hip_gen(CeedOperator op, CeedInt o CeedScalar *assembled_offset_array = &assembled_array[offset]; // Assemble diagonal - void *opargs[] = {(void *)&num_elem, &qf_data->d_c, &data->indices, &data->fields, &data->B, - &data->G, &data->W, &data->points, &assembled_offset_array}; - + void *opargs[] = {(void *)&num_elem, &qf_data->d_c, &data->indices, &data->fields, &data->B, + &data->G, &data->W, &data->points, &assembled_offset_array}; CeedInt block_sizes[3] = {data->thread_1d, (data->dim == 1 ? 1 : data->thread_1d), -1}; CeedCallBackend(BlockGridCalculate_Hip_gen(data->dim, num_elem, data->max_P_1d, data->Q_1d, block_sizes)); - block_sizes[2] = 1; - if (data->dim == 1) { - CeedInt grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num_elem) ? 1 : 0); - CeedInt sharedMem = block_sizes[2] * data->thread_1d * sizeof(CeedScalar); - - CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->assemble_full, NULL, grid, block_sizes[0], block_sizes[1], block_sizes[2], sharedMem, - &is_run_good, opargs)); - } else if (data->dim == 2) { - CeedInt grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num_elem) ? 1 : 0); - CeedInt sharedMem = block_sizes[2] * data->thread_1d * data->thread_1d * sizeof(CeedScalar); - - CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->assemble_full, NULL, grid, block_sizes[0], block_sizes[1], block_sizes[2], sharedMem, - &is_run_good, opargs)); - } else if (data->dim == 3) { - CeedInt grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num_elem) ? 1 : 0); - CeedInt sharedMem = block_sizes[2] * data->thread_1d * data->thread_1d * sizeof(CeedScalar); - - CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->assemble_full, NULL, grid, block_sizes[0], block_sizes[1], block_sizes[2], sharedMem, - &is_run_good, opargs)); - } + + CeedInt grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num_elem) ? 1 : 0); + CeedInt shared_mem = block_sizes[0] * block_sizes[1] * block_sizes[2] * sizeof(CeedScalar); + + CeedCallBackend(CeedTryRunKernelDimShared_Hip(ceed, data->assemble_full, NULL, grid, block_sizes[0], block_sizes[1], block_sizes[2], shared_mem, + &is_run_good, opargs)); CeedCallHip(ceed, hipDeviceSynchronize()); // Restore input arrays