Skip to content

Commit 629a26b

Browse files
zatkins-devzatkins-dev
committed
hip(kernels) - use nullptr instead of NULL for HIP/c++ source
1 parent 30b08b7 commit 629a26b

3 files changed

Lines changed: 16 additions & 16 deletions

File tree

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

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -302,7 +302,7 @@ static int CeedOperatorBuildKernelFieldData_Hip_gen(std::ostringstream &code, Ce
302302

303303
CeedCallBackend(CeedBasisIsCollocated(basis, &is_collocated));
304304
if ((is_active && skip_active_load) || (is_collocated && !is_at_points)) {
305-
code << tab << "CeedScalar *s_B" << var_suffix << " = NULL;\n";
305+
code << tab << "CeedScalar *s_B" << var_suffix << " = nullptr;\n";
306306
} else {
307307
code << tab << "__shared__ CeedScalar s_B" << var_suffix << "[" << P_name << "*" << Q_name << "];\n";
308308
code << tab << "LoadMatrix<" << P_name << ", " << Q_name << ">(data, B." << option_name << "[" << i << "], s_B" << var_suffix << ");\n";
@@ -341,7 +341,7 @@ static int CeedOperatorBuildKernelFieldData_Hip_gen(std::ostringstream &code, Ce
341341

342342
CeedCallBackend(CeedBasisIsCollocated(basis, &is_collocated));
343343
if ((is_active && skip_active_load) || (is_collocated && !is_at_points)) {
344-
code << tab << "CeedScalar *s_B" << var_suffix << " = NULL;\n";
344+
code << tab << "CeedScalar *s_B" << var_suffix << " = nullptr;\n";
345345
} else {
346346
code << tab << "__shared__ CeedScalar s_B" << var_suffix << "[" << P_name << "*" << Q_name << "];\n";
347347
code << tab << "LoadMatrix<" << P_name << ", " << Q_name << ">(data, B." << option_name << "[" << i << "], s_B" << var_suffix << ");\n";
@@ -357,7 +357,7 @@ static int CeedOperatorBuildKernelFieldData_Hip_gen(std::ostringstream &code, Ce
357357

358358
code << tab << "CeedScalar *s_G" << var_suffix << " = " << reuse_var << ";\n";
359359
} else if (is_active && skip_active_load) {
360-
code << tab << "CeedScalar *s_G" << var_suffix << " = NULL;\n";
360+
code << tab << "CeedScalar *s_G" << var_suffix << " = nullptr;\n";
361361
} else {
362362
code << tab << "__shared__ CeedScalar s_G" << var_suffix << "[" << Q_name << "*" << Q_name << "];\n";
363363
code << tab << "LoadMatrix<" << Q_name << ", " << Q_name << ">(data, G." << option_name << "[" << i << "], s_G" << var_suffix << ");\n";
@@ -373,7 +373,7 @@ static int CeedOperatorBuildKernelFieldData_Hip_gen(std::ostringstream &code, Ce
373373

374374
code << tab << "CeedScalar *s_G" << var_suffix << " = " << reuse_var << ";\n";
375375
} else if (is_active && skip_active_load) {
376-
code << tab << "CeedScalar *s_G" << var_suffix << " = NULL;\n";
376+
code << tab << "CeedScalar *s_G" << var_suffix << " = nullptr;\n";
377377
} else {
378378
code << tab << "__shared__ CeedScalar s_G" << var_suffix << "[" << Q_name << "*" << Q_name << "];\n";
379379
code << tab << "LoadMatrix<" << Q_name << ", " << Q_name << ">(data, G." << option_name << "[" << i << "], s_G" << var_suffix << ");\n";
@@ -384,7 +384,7 @@ static int CeedOperatorBuildKernelFieldData_Hip_gen(std::ostringstream &code, Ce
384384

385385
code << tab << "CeedScalar *s_G" << var_suffix << " = " << reuse_var << ";\n";
386386
} else if (is_active && skip_active_load) {
387-
code << tab << "CeedScalar *s_G" << var_suffix << " = NULL;\n";
387+
code << tab << "CeedScalar *s_G" << var_suffix << " = nullptr;\n";
388388
} else {
389389
code << tab << "__shared__ CeedScalar s_G" << var_suffix << "[" << P_name << "*" << Q_name << (is_tensor ? "" : "*dim")
390390
<< (is_tensor ? "" : var_suffix) << "];\n";

‎include/ceed/jit-source/hip/hip-ref-operator-assemble-diagonal.h‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -63,15 +63,15 @@ extern "C" __launch_bounds__(BLOCK_SIZE) __global__
6363
for (IndexType e_out = 0; e_out < NUM_EVAL_MODES_OUT; e_out++) {
6464
IndexType d_in = 0;
6565
CeedEvalMode eval_modes_in_prev = CEED_EVAL_NONE;
66-
const CeedScalar *b_t = NULL;
66+
const CeedScalar *b_t = nullptr;
6767

6868
GetBasisPointer(&b_t, eval_modes_out[e_out], identity, interp_out, grad_out, div_out, curl_out);
6969
if (e_out == 0 || eval_modes_out[e_out] != eval_modes_out_prev) d_out = 0;
7070
else b_t = &b_t[(++d_out) * NUM_QPTS * NUM_NODES];
7171
eval_modes_out_prev = eval_modes_out[e_out];
7272

7373
for (IndexType e_in = 0; e_in < NUM_EVAL_MODES_IN; e_in++) {
74-
const CeedScalar *b = NULL;
74+
const CeedScalar *b = nullptr;
7575

7676
GetBasisPointer(&b, eval_modes_in[e_in], identity, interp_in, grad_in, div_in, curl_in);
7777
if (e_in == 0 || eval_modes_in[e_in] != eval_modes_in_prev) d_in = 0;

‎include/ceed/jit-source/hip/hip-shared-basis-tensor.h‎

Lines changed: 9 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -407,21 +407,21 @@ extern "C" __launch_bounds__(BASIS_GRAD_BLOCK_SIZE) __global__
407407

408408
if (BASIS_DIM == 1) {
409409
ReadElementStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D, d_U, r_U);
410-
Grad1d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, NULL, s_G, r_V);
410+
Grad1d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, nullptr, s_G, r_V);
411411
if (e < num_elem) {
412412
WriteElementStrided1d<BASIS_NUM_COMP, BASIS_Q_1D>(data, elem, 1, BASIS_Q_1D * num_elem, BASIS_Q_1D, r_V, d_V);
413413
}
414414
} else if (BASIS_DIM == 2) {
415415
ReadElementStrided2d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * BASIS_P_1D * num_elem, BASIS_P_1D * BASIS_P_1D, d_U, r_U);
416-
GradTensorCollocatedNodes2d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, NULL, s_G, r_V);
416+
GradTensorCollocatedNodes2d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, nullptr, s_G, r_V);
417417
if (e < num_elem) {
418418
WriteElementStrided2d<BASIS_NUM_COMP * BASIS_DIM, BASIS_Q_1D>(data, elem, 1, BASIS_Q_1D * BASIS_Q_1D * num_elem, BASIS_Q_1D * BASIS_Q_1D, r_V,
419419
d_V);
420420
}
421421
} else if (BASIS_DIM == 3) {
422422
ReadElementStrided3d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * BASIS_P_1D * BASIS_P_1D * num_elem,
423423
BASIS_P_1D * BASIS_P_1D * BASIS_P_1D, d_U, r_U);
424-
GradTensorCollocatedNodes3d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, NULL, s_G, r_V);
424+
GradTensorCollocatedNodes3d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, nullptr, s_G, r_V);
425425
if (e < num_elem) {
426426
WriteElementStrided3d<BASIS_NUM_COMP * BASIS_DIM, BASIS_Q_1D>(data, elem, 1, BASIS_Q_1D * BASIS_Q_1D * BASIS_Q_1D * num_elem,
427427
BASIS_Q_1D * BASIS_Q_1D * BASIS_Q_1D, r_V, d_V);
@@ -522,21 +522,21 @@ extern "C" __launch_bounds__(BASIS_GRAD_BLOCK_SIZE) __global__
522522

523523
if (BASIS_DIM == 1) {
524524
ReadElementStrided1d<BASIS_NUM_COMP, BASIS_Q_1D>(data, elem, 1, BASIS_Q_1D * num_elem, BASIS_Q_1D, d_U, r_U);
525-
GradTranspose1d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, NULL, s_G, r_V);
525+
GradTranspose1d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, nullptr, s_G, r_V);
526526
if (e < num_elem) {
527527
WriteElementStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D, r_V, d_V);
528528
}
529529
} else if (BASIS_DIM == 2) {
530530
ReadElementStrided2d<BASIS_NUM_COMP * BASIS_DIM, BASIS_Q_1D>(data, elem, 1, BASIS_Q_1D * BASIS_Q_1D * num_elem, BASIS_Q_1D * BASIS_Q_1D, d_U,
531531
r_U);
532-
GradTransposeTensorCollocatedNodes2d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, NULL, s_G, r_V);
532+
GradTransposeTensorCollocatedNodes2d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, nullptr, s_G, r_V);
533533
if (e < num_elem) {
534534
WriteElementStrided2d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * BASIS_P_1D * num_elem, BASIS_P_1D * BASIS_P_1D, r_V, d_V);
535535
}
536536
} else if (BASIS_DIM == 3) {
537537
ReadElementStrided3d<BASIS_NUM_COMP * BASIS_DIM, BASIS_Q_1D>(data, elem, 1, BASIS_Q_1D * BASIS_Q_1D * BASIS_Q_1D * num_elem,
538538
BASIS_Q_1D * BASIS_Q_1D * BASIS_Q_1D, d_U, r_U);
539-
GradTransposeTensorCollocatedNodes3d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, NULL, s_G, r_V);
539+
GradTransposeTensorCollocatedNodes3d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, nullptr, s_G, r_V);
540540
if (e < num_elem) {
541541
WriteElementStrided3d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * BASIS_P_1D * BASIS_P_1D * num_elem,
542542
BASIS_P_1D * BASIS_P_1D * BASIS_P_1D, r_V, d_V);
@@ -637,21 +637,21 @@ extern "C" __launch_bounds__(BASIS_GRAD_BLOCK_SIZE) __global__
637637

