Skip to content

HIP/CUDA Gen: Use correct block size bounds in non-tensor dynamic dispatch calculations #1993

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
zatkins-dev wants to merge 3 commits into
base: main
Choose a base branch
from
+84 −146
Open

HIP/CUDA Gen: Use correct block size bounds in non-tensor dynamic dispatch calculations #1993

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
9c3dcd8
cuda: Fix block size determination for non-tensor to use max block si…
zatkins-dev Jul 17, 2026
aef039d
hip(gen) - fix bug in field reuse which caused RTC compile error
zatkins-dev Jul 17, 2026
dd9e5e1
hip(gen) - simplify and optimize block size determination to allow hi…
zatkins-dev Jul 17, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions backends/cuda-gen/ceed-cuda-gen-operator.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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));

@zatkins-dev zatkins-dev Jul 17, 2026

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We could also leave the cap on 512 by doing this:

Suggested change
CeedInt elems_per_block = CeedIntMin(cuda_data->device_prop.maxThreadsDim[2], CeedIntMax(max_threads_per_block / data->thread_1d, 1));
CeedInt elems_per_block = CeedIntMin(cuda_data->device_prop.maxThreadsDim[2], CeedIntMax(CeedIntMin(max_threads_per_block, 512) / data->thread_1d, 1));

That would ensure the previous behavior for the standard non-tensor elements while preventing the elems_per_block from being too high for the weird cases.

@jeremylt jeremylt Jul 20, 2026

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

That's probably a good safeguard here


grid = num_elem / elems_per_block + (num_elem % elems_per_block > 0);
block[2] = elems_per_block;
Expand Down Expand Up @@ -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;
Expand Down
71 changes: 15 additions & 56 deletions backends/hip-gen/ceed-hip-gen-operator-build.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
//------------------------------------------------------------------------------
Expand Down Expand Up @@ -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();

Expand Down Expand Up @@ -1479,6 +1450,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;
Expand Down Expand Up @@ -1679,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));
Expand Down Expand Up @@ -2157,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(),
Expand Down Expand Up @@ -2750,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));
Expand Down
1 change: 0 additions & 1 deletion backends/hip-gen/ceed-hip-gen-operator-build.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,7 +9,6 @@
#include <ceed.h>
#include <ceed/backend.h>

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);
Expand Down
Loading
Loading