diff --git a/src/fields/fft_poisson_solver/FFTPoissonSolverDirichletFast.cpp b/src/fields/fft_poisson_solver/FFTPoissonSolverDirichletFast.cpp index 75c7bffb21..0dd69ce288 100644 --- a/src/fields/fft_poisson_solver/FFTPoissonSolverDirichletFast.cpp +++ b/src/fields/fft_poisson_solver/FFTPoissonSolverDirichletFast.cpp @@ -102,7 +102,7 @@ void ToSine_Transpose_ToComplex (const Array2 in, return to_sine(in, j, i, n_data, sine_factor); }; -#if defined(AMREX_USE_CUDA) || defined(AMREX_USE_HIP) +#if defined(AMREX_USE_GPU) constexpr int tile_dim_x = 16; constexpr int tile_dim_x_ex = 34; constexpr int tile_dim_y = 32; @@ -118,13 +118,11 @@ void ToSine_Transpose_ToComplex (const Array2 in, const int num_blocks_x = (nx + tile_dim_x - 1)/tile_dim_x; const int num_blocks_y = (ny + tile_dim_y - 1)/tile_dim_y; - amrex::launch(num_blocks_x*num_blocks_y, amrex::Gpu::gpuStream(), - [=] AMREX_GPU_DEVICE() noexcept + amrex::LaunchRaw( + amrex::IntVectND<2>{num_blocks_x, num_blocks_y}, tile_dim_x_ex * tile_dim_y, + [=] AMREX_GPU_DEVICE(auto lh) noexcept { - __shared__ amrex::Real tile_ptr[tile_dim_x_ex * tile_dim_y]; - - const int block_y = blockIdx.x / num_blocks_x; - const int block_x = blockIdx.x - block_y*num_blocks_x; + const auto [block_x, block_y] = lh.blockIdxND(); const int tile_begin_x = 2 * block_x * tile_dim_x - 2; const int tile_begin_y = block_y * tile_dim_y; @@ -132,12 +130,13 @@ void ToSine_Transpose_ToComplex (const Array2 in, const int tile_end_x = tile_begin_x + tile_dim_x_ex; const int tile_end_y = tile_begin_y + tile_dim_y; - Array2 shared{{tile_ptr, {tile_begin_x, tile_begin_y, 0}, - {tile_end_x, tile_end_y, 1}, 1}}; + Array2 shared{{lh.shared_memory(), + {tile_begin_x, tile_begin_y, 0}, + {tile_end_x, tile_end_y, 1}, 1}}; { - const int thread_x = threadIdx.x / tile_dim_y; - const int thread_y = threadIdx.x - thread_x*tile_dim_y; + const auto [thread_y, thread_x] = + lh.template threadIdxND(); for (int tx = thread_x; tx < tile_dim_x_ex; tx += block_rows_x) { const int i = tile_begin_x + tx; @@ -149,11 +148,11 @@ void ToSine_Transpose_ToComplex (const Array2 in, } } - __syncthreads(); + lh.syncthreads(); { - const int thread_y = threadIdx.x / tile_dim_x; - const int thread_x = threadIdx.x - thread_y*tile_dim_x; + const auto [thread_x, thread_y] = + lh.template threadIdxND(); for (int ty = thread_y; ty < tile_dim_y; ty += block_rows_y) { const int i = block_x * tile_dim_x + thread_x; diff --git a/src/fields/fft_poisson_solver/FFTPoissonSolverDirichletQuick.cpp b/src/fields/fft_poisson_solver/FFTPoissonSolverDirichletQuick.cpp index 8b2ee2e621..4d36d76c54 100644 --- a/src/fields/fft_poisson_solver/FFTPoissonSolverDirichletQuick.cpp +++ b/src/fields/fft_poisson_solver/FFTPoissonSolverDirichletQuick.cpp @@ -96,7 +96,7 @@ dst2_out_mult_dst3_in (Array2> const& inout, int inline void dst2_out_t_in (Array2> const& in, Array2 const& out, int nx, int ny, const amrex::GpuComplex* omega) { -#if defined(AMREX_USE_CUDA) || defined(AMREX_USE_HIP) +#if defined(AMREX_USE_GPU) static constexpr int tile_dim_x = 16; static constexpr int tile_dim_y = 64; static constexpr int elem_per_thread = 2; @@ -108,58 +108,60 @@ dst2_out_t_in (Array2> const& in, Array2(num_blocks_x*num_blocks_y, amrex::Gpu::gpuStream(), - [=] AMREX_GPU_DEVICE() noexcept + amrex::LaunchRaw( + amrex::IntVectND<2>{num_blocks_y, num_blocks_x}, tile_dim_x * (tile_dim_y+1) * 2, + [=] AMREX_GPU_DEVICE(auto lh) noexcept { - __shared__ amrex::Real tile_real[tile_dim_x][tile_dim_y+1]; - __shared__ amrex::Real tile_imag[tile_dim_x][tile_dim_y+1]; + Array3 shared{{lh.shared_memory(), {0, 0, 0}, + {tile_dim_y+1, tile_dim_x, 1}, 2}}; - const int block_x = blockIdx.x / num_blocks_y; - const int block_y = blockIdx.x - block_x*num_blocks_y; + const auto [block_y, block_x] = lh.blockIdxND(); const int tile_start_x = tile_begin_x + block_x* tile_dim_x; const int tile_start_y = block_y * tile_dim_y; - int thread_y = threadIdx.x / tile_dim_x; - int thread_x = threadIdx.x - thread_y*tile_dim_x; - - #pragma unroll elem_per_thread - for (; thread_y < tile_dim_y; thread_y += block_rows_y) { - const int thread_xr = tile_dim_x - thread_x - 1; - const int iout = tile_start_x + thread_xr; - const int jout = tile_start_y + thread_y; - const int iin = nx-iout-1; - const int jin = 2*jout < ny ? 2*jout : 2*(ny-jout)-1; - - if (iout < nx && jout < ny) { - auto val = in(iin, jin) * omega[iin]; - if (2*jout < ny) { - val = -val; + auto [thread_x, thread_y] = lh.template threadIdxND(); + + #pragma unroll elem_per_thread + for (; thread_y < tile_dim_y; thread_y += block_rows_y) + { + const int thread_xr = tile_dim_x - thread_x - 1; + const int iout = tile_start_x + thread_xr; + const int jout = tile_start_y + thread_y; + const int iin = nx-iout-1; + const int jin = 2*jout < ny ? 2*jout : 2*(ny-jout)-1; + + if (iout < nx && jout < ny) { + auto val = in(iin, jin) * omega[iin]; + if (2*jout < ny) { + val = -val; + } + shared(thread_y, thread_xr, 0) = val.real(); + shared(thread_y, thread_xr, 1) = val.imag(); } - tile_real[thread_xr][thread_y] = val.real(); - tile_imag[thread_xr][thread_y] = val.imag(); } } - __syncthreads(); - - thread_x = threadIdx.x / tile_dim_y; - thread_y = threadIdx.x - thread_x*tile_dim_y; + lh.syncthreads(); - #pragma unroll elem_per_thread - for (; thread_x < tile_dim_x; thread_x += block_rows_x) { - const int iout = tile_start_x + thread_x; - const int jout = tile_start_y + thread_y; - const int iin = nx-iout-1; - const int iout2 = iin-1; - const bool do_iout2 = iout2 >= 0 && iout2 != iout; - - if (iout < nx && jout < ny) { - out(jout, iout) = -tile_real[thread_x][thread_y]; - if (do_iout2) { - out(jout, iout2) = tile_imag[thread_x][thread_y]; + auto [thread_y, thread_x] = lh.template threadIdxND(); + + #pragma unroll elem_per_thread + for (; thread_x < tile_dim_x; thread_x += block_rows_x) + { + const int iout = tile_start_x + thread_x; + const int jout = tile_start_y + thread_y; + const int iin = nx-iout-1; + const int iout2 = iin-1; + const bool do_iout2 = iout2 >= 0 && iout2 != iout; + + if (iout < nx && jout < ny) { + out(jout, iout) = -shared(thread_y, thread_x, 0); + if (do_iout2) { + out(jout, iout2) = shared(thread_y, thread_x, 1); + } } } } @@ -180,7 +182,7 @@ dst2_out_t_in (Array2> const& in, Array2 const& in, Array2> const& out, int nx, int ny, const amrex::GpuComplex* omega) { -#if defined(AMREX_USE_CUDA) || defined(AMREX_USE_HIP) +#if defined(AMREX_USE_GPU) static constexpr int tile_dim_x = 16; static constexpr int tile_dim_y = 64; static constexpr int elem_per_thread = 2; @@ -192,63 +194,65 @@ dst3_out_t_in (Array2 const& in, Array2(num_blocks_x*num_blocks_y, amrex::Gpu::gpuStream(), - [=] AMREX_GPU_DEVICE() noexcept + amrex::LaunchRaw( + amrex::IntVectND<2>{num_blocks_x, num_blocks_y}, tile_dim_y * (tile_dim_x+1) * 2, + [=] AMREX_GPU_DEVICE(auto lh) noexcept { - __shared__ amrex::Real tile_real[tile_dim_y][tile_dim_x+1]; - __shared__ amrex::Real tile_imag[tile_dim_y][tile_dim_x+1]; + Array3 shared{{lh.shared_memory(), {0, 0, 0}, + {tile_dim_x+1, tile_dim_y, 1}, 2}}; - const int block_y = blockIdx.x / num_blocks_x; - const int block_x = blockIdx.x - block_y*num_blocks_x; + const auto [block_x, block_y] = lh.blockIdxND(); const int tile_start_x = block_x * tile_dim_x; const int tile_start_y = block_y * tile_dim_y; - int thread_x = threadIdx.x / tile_dim_y; - int thread_y = threadIdx.x - thread_x*tile_dim_y; - - #pragma unroll elem_per_thread - for (; thread_x < tile_dim_x; thread_x += block_rows_x) { - const int thread_yp = (thread_y*2) % tile_dim_y + (thread_y*2) / tile_dim_y; - const int iout = tile_start_x + thread_x; - const int jout = tile_start_y + thread_yp; - const int jin = jout%2 == 0 ? jout/2 : ny-1-jout/2; - const int iin1 = nx-iout-1; - const int iin2 = iout - 1; - - if (iout < tile_end_x && jout < ny) { - amrex::GpuComplex val { - in(jin, iin1), - iout != 0 ? -in(jin, iin2) : 0 - }; - auto o = omega[iout]; - o.m_imag = - o.m_imag; - if (jout%2 != 0) { - val = -val; + auto [thread_y, thread_x] = lh.template threadIdxND(); + + #pragma unroll elem_per_thread + for (; thread_x < tile_dim_x; thread_x += block_rows_x) + { + const int thread_yp = (thread_y*2) % tile_dim_y + (thread_y*2) / tile_dim_y; + const int iout = tile_start_x + thread_x; + const int jout = tile_start_y + thread_yp; + const int jin = jout%2 == 0 ? jout/2 : ny-1-jout/2; + const int iin1 = nx-iout-1; + const int iin2 = iout - 1; + + if (iout < tile_end_x && jout < ny) { + amrex::GpuComplex val { + in(jin, iin1), + iout != 0 ? -in(jin, iin2) : 0 + }; + auto o = omega[iout]; + o.m_imag = - o.m_imag; + if (jout%2 != 0) { + val = -val; + } + val *= o; + shared(thread_x, thread_yp, 0) = val.real(); + shared(thread_x, thread_yp, 1) = val.imag(); } - val *= o; - tile_real[thread_yp][thread_x] = val.real(); - tile_imag[thread_yp][thread_x] = val.imag(); } } - __syncthreads(); + lh.syncthreads(); - thread_y = threadIdx.x / tile_dim_x; - thread_x = threadIdx.x - thread_y*tile_dim_x; - - #pragma unroll elem_per_thread - for (; thread_y < tile_dim_y; thread_y += block_rows_y) { - const int iout = tile_start_x + thread_x; - const int jout = tile_start_y + thread_y; - - if (iout < tile_end_x && jout < ny) { - out(iout, jout) = { - tile_real[thread_y][thread_x], - tile_imag[thread_y][thread_x] - }; + auto [thread_x, thread_y] = lh.template threadIdxND(); + + #pragma unroll elem_per_thread + for (; thread_y < tile_dim_y; thread_y += block_rows_y) + { + const int iout = tile_start_x + thread_x; + const int jout = tile_start_y + thread_y; + + if (iout < tile_end_x && jout < ny) { + out(iout, jout) = { + shared(thread_x, thread_y, 0), + shared(thread_x, thread_y, 1) + }; + } } } }); diff --git a/src/mg_solver/HpMultiGrid.cpp b/src/mg_solver/HpMultiGrid.cpp index cd8d2cbcbf..c8c8ef8353 100644 --- a/src/mg_solver/HpMultiGrid.cpp +++ b/src/mg_solver/HpMultiGrid.cpp @@ -415,7 +415,7 @@ void gsrb (int icolor, amrex::Box const& box, Array3 const& phi, // Optimized Gauss-Seidel update combined with compute residual: /////////////////////////////////// //////////////////////////////////////////////////////////////////////////////////////////////////// -#if defined(AMREX_USE_CUDA) || defined(AMREX_USE_HIP) +#if defined(AMREX_USE_GPU) // do multiple gsrb iterations in GPU shared memory with many ghost cells template @@ -455,17 +455,12 @@ void gsrb_shared (amrex::Box const& box, Array3 const& phi_out, int const jhi_loop = loop_box.bigEnd(1); const int num_blocks_x = (loop_box.length(0) + final_tilesize_x - 1)/final_tilesize_x; const int num_blocks_y = (loop_box.length(1) + final_tilesize_y - 1)/final_tilesize_y; - amrex::Math::FastDivmodU64 num_blocks_divmod {static_cast(num_blocks_x)}; - amrex::launch(num_blocks_x*num_blocks_y, amrex::Gpu::gpuStream(), - [=] AMREX_GPU_DEVICE() noexcept + amrex::LaunchRaw( + amrex::IntVectND<2>{num_blocks_x, num_blocks_y}, num_cells_in_tile, + [=] AMREX_GPU_DEVICE(auto lh) noexcept { - // allocate static shared memory - __shared__ amrex::Real phi_ptr[num_cells_in_tile]; - - std::uint64_t remainder = 0; - const int iblock_y = num_blocks_divmod.divmod(remainder, blockIdx.x); - const int iblock_x = remainder; + const auto [iblock_x, iblock_y] = lh.blockIdxND(); const int tile_begin_x = iblock_x * final_tilesize_x - edge_offset - 1 + ilo_loop; const int tile_begin_y = iblock_y * final_tilesize_y - edge_offset - 1 + jlo_loop; @@ -474,36 +469,39 @@ void gsrb_shared (amrex::Box const& box, Array3 const& phi_out, const int tile_end_y = tile_begin_y + tilesize_array_y; // make Array3 reference shared memory tile - Array3 phi_shared({phi_ptr, {tile_begin_x, tile_begin_y, 0}, - {tile_end_x, tile_end_y, 1}, num_comps}); + Array3 phi_shared({lh.shared_memory(), + {tile_begin_x, tile_begin_y, 0}, + {tile_end_x, tile_end_y, 1}, num_comps}); if (zero_init) { // initialize shared memory to zero - for (int s = threadIdx.x; s < num_cells_in_tile; s+=blockDim.x) { - phi_ptr[s] = amrex::Real(0.); + for (int s = lh.threadIdx1D(); s < num_cells_in_tile; s+=lh.blockDim1D()) { + phi_shared.p[s] = amrex::Real(0.); } } else { // initialize shared memory to phi_in inside the domain, outside zero - for (int s = threadIdx.x; s < tilesize_array_x*tilesize_array_y; s+=blockDim.x) { - int sy = s / tilesize_array_x; - int sx = s - sy * tilesize_array_x; - sx += tile_begin_x; - sy += tile_begin_y; - if (ilo_loop <= sx && sx <= ihi_loop && - jlo_loop <= sy && sy <= jhi_loop) { + constexpr int tilesize_array = tilesize_array_x * tilesize_array_y; + int s = lh.threadIdx1D(); + // loop with compile time known bounds to help compiler + for (int sc = 0; sc < tilesize_array; sc+=lh.blockDim1D()) { + // do runtime check only on last iteration + if (sc + lh.blockDim1D() - 1 < tilesize_array || + s < tilesize_array) { + int sy = s / tilesize_array_x; + int sx = s - sy * tilesize_array_x; + sx += tile_begin_x; + sy += tile_begin_y; + const bool is_inside = ilo_loop <= sx && sx <= ihi_loop && + jlo_loop <= sy && sy <= jhi_loop; for (int n=0; n(); ithread_y *= 2; const int i = tile_begin_x + 1 + ithread_x; @@ -525,7 +523,7 @@ void gsrb_shared (amrex::Box const& box, Array3 const& phi_out, } } - __syncthreads(); + lh.syncthreads(); for (int icolor=0; icolor const& phi_out, rhs_loc[0], facx, facy); } } - __syncthreads(); + lh.syncthreads(); } for (int nj=0; nj<=1; ++nj) { @@ -598,7 +596,7 @@ void gsrb_shared (amrex::Box const& box, Array3 const& phi_out, }); } -#elif !defined(AMREX_USE_GPU) +#else // do multiple gsrb iterations in CPU cached memory with many ghost cells template @@ -770,7 +768,7 @@ void gsrb_4_residual (int system_type, amrex::Box const& box, // Compute 4 gsrb (Gauss-Seidel red-black) iterations using rhs and acf and // store the result in phi_out. // If do_compute_residual is set, store the residual in res. -#if defined(AMREX_USE_CUDA) || defined(AMREX_USE_HIP) +#if defined(AMREX_USE_GPU) if (system_type == 1) { if (box.cellCentered()) { gsrb_shared<1, zero_init, do_compute_residual, true>( @@ -856,12 +854,6 @@ void gsrb_4_residual (int system_type, amrex::Box const& box, #if defined(AMREX_USE_GPU) -#if defined(AMREX_USE_DPCPP) -#define HPMG_SYNCTHREADS item.barrier(sycl::access::fence_space::global_and_local) -#else -#define HPMG_SYNCTHREADS __syncthreads() -#endif - template void bottomsolve_gpu (amrex::Real dx0, amrex::Real dy0, Array3 const* acf, Array3 const* res, Array3 const* cor, @@ -876,24 +868,14 @@ void bottomsolve_gpu (amrex::Real dx0, amrex::Real dy0, Array3 cons // Currently, this function does not use shared memory. static_assert(n_cell_single*n_cell_single <= 1024, "n_cell_single is too big"); -#if defined(AMREX_USE_DPCPP) - amrex::launch(system_type == 1 ? 2 : 1, 1024, amrex::Gpu::gpuStream(), - [=] (sycl::nd_item<1> const& item) noexcept -#else - amrex::launch_global<1024><<>>( - [=] AMREX_GPU_DEVICE () noexcept -#endif + amrex::LaunchRaw<1024>(amrex::IntVectND<1>{system_type == 1 ? 2 : 1}, + [=] AMREX_GPU_DEVICE (auto lh) noexcept { amrex::Real facx = amrex::Real(1.)/(dx0*dx0); amrex::Real facy = amrex::Real(1.)/(dy0*dy0); const int lenx = lev_domain[0].length(0) - 2*corner_offset; -#if defined(AMREX_USE_DPCPP) - const int icell = item.get_local_linear_id(); - const int n = system_type == 1 ? item.get_group_linear_id() : 0; -#else - const int icell = threadIdx.x; - const int n = system_type == 1 ? blockIdx.x : 0; -#endif + const int icell = lh.threadIdx1D(); + const int n = system_type == 1 ? lh.blockIdx1D() : 0; int j = icell / lenx; int i = icell - j*lenx; j += lev_domain[0].smallEnd(1) + corner_offset; @@ -912,7 +894,7 @@ void bottomsolve_gpu (amrex::Real dx0, amrex::Real dy0, Array3 cons cor[ilev](i,j,0) = amrex::Real(0.); } } - HPMG_SYNCTHREADS; + lh.syncthreads(); // do 4 Gauss-Seidel red-black iterations for (int is = 0; is < 4; ++is) { @@ -925,7 +907,7 @@ void bottomsolve_gpu (amrex::Real dx0, amrex::Real dy0, Array3 cons facx, facy); } } - HPMG_SYNCTHREADS; + lh.syncthreads(); } // calculate the residual @@ -938,7 +920,7 @@ void bottomsolve_gpu (amrex::Real dx0, amrex::Real dy0, Array3 cons res[ilev], acf[ilev], facx, facy); } - HPMG_SYNCTHREADS; + lh.syncthreads(); // interpolate residual to next level is_active = @@ -967,7 +949,7 @@ void bottomsolve_gpu (amrex::Real dx0, amrex::Real dy0, Array3 cons } } } - HPMG_SYNCTHREADS; + lh.syncthreads(); facx *= amrex::Real(0.25); facy *= amrex::Real(0.25); @@ -987,7 +969,7 @@ void bottomsolve_gpu (amrex::Real dx0, amrex::Real dy0, Array3 cons cor[ilev](i,j,0) = amrex::Real(0.); } } - HPMG_SYNCTHREADS; + lh.syncthreads(); // do NS Gauss-Seidel red-black iterations for (int is = 0; is < NS; ++is) { @@ -1000,7 +982,7 @@ void bottomsolve_gpu (amrex::Real dx0, amrex::Real dy0, Array3 cons facx, facy); } } - HPMG_SYNCTHREADS; + lh.syncthreads(); } } @@ -1040,7 +1022,7 @@ void bottomsolve_gpu (amrex::Real dx0, amrex::Real dy0, Array3 cons // do 4 Gauss-Seidel red-black iterations for (int is = 0; is < 4; ++is) { - HPMG_SYNCTHREADS; + lh.syncthreads(); if (is_active) { if ((i+j+is)%2 == 0) { fgs(i, j, n, @@ -1654,26 +1636,15 @@ namespace { void avgdown_acf (Array3 const* acf, amrex::BoxND<2> const* lev_domain, int ncomp, int nlevels, F&& f) { -#if defined(AMREX_USE_DPCPP) - amrex::launch(1, 1024, amrex::Gpu::gpuStream(), - [=] (sycl::nd_item<1> const& item) noexcept -#else - amrex::launch_global<1024><<<1, 1024, 0, amrex::Gpu::gpuStream()>>>( - [=] AMREX_GPU_DEVICE () noexcept -#endif + amrex::LaunchRaw<1024>(amrex::IntVectND<1>{1}, + [=] AMREX_GPU_DEVICE (auto lh) noexcept { for (int ilev = 1; ilev < nlevels; ++ilev) { const int lenx = lev_domain[ilev].length(0); const int leny = lev_domain[ilev].length(1); const int ncells = lenx*leny; -#if defined(AMREX_USE_DPCPP) - for (int icell = item.get_local_range(0)*item.get_group_linear_id() - + item.get_local_linear_id(), - stride = item.get_local_range(0)*item.get_group_range(0); -#else - for (int icell = blockDim.x*blockIdx.x+threadIdx.x, stride = blockDim.x*gridDim.x; -#endif - icell < ncells; icell += stride) { + + for (int icell = lh.threadIdx1D(); icell < ncells; icell += lh.blockDim1D()) { int j = icell / lenx; int i = icell - j*lenx; j += lev_domain[ilev].smallEnd(1); @@ -1682,7 +1653,7 @@ namespace { f(i,j,n,acf[ilev],acf[ilev-1],lev_domain[ilev]); } } - HPMG_SYNCTHREADS; + lh.syncthreads(); } }); } diff --git a/src/particles/deposition/DepositionUtil.H b/src/particles/deposition/DepositionUtil.H index bda9ae156f..2495d15239 100644 --- a/src/particles/deposition/DepositionUtil.H +++ b/src/particles/deposition/DepositionUtil.H @@ -104,20 +104,14 @@ SharedMemoryDeposition (int num_particles, } }); - amrex::Math::FastDivmodU64 num_tiles_divmod {static_cast(ntile_x)}; - // launch shared memory kernel to deposit the charge / current // use one block per tile and one thread per cell - amrex::launch(ntile_x * ntile_y, amrex::Gpu::gpuStream(), - [=] AMREX_GPU_DEVICE () { - // allocate static shared memory - __shared__ amrex::Real shared_ptr[tile_s_x * tile_s_y * (max_cache + max_depos)]; - - const int tile_id = blockIdx.x; + amrex::LaunchRaw( + amrex::IntVectND<2>{ntile_x, ntile_y}, tile_s_x * tile_s_y * (max_cache + max_depos), + [=] AMREX_GPU_DEVICE (auto lh) { + const int tile_id = lh.blockIdx1D(); - std::uint64_t remainder = 0; - const int tile_id_y = num_tiles_divmod.divmod(remainder, tile_id); - const int tile_id_x = remainder; + const auto [tile_id_x, tile_id_y] = lh.blockIdxND(); const int tile_begin_x = lo_x + tile_id_x * tile_x; const int tile_begin_y = lo_y + tile_id_y * tile_y; @@ -127,13 +121,13 @@ SharedMemoryDeposition (int num_particles, // make Array3 reference shared memory tile Array3 shared_arr{{ - shared_ptr, + lh.shared_memory(), {tile_begin_x, tile_begin_y, 0}, {tile_end_x, tile_end_y, 1}, max_cache + max_depos }}; - for (int s = threadIdx.x; s < tile_s_x * tile_s_y; s+=threads_per_tile) { + for (int s = lh.threadIdx1D(); s < tile_s_x * tile_s_y; s+=threads_per_tile) { int sy = s / tile_s_x; int sx = s - sy * tile_s_x; sx += tile_begin_x; @@ -166,19 +160,19 @@ SharedMemoryDeposition (int num_particles, loc_idx_depos[n] = (dynamic_comps && idx_depos[n]==-1) ? -1 : n+int(max_cache); } - __syncthreads(); + lh.syncthreads(); // deposit the charge / current of all particles in the linked list - int current_idx = p_ll_start[tile_id * threads_per_tile + threadIdx.x]; + int current_idx = p_ll_start[tile_id * threads_per_tile + lh.threadIdx1D()]; while (current_idx != ll_guard) { do_deposit(current_idx, ptd, shared_arr, loc_idx_cache, loc_idx_depos); current_idx = p_ll_next[current_idx]; } - __syncthreads(); + lh.syncthreads(); // add to local charge / current to the global one - for (int s = threadIdx.x; s < tile_s_x * tile_s_y; s+=threads_per_tile) { + for (int s = lh.threadIdx1D(); s < tile_s_x * tile_s_y; s+=threads_per_tile) { int sy = s / tile_s_x; int sx = s - sy * tile_s_x; sx += tile_begin_x; @@ -192,8 +186,6 @@ SharedMemoryDeposition (int num_particles, } } } - - } );