638638
if (BASIS_DIM == 1) {
639639
ReadElementStrided1d<BASIS_NUM_COMP, BASIS_Q_1D>(data, elem, 1, BASIS_Q_1D * num_elem, BASIS_Q_1D, d_U, r_U);
640-
GradTranspose1d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, NULL, s_G, r_V);
640+
GradTranspose1d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, nullptr, s_G, r_V);
641641
if (e < num_elem) {
642642
SumElementStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D, r_V, d_V);
643643
}
644644
} else if (BASIS_DIM == 2) {
645645
ReadElementStrided2d<BASIS_NUM_COMP * BASIS_DIM, BASIS_Q_1D>(data, elem, 1, BASIS_Q_1D * BASIS_Q_1D * num_elem, BASIS_Q_1D * BASIS_Q_1D, d_U,
646646
r_U);
647-
GradTransposeTensorCollocatedNodes2d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, NULL, s_G, r_V);
647+
GradTransposeTensorCollocatedNodes2d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, nullptr, s_G, r_V);
648648
if (e < num_elem) {
649649
SumElementStrided2d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * BASIS_P_1D * num_elem, BASIS_P_1D * BASIS_P_1D, r_V, d_V);
650650
}
651651
} else if (BASIS_DIM == 3) {
652652
ReadElementStrided3d<BASIS_NUM_COMP * BASIS_DIM, BASIS_Q_1D>(data, elem, 1, BASIS_Q_1D * BASIS_Q_1D * BASIS_Q_1D * num_elem,
653653
BASIS_Q_1D * BASIS_Q_1D * BASIS_Q_1D, d_U, r_U);
654-
GradTransposeTensorCollocatedNodes3d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, NULL, s_G, r_V);
654+
GradTransposeTensorCollocatedNodes3d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, nullptr, s_G, r_V);
655655
if (e < num_elem) {
656656
SumElementStrided3d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * BASIS_P_1D * BASIS_P_1D * num_elem,
657657
BASIS_P_1D * BASIS_P_1D * BASIS_P_1D, r_V, d_V);

0 commit comments

Comments
 (0)