OilfieldToolsNet/LBPM
4
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Refactoring flux bC for use with cooperative groups

Browse Source
This commit is contained in:
James E McClure
2018-02-20 21:31:37 -05:00
parent 030f997516
commit 73a2347bef
1 changed files with 18 additions and 6 deletions
Download Patch File Download Diff File Expand all files Collapse all files

24
gpu/D3Q19.cu
View File

@@ -1960,9 +1960,15 @@ __global__ void dvc_ScaLBL_D3Q19_AAeven_Flux_BC_z(int *list, double *dist, doub
sum = factor*(f0+f1+f2+f3+f4+f7+f8+f9+f10 + 2*(f6+f12+f13+f16+f17)); sum = factor*(f0+f1+f2+f3+f4+f7+f8+f9+f10 + 2*(f6+f12+f13+f16+f17));
} }
sum = blockReduceSum(sum); //sum = blockReduceSum(sum);
if (threadIdx.x==0) //if (threadIdx.x==0)
atomicAdd(dvcsum, sum); // atomicAdd(dvcsum, sum);
extern __shared__ double temp[];
thread_group g = this_thread_block();
double block_sum = reduce_sum(g, temp, sum);
if (g.thread_rank() == 0) atomicAdd(dvcsum, block_sum);
} }
@@ -2029,9 +2035,15 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_Flux_BC_z(int *d_neighborList, int *list
} }
sum = blockReduceSum(sum); //sum = blockReduceSum(sum);
if (threadIdx.x==0) //if (threadIdx.x==0)
atomicAdd(dvcsum, sum); // atomicAdd(dvcsum, sum);
extern __shared__ double temp[];
thread_group g = this_thread_block();
double block_sum = reduce_sum(g, temp, sum);
if (g.thread_rank() == 0) atomicAdd(dvcsum, block_sum);
} }