Merge pull request #1628 from LLNL/artv3/sycl-thread-grid-layout-fix

Swap ordering of thread configuration in Sycl
LLNL · Apr 30, 2024 · c315ddd · c315ddd
2 parents e7e2477 + aef818c
commit c315ddd
Show file tree

Hide file tree

Showing 8 changed files with 42 additions and 26 deletions.
diff --git a/docs/sphinx/user_guide/feature/policies.rst b/docs/sphinx/user_guide/feature/policies.rst
@@ -524,10 +524,26 @@ write more explicit policies.
             ignored. For example in cuda_thread_x_direct block_size is
             unspecified so a runtime number of threads is used, but grid_size is
             ignored so blocks are ignored when getting indices.
-
+	    
 GPU Policies for SYCL
 ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
 
+.. note:: SYCL uses C++-style ordering in which the right
+	  most index corresponds to having unit stride.
+	  In a three-dimensional compute grid this means
+	  that dimension 2 has the unit stride while
+	  dimension 0 has the longest stride. This is
+	  important to note as the ordering is reverse
+	  compared to the CUDA and HIP programming models.
+	  CUDA and HIP employ a x/y/z ordering in which
+	  dimension x has the unit stride.
+
+	  When using RAJA::launch, thread and team configuration
+	  follows CUDA and HIP programming models and is always
+	  configured in three-dimensions. This means that dimension
+	  2 always exists and should be used as one would use the
+	  x dimension for CUDA and HIP.
+
  ======================================== ============= ==============================
  SYCL Execution Policies                  Works with    Brief description
  ======================================== ============= ==============================

