OilfieldToolsNet/opm-simulators
4
0
Fork 0
mirror of https://github.com/OPM/opm-simulators.git synced 2025-02-25 18:55:30 -06:00
Code Issues Packages Projects Releases Wiki Activity

Only compile HIP kernel with hipcc

Browse Source
This commit is contained in:
Tong Dong Qiu 2023-04-14 13:00:29 +02:00 committed by Razvan Nane
parent ef6be5859e
commit abce3a897c
1 changed files with 8 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

9
opm/simulators/linalg/bda/rocsparseWellContributions.cpp
View File

@ -54,6 +54,7 @@
namespace Opm namespace Opm
{ {
#ifdef __HIP__
/// HIP kernel to apply the standard wellcontributions /// HIP kernel to apply the standard wellcontributions
__global__ void stdwell_apply( __global__ void stdwell_apply(
const double *Cnnzs, const double *Cnnzs,
@ -135,9 +136,12 @@ __global__ void stdwell_apply(
y[colIdx*dim + c] -= temp; y[colIdx*dim + c] -= temp;
} }
} }
#endif
void WellContributionsRocsparse::apply_stdwells(double *d_x, double *d_y){ void WellContributionsRocsparse::apply_stdwells([[maybe_unused]] double *d_x,
[[maybe_unused]] double *d_y){
#ifdef __HIP__
unsigned gridDim = num_std_wells; unsigned gridDim = num_std_wells;
unsigned blockDim = 32; unsigned blockDim = 32;
unsigned shared_mem_size = (blockDim + 2 * dim_wells) * sizeof(double); // shared memory for localSum, z1 and z2 unsigned shared_mem_size = (blockDim + 2 * dim_wells) * sizeof(double); // shared memory for localSum, z1 and z2
@ -147,6 +151,9 @@ void WellContributionsRocsparse::apply_stdwells(double *d_x, double *d_y){
d_x, d_y, dim, dim_wells, d_val_pointers_hip d_x, d_y, dim, dim_wells, d_val_pointers_hip
); );
HIP_CHECK(hipStreamSynchronize(stream)); HIP_CHECK(hipStreamSynchronize(stream));
#else
OPM_THROW(std::logic_error, "Error separate wellcontributions for rocsparse only supported when compiling with hipcc");
#endif
} }
void WellContributionsRocsparse::apply_mswells(double *d_x, double *d_y){ void WellContributionsRocsparse::apply_mswells(double *d_x, double *d_y){