Skip to content

Commit fd85cfe

Browse files
zatkins-devzatkins-dev
authored
Merge pull request #1993 from CEED/zach/gen-block-size-fix
HIP/CUDA Gen: Use correct block size bounds in non-tensor dynamic dispatch calculations
2 parents 0e106d7 + 3723f5f commit fd85cfe

4 files changed

Lines changed: 86 additions & 146 deletions

File tree

‎backends/cuda-gen/ceed-cuda-gen-operator.c‎

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -218,7 +218,8 @@ static int CeedOperatorApplyAddCore_Cuda_gen(CeedOperator op, CUstream stream, c
218218
CeedCallBackend(BlockGridCalculate(num_elem, min_grid_size / cuda_data->device_prop.multiProcessorCount, max_threads_per_block,
219219
cuda_data->device_prop.maxThreadsDim[2], cuda_data->device_prop.warpSize, block, &grid));
220220
} else {
221-
CeedInt elems_per_block = CeedIntMin(cuda_data->device_prop.maxThreadsDim[2], CeedIntMax(512 / data->thread_1d, 1));
221+
CeedInt elems_per_block =
222+
CeedIntMin(cuda_data->device_prop.maxThreadsDim[2], CeedIntMax(CeedIntMin(max_threads_per_block, 512) / data->thread_1d, 1));
222223

223224
grid = num_elem / elems_per_block + (num_elem % elems_per_block > 0);
224225
block[2] = elems_per_block;
@@ -476,7 +477,8 @@ static int CeedOperatorLinearAssembleQFunctionCore_Cuda_gen(CeedOperator op, boo
476477
CeedCallBackend(BlockGridCalculate(num_elem, min_grid_size / cuda_data->device_prop.multiProcessorCount, max_threads_per_block,
477478
cuda_data->device_prop.maxThreadsDim[2], cuda_data->device_prop.warpSize, block, &grid));
478479
} else {
479-
CeedInt elems_per_block = CeedIntMin(cuda_data->device_prop.maxThreadsDim[2], CeedIntMax(512 / data->thread_1d, 1));
480+
CeedInt elems_per_block =
481+
CeedIntMin(cuda_data->device_prop.maxThreadsDim[2], CeedIntMax(CeedIntMin(max_threads_per_block, 512) / data->thread_1d, 1));
480482

481483
grid = num_elem / elems_per_block + (num_elem % elems_per_block > 0);
482484
block[2] = elems_per_block;

‎backends/hip-gen/ceed-hip-gen-operator-build.cpp‎

Lines changed: 15 additions & 56 deletions
Original file line numberDiff line numberDiff line change
@@ -30,34 +30,6 @@ struct FieldReuse_Hip {
3030
CeedEvalMode eval_mode;
3131
};
3232

33-
//------------------------------------------------------------------------------
34-
// Calculate the block size used for launching the operator kernel
35-
//------------------------------------------------------------------------------
36-
extern "C" int BlockGridCalculate_Hip_gen(const CeedInt dim, const CeedInt num_elem, const CeedInt P_1d, const CeedInt Q_1d, CeedInt *block_sizes) {
37-
const CeedInt thread_1d = CeedIntMax(Q_1d, P_1d);
38-
if (dim == 1) {
39-
CeedInt elems_per_block = 64 * thread_1d > 256 ? 256 / thread_1d : 64;
40-
41-
elems_per_block = elems_per_block > 0 ? elems_per_block : 1;
42-
block_sizes[0] = thread_1d;
43-
block_sizes[1] = 1;
44-
block_sizes[2] = elems_per_block;
45-
} else if (dim == 2) {
46-
const CeedInt elems_per_block = thread_1d < 4 ? 16 : 2;
47-
48-
block_sizes[0] = thread_1d;
49-
block_sizes[1] = thread_1d;
50-
block_sizes[2] = elems_per_block;
51-
} else if (dim == 3) {
52-
const CeedInt elems_per_block = thread_1d < 6 ? 4 : (thread_1d < 8 ? 2 : 1);
53-
54-
block_sizes[0] = thread_1d;
55-
block_sizes[1] = thread_1d;
56-
block_sizes[2] = elems_per_block;
57-
}
58-
return CEED_ERROR_SUCCESS;
59-
}
60-
6133
//------------------------------------------------------------------------------
6234
// Determine type of operator
6335
//------------------------------------------------------------------------------
@@ -1386,8 +1358,7 @@ extern "C" int CeedOperatorBuildKernel_Hip_gen(CeedOperator op, bool *is_good_bu
13861358
code << tab << "// s_B_[in,out]_i: Interpolation matrix, shared memory\n";
13871359
code << tab << "// s_G_[in,out]_i: Gradient matrix, shared memory\n";
13881360
code << tab << "// -----------------------------------------------------------------------------\n";
1389-
code << tab << "extern \"C\" __launch_bounds__(BLOCK_SIZE)\n";
1390-
code << "__global__ void " << operator_name
1361+
code << tab << "extern \"C\" __global__ void " << operator_name
13911362
<< "(CeedInt num_elem, void* ctx, FieldsInt_Hip indices, Fields_Hip fields, Fields_Hip B, Fields_Hip G, CeedScalar* W, Points_Hip points) {\n";
13921363
tab.push();
13931364

@@ -1479,6 +1450,7 @@ extern "C" int CeedOperatorBuildKernel_Hip_gen(CeedOperator op, bool *is_good_bu
14791450

14801451
CeedCallBackend(CeedQFunctionFieldGetEvalMode(qf_output_fields[i], &eval_mode_i));
14811452
CeedCallBackend(CeedOperatorFieldGetBasis(op_output_fields[i], &basis_i));
1453+
CeedCallBackend(CeedBasisIsTensor(basis_i, &is_tensor));
14821454
for (CeedInt j = 0; (output_matrix_reuse[i].index == -1) && (j < num_input_fields); j++) {
14831455
CeedEvalMode eval_mode_j;
14841456
CeedBasis basis_j;
@@ -1679,18 +1651,14 @@ extern "C" int CeedOperatorBuildKernel_Hip_gen(CeedOperator op, bool *is_good_bu
16791651
code << tab << "}\n";
16801652
code << tab << "// -----------------------------------------------------------------------------\n\n";
16811653

1682-
CeedInt block_sizes[3] = {0, 0, 0};
1683-
CeedInt num_elem;
1684-
16851654
// Compile
1686-
CeedCallBackend(CeedOperatorGetNumElements(op, &num_elem));
1687-
CeedCallBackend(BlockGridCalculate_Hip_gen(is_all_tensor ? max_dim : 1, num_elem, data->max_P_1d, is_all_tensor ? Q_1d : Q, block_sizes));
16881655
{
1689-
bool is_compile_good = false;
1656+
bool is_compile_good = false;
1657+
const CeedInt T_1d = CeedIntMax(is_all_tensor ? Q_1d : Q, data->max_P_1d);
16901658

1691-
data->thread_1d = block_sizes[0];
1659+
data->thread_1d = T_1d;
16921660
CeedCallBackend(CeedTryCompile_Hip(ceed, code.str().c_str(), (std::string("operator_") + qfunction_name).c_str(), &is_compile_good, &data->module,
1693-
2, "OP_T_1D", block_sizes[0], "BLOCK_SIZE", block_sizes[0] * block_sizes[1] * block_sizes[2]));
1661+
1, "OP_T_1D", T_1d));
16941662
if (is_compile_good) {
16951663
*is_good_build = true;
16961664
CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, operator_name.c_str(), &data->op));
@@ -2157,19 +2125,15 @@ static int CeedOperatorBuildKernelAssemblyAtPoints_Hip_gen(CeedOperator op, bool
21572125
code << tab << "}\n";
21582126
code << tab << "// -----------------------------------------------------------------------------\n\n";
21592127

2160-
CeedInt block_sizes[3] = {0, 0, 0};
2161-
CeedInt num_elem;
2162-
21632128
// Compile
2164-
CeedCallBackend(CeedOperatorGetNumElements(op, &num_elem));
2165-
CeedCallBackend(BlockGridCalculate_Hip_gen(max_dim, num_elem, data->max_P_1d, Q_1d, block_sizes));
21662129
{
2167-
bool is_compile_good = false;
2130+
bool is_compile_good = false;
2131+
const CeedInt T_1d = CeedIntMax(is_all_tensor ? Q_1d : Q, data->max_P_1d);
21682132

2169-
data->thread_1d = block_sizes[0];
2133+
data->thread_1d = T_1d;
21702134
CeedCallBackend(CeedTryCompile_Hip(ceed, code.str().c_str(), (std::string("operator_assembly_at_points") + qfunction_name).c_str(),
2171-
&is_compile_good, is_full ? &data->module_assemble_full : &data->module_assemble_diagonal, 2, "OP_T_1D",
2172-
block_sizes[0], "BLOCK_SIZE", block_sizes[0] * block_sizes[1] * block_sizes[2]));
2135+
&is_compile_good, is_full ? &data->module_assemble_full : &data->module_assemble_diagonal, 1, "OP_T_1D",
2136+
T_1d));
21732137
if (is_compile_good) {
21742138
*is_good_build = true;
21752139
CeedCallBackend(CeedGetKernel_Hip(ceed, is_full ? data->module_assemble_full : data->module_assemble_diagonal, operator_name.c_str(),
@@ -2750,19 +2714,14 @@ extern "C" int CeedOperatorBuildKernelLinearAssembleQFunction_Hip_gen(CeedOperat
27502714
code << tab << "}\n";
27512715
code << tab << "// -----------------------------------------------------------------------------\n\n";
27522716

2753-
CeedInt block_sizes[3] = {0, 0, 0};
2754-
CeedInt num_elem;
2755-
27562717
// Compile
2757-
CeedCallBackend(CeedOperatorGetNumElements(op, &num_elem));
2758-
CeedCallBackend(BlockGridCalculate_Hip_gen(max_dim, num_elem, data->max_P_1d, Q_1d, block_sizes));
27592718
{
2760-
bool is_compile_good = false;
2719+
bool is_compile_good = false;
2720+
const CeedInt T_1d = CeedIntMax(is_all_tensor ? Q_1d : Q, data->max_P_1d);
27612721

2762-
data->thread_1d = block_sizes[0];
2722+
data->thread_1d = T_1d;
27632723
CeedCallBackend(CeedTryCompile_Hip(ceed, code.str().c_str(), (std::string("operator_assembly") + qfunction_name).c_str(), &is_compile_good,
2764-
&data->module_assemble_qfunction, 2, "OP_T_1D", block_sizes[0], "BLOCK_SIZE",
2765-
block_sizes[0] * block_sizes[1] * block_sizes[2]));
2724+
&data->module_assemble_qfunction, 1, "OP_T_1D", T_1d));
27662725
if (is_compile_good) {
27672726
*is_good_build = true;
27682727
CeedCallBackend(CeedGetKernel_Hip(ceed, data->module_assemble_qfunction, operator_name.c_str(), &data->assemble_qfunction));

‎backends/hip-gen/ceed-hip-gen-operator-build.h‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -9,7 +9,6 @@
99
#include <ceed.h>
1010
#include <ceed/backend.h>
1111

12-
CEED_INTERN int BlockGridCalculate_Hip_gen(CeedInt dim, CeedInt num_elem, CeedInt P_1d, CeedInt Q_1d, CeedInt *block_sizes);
1312
CEED_INTERN int CeedOperatorBuildKernel_Hip_gen(CeedOperator op, bool *is_good_build);
1413
CEED_INTERN int CeedOperatorBuildKernelFullAssemblyAtPoints_Hip_gen(CeedOperator op, bool *is_good_build);
1514
CEED_INTERN int CeedOperatorBuildKernelDiagonalAssemblyAtPoints_Hip_gen(CeedOperator op, bool *is_good_build);

0 commit comments

Comments
 (0)