Initial implementation of gfx942 (kokkos#6358)

* Initial implementation of gfx942 Change-Id: Id31ca3ba5356d021cade2abc3e3f51f9f3b4d211 * remove VEGA arch Change-Id: I1454bb0b91518bfcf7a04506e40b98387cdf8ed9 * apply formatting Change-Id: Id9c03fe451d1d28a3c23a77f161a2600f016c7e4 * Fix conditional Co-authored-by: Daniel Arndt <arndtd@ornl.gov> * More cmake fixes Co-authored-by: Damien L-G <dalg24+github@gmail.com> * remove unneeded for old naming schema Change-Id: Ibd028fddeedf8e0fdda50b72625ab62cee6fa71e --------- Co-authored-by: Nicholas Curtis <nicurtis@amd.com> Co-authored-by: Daniel Arndt <arndtd@ornl.gov> Co-authored-by: Damien L-G <dalg24+github@gmail.com>
nliber · Aug 16, 2023 · 04d5c55 · 04d5c55
1 parent 8d8b24a
commit 04d5c55
Show file tree

Hide file tree

Showing 8 changed files with 20 additions and 4 deletions.
diff --git a/Makefile.kokkos b/Makefile.kokkos
@@ -13,7 +13,7 @@ KOKKOS_DEVICES ?= "Threads"
 # NVIDIA:   Kepler,Kepler30,Kepler32,Kepler35,Kepler37,Maxwell,Maxwell50,Maxwell52,Maxwell53,Pascal60,Pascal61,Volta70,Volta72,Turing75,Ampere80,Ampere86,Ada89,Hopper90
 # ARM:      ARMv80,ARMv81,ARMv8-ThunderX,ARMv8-TX2,A64FX
 # IBM:      BGQ,Power7,Power8,Power9
-# AMD-GPUS: GFX906,GFX908,GFX90A,GFX1030, GFX1100
+# AMD-GPUS: GFX906,GFX908,GFX90A,GFX942,GFX1030,GFX1100
 # AMD-CPUS: AMDAVX,Zen,Zen2,Zen3
 # Intel-GPUs: Gen9,Gen11,Gen12LP,DG1,XeHP,PVC
 KOKKOS_ARCH ?= ""
@@ -1091,6 +1091,11 @@ ifeq ($(KOKKOS_INTERNAL_USE_ARCH_AMD_GFX90A), 1)
   tmp := $(call kokkos_append_header,"$H""define KOKKOS_ARCH_AMD_GPU")
   KOKKOS_INTERNAL_HIP_ARCH_FLAG := --offload-arch=gfx90a
 endif
+ifeq ($(KOKKOS_INTERNAL_USE_ARCH_AMD_GFX942), 1)
+  tmp := $(call kokkos_append_header,"$H""define KOKKOS_ARCH_AMD_GFX942")
+  tmp := $(call kokkos_append_header,"$H""define KOKKOS_ARCH_AMD_GPU")
+  KOKKOS_INTERNAL_HIP_ARCH_FLAG := --offload-arch=gfx942
+endif
 ifeq ($(KOKKOS_INTERNAL_USE_ARCH_AMD_GFX1030), 1)
   tmp := $(call kokkos_append_header,"$H""define KOKKOS_ARCH_AMD_GFX1030")
   tmp := $(call kokkos_append_header,"$H""define KOKKOS_ARCH_AMD_GPU")

diff --git a/cmake/KokkosCore_config.h.in b/cmake/KokkosCore_config.h.in
@@ -113,6 +113,7 @@
 #cmakedefine KOKKOS_ARCH_AMD_GFX906
 #cmakedefine KOKKOS_ARCH_AMD_GFX908
 #cmakedefine KOKKOS_ARCH_AMD_GFX90A
+#cmakedefine KOKKOS_ARCH_AMD_GFX942
 #cmakedefine KOKKOS_ARCH_AMD_GFX1030
 #cmakedefine KOKKOS_ARCH_AMD_GFX1100
 #cmakedefine KOKKOS_ARCH_AMD_GPU

diff --git a/cmake/kokkos_arch.cmake b/cmake/kokkos_arch.cmake
@@ -94,6 +94,9 @@ IF(Kokkos_ENABLE_HIP OR Kokkos_ENABLE_OPENMPTARGET OR Kokkos_ENABLE_OPENACC OR K
 ENDIF()
 
 # AMD archs ordered in decreasing priority of autodetection
+LIST(APPEND SUPPORTED_AMD_GPUS       MI300)
+LIST(APPEND SUPPORTED_AMD_ARCHS      AMD_GFX942)
+LIST(APPEND CORRESPONDING_AMD_FLAGS  gfx942)
 LIST(APPEND SUPPORTED_AMD_GPUS       MI200    MI200       MI100    MI100)
 LIST(APPEND SUPPORTED_AMD_ARCHS      VEGA90A  AMD_GFX90A  VEGA908  AMD_GFX908)
 LIST(APPEND CORRESPONDING_AMD_FLAGS  gfx90a   gfx90a      gfx908   gfx908)

diff --git a/core/src/HIP/Kokkos_HIP_Instance.hpp b/core/src/HIP/Kokkos_HIP_Instance.hpp
@@ -30,7 +30,7 @@ namespace Impl {
 
 struct HIPTraits {
 #if defined(KOKKOS_ARCH_AMD_GFX906) || defined(KOKKOS_ARCH_AMD_GFX908) || \
-    defined(KOKKOS_ARCH_AMD_GFX90A)
+    defined(KOKKOS_ARCH_AMD_GFX90A) || defined(KOKKOS_ARCH_AMD_GFX942)
   static int constexpr WarpSize       = 64;
   static int constexpr WarpIndexMask  = 0x003f; /* hexadecimal for 63 */
   static int constexpr WarpIndexShift = 6;      /* WarpSize == 1 << WarpShift*/

diff --git a/core/src/HIP/Kokkos_HIP_ReduceScan.hpp b/core/src/HIP/Kokkos_HIP_ReduceScan.hpp
@@ -90,6 +90,7 @@ struct HIPReductionsFunctor<FunctorType, true> {
       }
       scalar_intra_warp_reduction(functor, value, false, warp_size,
                                   *my_global_team_buffer_element);
+      __threadfence();
     }
   }
 
@@ -182,7 +183,10 @@ struct HIPReductionsFunctor<FunctorType, false> {
       scalar_intra_warp_reduction(
           functor, my_shared_team_buffer_element, false,
           blockDim.x * blockDim.y / HIPTraits::WarpSize);
-      if (threadIdx.x + threadIdx.y == 0) *result = *shared_team_buffer_element;
+      if (threadIdx.x + threadIdx.y == 0) {
+        *result = *shared_team_buffer_element;
+        if (skip) __threadfence();
+      }
     }
   }
 
@@ -382,6 +386,7 @@ __device__ bool hip_single_inter_block_reduce_scan_impl(
     for (size_t i = threadIdx.y; i < word_count.value; i += blockDim.y) {
       global[i] = shared[i];
     }
+    __threadfence();
   }
 
   // Contributing blocks note that their contribution has been completed via an

diff --git a/core/src/HIP/Kokkos_HIP_Shuffle_Reduce.hpp b/core/src/HIP/Kokkos_HIP_Shuffle_Reduce.hpp
@@ -118,6 +118,7 @@ __device__ inline bool hip_inter_block_shuffle_reduction(
     pointer_type global =
         reinterpret_cast<pointer_type>(m_scratch_space) + blockIdx.x;
     *global = value;
+    __threadfence();
   }
 
   // One warp of last block performs inter block reduction through loading the

diff --git a/core/unit_test/TestMathematicalFunctions.hpp b/core/unit_test/TestMathematicalFunctions.hpp
@@ -795,7 +795,7 @@ TEST(TEST_CATEGORY, mathematical_functions_exponential_functions) {
 // FIXME_OPENMPTARGET FIXME_AMD
 #if defined(KOKKOS_ENABLE_OPENMPTARGET) &&                                 \
     (defined(KOKKOS_ARCH_AMD_GFX906) || defined(KOKKOS_ARCH_AMD_GFX908) || \
-     defined(KOKKOS_ARCH_AMD_GFX90A))
+     defined(KOKKOS_ARCH_AMD_GFX90A) || defined(KOKKOS_ARCH_AMD_GFX942))
 
   TEST_MATH_FUNCTION(log2)({1, 23, 456, 7890});
 #endif

diff --git a/generate_makefile.bash b/generate_makefile.bash
@@ -160,6 +160,7 @@ display_help_text() {
       echo "                 AMD_GFX906      = AMD GPU MI50/MI60 GFX906"
       echo "                 AMD_GFX908      = AMD GPU MI100 GFX908"
       echo "                 AMD_GFX90A      = AMD GPU MI200 GFX90A"
+      echo "                 AMD_GFX942      = AMD GPU MI300 GFX942"
       echo "                 AMD_GFX1030     = AMD GPU V620/W6800 GFX1030"
       echo "                 AMD_GFX1100     = AMD GPU RX 7900 XT(X) GFX1100"
       echo "               [ARM]"