diff --git a/include/RAJA/policy/sycl/launch.hpp b/include/RAJA/policy/sycl/launch.hpp
@@ -56,13 +56,13 @@ struct LaunchExecute<RAJA::sycl_launch_t<async, 0>> {
     // Compute the number of blocks and threads
     //
 
-    const ::sycl::range<3> blockSize(params.threads.value[0],
+    const ::sycl::range<3> blockSize(params.threads.value[2],
 				     params.threads.value[1],
-				     params.threads.value[2]);
+				     params.threads.value[0]);
 
-    const ::sycl::range<3> gridSize(params.threads.value[0] * params.teams.value[0],
+    const ::sycl::range<3> gridSize(params.threads.value[2] * params.teams.value[2],
 				    params.threads.value[1] * params.teams.value[1],
-				    params.threads.value[2] * params.teams.value[2]);
+				    params.threads.value[0] * params.teams.value[0]);
 
     // Only launch kernel if we have something to iterate over
     constexpr size_t zero = 0;
@@ -138,13 +138,13 @@ struct LaunchExecute<RAJA::sycl_launch_t<async, 0>> {
     // Compute the number of blocks and threads
     //
 
-    const ::sycl::range<3> blockSize(params.threads.value[0],
+    const ::sycl::range<3> blockSize(params.threads.value[2],
 				     params.threads.value[1],
-				     params.threads.value[2]);
+				     params.threads.value[0]);
 
-    const ::sycl::range<3> gridSize(params.threads.value[0] * params.teams.value[0],
+    const ::sycl::range<3> gridSize(params.threads.value[2] * params.teams.value[2],
 				    params.threads.value[1] * params.teams.value[1],
-				    params.threads.value[2] * params.teams.value[2]);
+				    params.threads.value[0] * params.teams.value[0]);
 
     // Only launch kernel if we have something to iterate over
     constexpr size_t zero = 0;

diff --git a/test/include/RAJA_test-launch-direct-teams-threads-1D-execpol.hpp b/test/include/RAJA_test-launch-direct-teams-threads-1D-execpol.hpp
@@ -81,8 +81,8 @@ using Hip_launch_policies = camp::list<hip_direct_policies>;
 using sycl_direct_policies =
   camp::list<
              RAJA::LaunchPolicy<RAJA::sycl_launch_t<true>>,
-             RAJA::LoopPolicy<RAJA::sycl_group_0_direct>,
-             RAJA::LoopPolicy<RAJA::sycl_local_0_direct>
+             RAJA::LoopPolicy<RAJA::sycl_group_2_direct>,
+             RAJA::LoopPolicy<RAJA::sycl_local_2_direct>
             >;
 
 using Sycl_launch_policies = camp::list<sycl_direct_policies>;

diff --git a/test/include/RAJA_test-launch-direct-teams-threads-3D-execpol.hpp b/test/include/RAJA_test-launch-direct-teams-threads-3D-execpol.hpp
@@ -100,12 +100,12 @@ using Hip_launch_policies = camp::list<hip_direct_policies>;
 using sycl_direct_policies = 
   camp::list<
              RAJA::LaunchPolicy<RAJA::sycl_launch_t<true>>,
-             RAJA::LoopPolicy<RAJA::sycl_group_2_direct>,
+             RAJA::LoopPolicy<RAJA::sycl_group_0_direct>, //slowest
              RAJA::LoopPolicy<RAJA::sycl_group_1_direct>,
-             RAJA::LoopPolicy<RAJA::sycl_group_0_direct>,
-             RAJA::LoopPolicy<RAJA::sycl_local_2_direct>,
+             RAJA::LoopPolicy<RAJA::sycl_group_2_direct>, //fastest
+             RAJA::LoopPolicy<RAJA::sycl_local_0_direct>,
              RAJA::LoopPolicy<RAJA::sycl_local_1_direct>,
-             RAJA::LoopPolicy<RAJA::sycl_local_0_direct>
+             RAJA::LoopPolicy<RAJA::sycl_local_2_direct>
             >;
 
 using Sycl_launch_policies = camp::list<sycl_direct_policies>;

diff --git a/test/include/RAJA_test-launch-execpol.hpp b/test/include/RAJA_test-launch-execpol.hpp
@@ -68,7 +68,7 @@ using Hip_launch_policies = camp::list<
 
 using sycl_policies = camp::list<
   RAJA::LaunchPolicy<RAJA::sycl_launch_t<true>>,
-  RAJA::LoopPolicy<RAJA::sycl_global_item_0>>;
+  RAJA::LoopPolicy<RAJA::sycl_global_item_2>>;
 
 using Sycl_launch_policies = camp::list<
       sycl_policies

diff --git a/test/include/RAJA_test-launch-loop-teams-threads-1D-execpol.hpp b/test/include/RAJA_test-launch-loop-teams-threads-1D-execpol.hpp
@@ -75,8 +75,8 @@ using Hip_launch_policies = camp::list<
 #if defined(RAJA_ENABLE_SYCL)
 using sycl_loop_policies = camp::list<
   RAJA::LaunchPolicy<RAJA::sycl_launch_t<true>>,
-  RAJA::LoopPolicy<RAJA::sycl_group_0_loop>,
-  RAJA::LoopPolicy<RAJA::sycl_local_0_loop>
+  RAJA::LoopPolicy<RAJA::sycl_group_2_loop>,
+  RAJA::LoopPolicy<RAJA::sycl_local_2_loop>
   >;
 
 using Sycl_launch_policies = camp::list<  

diff --git a/test/include/RAJA_test-launch-loop-teams-threads-3D-execpol.hpp b/test/include/RAJA_test-launch-loop-teams-threads-3D-execpol.hpp
@@ -95,12 +95,12 @@ using Hip_launch_policies = camp::list<
 #if defined(RAJA_ENABLE_SYCL)
 using sycl_loop_policies = camp::list<
   RAJA::LaunchPolicy<RAJA::sycl_launch_t<true>>,
-  RAJA::LoopPolicy<RAJA::sycl_group_2_loop>,
+  RAJA::LoopPolicy<RAJA::sycl_group_0_loop>, //slowest index
   RAJA::LoopPolicy<RAJA::sycl_group_1_loop>,
-  RAJA::LoopPolicy<RAJA::sycl_group_0_loop>,
-  RAJA::LoopPolicy<RAJA::sycl_local_2_loop>,
+  RAJA::LoopPolicy<RAJA::sycl_group_2_loop>, //fastest index
+  RAJA::LoopPolicy<RAJA::sycl_local_0_loop>,
   RAJA::LoopPolicy<RAJA::sycl_local_1_loop>,
-  RAJA::LoopPolicy<RAJA::sycl_local_0_loop>
+  RAJA::LoopPolicy<RAJA::sycl_local_2_loop>
   >;
 
 using Sycl_launch_policies = camp::list<  

diff --git a/test/include/RAJA_test-launch-runtime-execpol.hpp b/test/include/RAJA_test-launch-runtime-execpol.hpp
@@ -52,8 +52,8 @@ using Sequential_launch_policies = camp::list<seq_hip_policies>;
 using seq_sycl_policies =
   camp::list<
              RAJA::LaunchPolicy<RAJA::seq_launch_t,RAJA::sycl_launch_t<true>>,
-             RAJA::LoopPolicy<RAJA::seq_exec, RAJA::sycl_group_0_direct>,
-             RAJA::LoopPolicy<RAJA::seq_exec,RAJA::sycl_local_0_loop>
+             RAJA::LoopPolicy<RAJA::seq_exec, RAJA::sycl_group_2_direct>,
+             RAJA::LoopPolicy<RAJA::seq_exec,RAJA::sycl_local_2_loop>
             >;
 
 using Sequential_launch_policies = camp::list<seq_sycl_policies>;
@@ -110,8 +110,8 @@ using OpenMP_launch_policies = camp::list<omp_hip_policies>;
 using omp_sycl_policies =
   camp::list<
              RAJA::LaunchPolicy<RAJA::omp_launch_t,RAJA::sycl_launch_t<false>>,
-             RAJA::LoopPolicy<RAJA::omp_for_exec, RAJA::sycl_group_0_direct>,
-             RAJA::LoopPolicy<RAJA::seq_exec,RAJA::sycl_local_0_loop>
+             RAJA::LoopPolicy<RAJA::omp_for_exec, RAJA::sycl_group_2_direct>,
+             RAJA::LoopPolicy<RAJA::seq_exec,RAJA::sycl_local_2_loop>
             >;
 
 using OpenMP_launch_policies = camp::list<omp_sycl_policies>;