Skip to content

Testing direct vs loop stride #500

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Draft
artv3 wants to merge 7 commits into
base: develop
Choose a base branch
from
Draft

Testing direct vs loop stride #500

Show file tree
Hide file tree
Changes from 6 commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
179288f
test with direct policies
artv3 Feb 3, 2025
7926f9b
clean up pass
artv3 Feb 3, 2025
5a2b824
add direct macros -- for testing
artv3 Feb 3, 2025
9d810b6
use direct methods
artv3 Feb 4, 2025
6bacf6b
revert mass
artv3 Feb 4, 2025
c3100c4
use direct for raja too
artv3 Feb 4, 2025
0136ceb
Merge branch 'develop' into artv3/hip-direct-fem-test
artv3 Sep 30, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
58 changes: 29 additions & 29 deletions src/apps/DIFFUSION3DPA-Hip.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -23,7 +23,7 @@ namespace rajaperf {
namespace apps {

template < size_t block_size >
__launch_bounds__(block_size)
__launch_bounds__(block_size)
__global__ void Diffusion3DPA(const Real_ptr Basis,
const Real_ptr dBasis, const Real_ptr D,
const Real_ptr X, Real_ptr Y, bool symmetric) {
Expand All @@ -32,73 +32,73 @@ __global__ void Diffusion3DPA(const Real_ptr Basis,

DIFFUSION3DPA_0_GPU;

GPU_FOREACH_THREAD(dz, z, DPA_D1D) {
GPU_FOREACH_THREAD(dy, y, DPA_D1D) {
GPU_FOREACH_THREAD(dx, x, DPA_D1D) {
GPU_FOREACH_THREAD_DIRECT(dz, z, DPA_D1D) {
GPU_FOREACH_THREAD_DIRECT(dy, y, DPA_D1D) {
GPU_FOREACH_THREAD_DIRECT(dx, x, DPA_D1D) {
DIFFUSION3DPA_1;
}
}
}

if (threadIdx.z == 0) {
GPU_FOREACH_THREAD(dy, y, DPA_D1D) {
GPU_FOREACH_THREAD(qx, x, DPA_Q1D) {
GPU_FOREACH_THREAD_DIRECT(dy, y, DPA_D1D) {
GPU_FOREACH_THREAD_DIRECT(qx, x, DPA_Q1D) {
DIFFUSION3DPA_2;
}
}
}
__syncthreads();
GPU_FOREACH_THREAD(dz, z, DPA_D1D) {
GPU_FOREACH_THREAD(dy, y, DPA_D1D) {
GPU_FOREACH_THREAD(qx, x, DPA_Q1D) {
GPU_FOREACH_THREAD_DIRECT(dz, z, DPA_D1D) {
GPU_FOREACH_THREAD_DIRECT(dy, y, DPA_D1D) {
GPU_FOREACH_THREAD_DIRECT(qx, x, DPA_Q1D) {
DIFFUSION3DPA_3;
}
}
}
__syncthreads();
GPU_FOREACH_THREAD(dz, z, DPA_D1D) {
GPU_FOREACH_THREAD(qy, y, DPA_Q1D) {
GPU_FOREACH_THREAD(qx, x, DPA_Q1D) {
GPU_FOREACH_THREAD_DIRECT(dz, z, DPA_D1D) {
GPU_FOREACH_THREAD_DIRECT(qy, y, DPA_Q1D) {
GPU_FOREACH_THREAD_DIRECT(qx, x, DPA_Q1D) {
DIFFUSION3DPA_4;
}
}
}
__syncthreads();
GPU_FOREACH_THREAD(qz, z, DPA_Q1D) {
GPU_FOREACH_THREAD(qy, y, DPA_Q1D) {
GPU_FOREACH_THREAD(qx, x, DPA_Q1D) {
GPU_FOREACH_THREAD_DIRECT(qz, z, DPA_Q1D) {
GPU_FOREACH_THREAD_DIRECT(qy, y, DPA_Q1D) {
GPU_FOREACH_THREAD_DIRECT(qx, x, DPA_Q1D) {
DIFFUSION3DPA_5;
}
}
}
__syncthreads();
if (threadIdx.z == 0) {
GPU_FOREACH_THREAD(d, y, DPA_D1D) {
GPU_FOREACH_THREAD(q, x, DPA_Q1D) {
GPU_FOREACH_THREAD_DIRECT(d, y, DPA_D1D) {
GPU_FOREACH_THREAD_DIRECT(q, x, DPA_Q1D) {
DIFFUSION3DPA_6;
}
}
}
__syncthreads();
GPU_FOREACH_THREAD(qz, z, DPA_Q1D) {
GPU_FOREACH_THREAD(qy, y, DPA_Q1D) {
GPU_FOREACH_THREAD(dx, x, DPA_D1D) {
GPU_FOREACH_THREAD_DIRECT(qz, z, DPA_Q1D) {
GPU_FOREACH_THREAD_DIRECT(qy, y, DPA_Q1D) {
GPU_FOREACH_THREAD_DIRECT(dx, x, DPA_D1D) {
DIFFUSION3DPA_7;
}
}
}
__syncthreads();
GPU_FOREACH_THREAD(qz, z, DPA_Q1D) {
GPU_FOREACH_THREAD(dy, y, DPA_D1D) {
GPU_FOREACH_THREAD(dx, x, DPA_D1D) {
GPU_FOREACH_THREAD_DIRECT(qz, z, DPA_Q1D) {
GPU_FOREACH_THREAD_DIRECT(dy, y, DPA_D1D) {
GPU_FOREACH_THREAD_DIRECT(dx, x, DPA_D1D) {
DIFFUSION3DPA_8;
}
}
}
__syncthreads();
GPU_FOREACH_THREAD(dz, z, DPA_D1D) {
GPU_FOREACH_THREAD(dy, y, DPA_D1D) {
GPU_FOREACH_THREAD(dx, x, DPA_D1D) {
GPU_FOREACH_THREAD_DIRECT(dz, z, DPA_D1D) {
GPU_FOREACH_THREAD_DIRECT(dy, y, DPA_D1D) {
GPU_FOREACH_THREAD_DIRECT(dx, x, DPA_D1D) {
DIFFUSION3DPA_9;
}
}
Expand Down Expand Up @@ -144,13 +144,13 @@ void DIFFUSION3DPA::runHipVariantImpl(VariantID vid) {
RAJA::LoopPolicy<RAJA::hip_block_x_direct>;

using inner_x =
RAJA::LoopPolicy<RAJA::hip_thread_size_x_loop<DPA_Q1D>>;
RAJA::LoopPolicy<RAJA::hip_thread_size_x_direct<DPA_Q1D>>;

using inner_y =
RAJA::LoopPolicy<RAJA::hip_thread_size_y_loop<DPA_Q1D>>;
RAJA::LoopPolicy<RAJA::hip_thread_size_y_direct<DPA_Q1D>>;

using inner_z =
RAJA::LoopPolicy<RAJA::hip_thread_size_z_loop<DPA_Q1D>>;
RAJA::LoopPolicy<RAJA::hip_thread_size_z_direct<DPA_Q1D>>;

startTimer();
for (RepIndex_type irep = 0; irep < run_reps; ++irep) {
Expand Down
5 changes: 5 additions & 0 deletions src/apps/FEM_MACROS.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,6 +24,11 @@
for (int i = threadIdx.k; i < N; i += blockDim.k)
#endif

#if defined(RAJA_ENABLE_CUDA) || defined(RAJA_ENABLE_HIP)
#define GPU_FOREACH_THREAD_DIRECT(i, k, N) \
if(int i = threadIdx.k; i < N)
Copy link

@khuck-AMD khuck-AMD Jul 16, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I am confused about this line, @artv3 - did you intend for this to be an if statement?

Copy link

@khuck-AMD khuck-AMD Jul 17, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

never mind, I seem to understand what the code is doing now

#endif

#if defined(RAJA_ENABLE_SYCL)
#define SYCL_FOREACH_THREAD(i, k, N) \
for (int i = itm.get_local_id(k); i < N; i += itm.get_local_range(k))
Expand Down