OilfieldToolsNet/opm-simulators
4
0
Fork 0
mirror of https://github.com/OPM/opm-simulators.git synced 2024-12-01 13:29:08 -06:00
Code Issues Packages Projects Releases Wiki Activity

Fix stdwell_apply opencl kernel

Browse Source
This commit is contained in:
Tong Dong Qiu 2021-12-15 15:19:17 +01:00 committed by Markus Blatt
parent 994260aaea
commit 3decb73561
3 changed files with 16 additions and 13 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
opm/simulators/linalg/bda/cuWellContributions.cu
View File

@ -80,11 +80,10 @@ __global__ void apply_well_contributions(
}
// merge all blocks into 1 dim*dim_wells block
// since NORNE has only 2 parallel blocks, do not use a loop
// since 3*4 blocks has give 2 parallel blocks, do not use a loop
// 0x00ffffff contains 24 ones, representing the two blocks that are added
temp += __shfl_down_sync(0x00ffffff, temp, dim * dim_wells);
b = idx_t / vals_per_block + val_pointers[idx_b];
// merge all (dim) columns of 1 block, results in a single 1*dim_wells vector, which is used to multiply with invD
if (idx_t < vals_per_block) {
// should be a loop as well, now only works for dim == 3

12
opm/simulators/linalg/bda/opencl/kernels/stdwell_apply.cl
View File

@ -37,15 +37,17 @@ __kernel void stdwell_apply(
b += numBlocksPerWarp;
}
// merge all blocks in this workgroup into 1 block
// if numBlocksPerWarp >= 3, should use loop
if(wiId < valsPerBlock){
localSum[wiId] += localSum[wiId + valsPerBlock];
for (int i = 1; i < numBlocksPerWarp; ++i) {
localSum[wiId] += localSum[wiId + i*valsPerBlock];
}
}
b = wiId/valsPerBlock + val_pointers[wgId];
if(c == 0 && wiId < valsPerBlock){
for(unsigned int stride = 2; stride > 0; stride >>= 1){
localSum[wiId] += localSum[wiId + stride];
for(unsigned int i = dim - 1; i > 0; --i){
localSum[wiId] += localSum[wiId + i];
}
z1[r] = localSum[wiId];
}

12
opm/simulators/linalg/bda/opencl/kernels/stdwell_apply_no_reorder.cl
View File

@ -35,15 +35,17 @@ __kernel void stdwell_apply_no_reorder(
b += numBlocksPerWarp;
}
// merge all blocks in this workgroup into 1 block
// if numBlocksPerWarp >= 3, should use loop
if(wiId < valsPerBlock){
localSum[wiId] += localSum[wiId + valsPerBlock];
for (int i = 1; i < numBlocksPerWarp; ++i) {
localSum[wiId] += localSum[wiId + i*valsPerBlock];
}
}
b = wiId/valsPerBlock + val_pointers[wgId];
if(c == 0 && wiId < valsPerBlock){
for(unsigned int stride = 2; stride > 0; stride >>= 1){
localSum[wiId] += localSum[wiId + stride];
for(unsigned int i = dim - 1; i > 0; --i){
localSum[wiId] += localSum[wiId + i];
}
z1[r] = localSum[wiId];
}