EXSWHTEC-298 - Extend tests for atomic exchange operations #290

Change-Id: I19b4cb6177f3caa74c3f889de107e349ebb1d37c [ROCm/hip-tests commit: 73e8f26cf4]
2023-12-28 17:11:02 +01:00
Parent 5962202338
@@ -40,6 +40,7 @@ set(TEST_SRC
    __hip_atomic_fetch_and.cc
    __hip_atomic_fetch_or.cc
    __hip_atomic_fetch_xor.cc
+    __hip_atomic_exchange.cc
 )

 if(HIP_PLATFORM MATCHES "nvidia")
@@ -0,0 +1,136 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "atomicExch_common.hh"
+
+/**
+ * @addtogroup __hip_atomic_exchange __hip_atomic_exchange
+ * @{
+ * @ingroup AtomicsTest
+ * ________________________
+ * Test cases from other modules:
+ *    - @ref Unit_AtomicBuiltins_Negative_Parameters_RTC
+ */
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Executes a single kernel on a single device wherein all threads will perform an atomic
+ * exchange into a runtime determined memory location. Each thread will exchange its own grid wide
+ * linear index + offset into the memory location, storing the return value into a separate output
+ * array slot corresponding to it. Once complete, the union of output array and exchange memory is
+ * validated to contain all values in the range [0, number_of_threads +
+ * number_of_exchange_memory_slots). Several memory access patterns are tested:
+ *      -# All threads exchange to a single memory location
+ *      -# Each thread exchanges into an array containing warp_size elements, using tid % warp_size
+ *         for indexing
+ *      -# Same as the above, but the exchange elements are spread out by L1 cache line size bytes.
+ *
+ *    - The test is run for:
+ *      - All overloads of atomicExch
+ *      - hipMalloc, hipMallocManaged, hipHostMalloc and hipHostRegister allocated exchange memory
+ *      - Exchange memory located in shared memory
+ *      - WAVEFRONT memory scope
+ * Test source
+ * ------------------------
+ *    - unit/atomics/__hip_atomic_exchange.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit___hip_atomic_exchange_Positive_Wavefront", "", int, unsigned int,
+                   unsigned long, unsigned long long, float, double) {
+  int warp_size = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
+  const auto cache_line_size = 128u;
+
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Same address " << current) {
+      AtomicExchSingleDeviceSingleKernelTest<TestType, AtomicScopes::builtin,
+                                             __HIP_MEMORY_SCOPE_WAVEFRONT>(1, sizeof(TestType));
+    }
+
+    DYNAMIC_SECTION("Adjacent addresses " << current) {
+      AtomicExchSingleDeviceSingleKernelTest<TestType, AtomicScopes::builtin,
+                                             __HIP_MEMORY_SCOPE_WAVEFRONT>(warp_size,
+                                                                           sizeof(TestType));
+    }
+
+    DYNAMIC_SECTION("Scattered addresses " << current) {
+      AtomicExchSingleDeviceSingleKernelTest<TestType, AtomicScopes::builtin,
+                                             __HIP_MEMORY_SCOPE_WAVEFRONT>(warp_size,
+                                                                           cache_line_size);
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Executes a single kernel on a single device wherein all threads will perform an atomic
+ * exchange into a runtime determined memory location. Each thread will exchange its own grid wide
+ * linear index + offset into the memory location, storing the return value into a separate output
+ * array slot corresponding to it. Once complete, the union of output array and exchange memory is
+ * validated to contain all values in the range [0, number_of_threads +
+ * number_of_exchange_memory_slots). Several memory access patterns are tested:
+ *      -# All threads exchange to a single memory location
+ *      -# Each thread exchanges into an array containing warp_size elements, using tid % warp_size
+ *         for indexing
+ *      -# Same as the above, but the exchange elements are spread out by L1 cache line size bytes.
+ *
+ *    - The test is run for:
+ *      - All overloads of atomicExch
+ *      - hipMalloc, hipMallocManaged, hipHostMalloc and hipHostRegister allocated exchange memory
+ *      - Exchange memory located in shared memory
+ *      - WORKGROUP memory scope
+ * Test source
+ * ------------------------
+ *    - unit/atomics/__hip_atomic_exchange.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit___hip_atomic_exchange_Positive_Workgroup", "", int, unsigned int,
+                   unsigned long, unsigned long long, float, double) {
+  int warp_size = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
+  const auto cache_line_size = 128u;
+
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Same address " << current) {
+      AtomicExchSingleDeviceSingleKernelTest<TestType, AtomicScopes::builtin,
+                                             __HIP_MEMORY_SCOPE_WORKGROUP>(1, sizeof(TestType));
+    }
+
+    DYNAMIC_SECTION("Adjacent addresses " << current) {
+      AtomicExchSingleDeviceSingleKernelTest<TestType, AtomicScopes::builtin,
+                                             __HIP_MEMORY_SCOPE_WORKGROUP>(warp_size,
+                                                                           sizeof(TestType));
+    }
+
+    DYNAMIC_SECTION("Scattered addresses " << current) {
+      AtomicExchSingleDeviceSingleKernelTest<TestType, AtomicScopes::builtin,
+                                             __HIP_MEMORY_SCOPE_WORKGROUP>(warp_size,
+                                                                           cache_line_size);
+    }
+  }
+}
@@ -22,24 +22,26 @@ THE SOFTWARE.

 #pragma once

-#include <numeric>
-
+#include <cmd_options.hh>
 #include <hip_test_common.hh>
 #include <resource_guards.hh>
 #include <hip/hip_cooperative_groups.h>
-#include <cmd_options.hh>

-enum class AtomicScopes { device, system };
+enum class AtomicScopes { device, system, builtin };

-template <typename T, AtomicScopes scope> __device__ T perform_atomic_exch(T* address, T val) {
+template <typename T, AtomicScopes scope, int memory_scope = __HIP_MEMORY_SCOPE_AGENT>
+__device__ T perform_atomic_exch(T* address, T val) {
  if constexpr (scope == AtomicScopes::device) {
    return atomicExch(address, val);
  } else if (scope == AtomicScopes::system) {
    return atomicExch_system(address, val);
+  } else if (scope == AtomicScopes::builtin) {
+    return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, memory_scope);
  }
 }

-template <typename T, bool use_shared_mem, AtomicScopes scope>
+template <typename T, bool use_shared_mem, AtomicScopes scope,
+          int memory_scope = __HIP_MEMORY_SCOPE_AGENT>
 __global__ void atomic_exch_kernel_compile_time(T* const global_mem, T* const old_vals) {
  __shared__ T shared_mem;

@@ -52,7 +54,7 @@ __global__ void atomic_exch_kernel_compile_time(T* const global_mem, T* const ol
    __syncthreads();
  }

-  old_vals[tid] = perform_atomic_exch<T, scope>(mem, static_cast<T>(tid + 1));
+  old_vals[tid] = perform_atomic_exch<T, scope, memory_scope>(mem, static_cast<T>(tid + 1));

  if constexpr (use_shared_mem) {
    __syncthreads();
@@ -67,7 +69,16 @@ __host__ __device__ T* pitched_offset(T* const ptr, const unsigned int pitch,
  return reinterpret_cast<T*>(byte_ptr + idx * pitch);
 }

-template <typename T, bool use_shared_mem, AtomicScopes scope>
+__device__ void generate_memory_traffic(uint8_t* const begin_addr, uint8_t* const end_addr) {
+  for (volatile uint8_t* addr = begin_addr; addr != end_addr; ++addr) {
+    uint8_t val = *addr;
+    val ^= 0xAB;
+    *addr = val;
+  }
+}
+
+template <typename T, bool use_shared_mem, AtomicScopes scope,
+          int memory_scope = __HIP_MEMORY_SCOPE_AGENT>
 __global__ void atomic_exch_kernel(T* const global_mem, T* const old_vals, const unsigned int width,
                                   const unsigned pitch, const T base_val = 0) {
  extern __shared__ uint8_t shared_mem[];
@@ -84,8 +95,18 @@ __global__ void atomic_exch_kernel(T* const global_mem, T* const old_vals, const
    __syncthreads();
  }

-  old_vals[tid] = perform_atomic_exch<T, scope>(pitched_offset(mem, pitch, tid % width),
-                                                base_val + static_cast<T>(tid + width));
+  const auto n = cooperative_groups::this_grid().size() - width;
+
+  T* atomic_addr = pitched_offset(mem, pitch, tid % width);
+
+  if (tid < n) {
+    old_vals[tid] = perform_atomic_exch<T, scope, memory_scope>(
+        pitched_offset(mem, pitch, tid % width), base_val + static_cast<T>(tid + width));
+  } else {
+    uint8_t* const begin_addr = reinterpret_cast<uint8_t*>(atomic_addr + 1);
+    uint8_t* const end_addr = reinterpret_cast<uint8_t*>(atomic_addr) + pitch;
+    generate_memory_traffic(begin_addr, end_addr);
+  }

  if constexpr (use_shared_mem) {
    __syncthreads();
@@ -255,14 +276,16 @@ class AtomicExchCRTP {
  }
 };

-template <typename T, bool use_shared_mem, AtomicScopes scope>
+template <typename T, bool use_shared_mem, AtomicScopes scope,
+          int memory_scope = __HIP_MEMORY_SCOPE_AGENT>
 class AtomicExch
    : public AtomicExchCRTP<AtomicExch<T, use_shared_mem, scope>, T, use_shared_mem, scope> {
 public:
  void LaunchKernel(const unsigned int shared_mem_size, const hipStream_t stream, T* const mem,
                    T* const old_vals, const T base_val, const AtomicExchParams& p) const {
-    atomic_exch_kernel<T, use_shared_mem, scope><<<p.blocks, p.threads, shared_mem_size, stream>>>(
-        mem, old_vals, p.width, p.pitch, base_val);
+    atomic_exch_kernel<T, use_shared_mem, scope, memory_scope>
+        <<<p.blocks, p.threads, shared_mem_size, stream>>>(mem, old_vals, p.width, p.pitch,
+                                                           base_val);
  }

  void ValidateResults(std::vector<T>& old_vals) const {
@@ -281,23 +304,39 @@ inline dim3 GenerateAtomicExchBlockDimensions() {
  return GENERATE_COPY(dim3(sm_count), dim3(sm_count + sm_count / 2));
 }

-template <typename TestType, AtomicScopes scope>
+template <typename TestType, AtomicScopes scope, int memory_scope = __HIP_MEMORY_SCOPE_AGENT>
 void AtomicExchSingleDeviceSingleKernelTest(const unsigned int width, const unsigned int pitch) {
  AtomicExchParams params;
  params.num_devices = 1;
  params.kernel_count = 1;
-  params.threads = GenerateAtomicExchThreadDimensions();
+  if constexpr (scope == AtomicScopes::builtin && memory_scope == __HIP_MEMORY_SCOPE_SINGLETHREAD) {
+    params.threads = 1;
+  } else if constexpr (scope == AtomicScopes::builtin &&
+                       memory_scope == __HIP_MEMORY_SCOPE_WAVEFRONT) {
+    int warp_size = 0;
+    HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
+    params.threads = dim3(warp_size);
+  } else {
+    params.threads = GenerateAtomicExchThreadDimensions();
+  }
  params.width = width;
  params.pitch = pitch;

  SECTION("Global memory") {
-    params.blocks = GenerateAtomicExchBlockDimensions();
+    if constexpr (scope == AtomicScopes::builtin &&
+                  (memory_scope == __HIP_MEMORY_SCOPE_SINGLETHREAD ||
+                   memory_scope == __HIP_MEMORY_SCOPE_WAVEFRONT ||
+                   memory_scope == __HIP_MEMORY_SCOPE_WORKGROUP)) {
+      params.blocks = dim3(1);
+    } else {
+      params.blocks = GenerateAtomicExchBlockDimensions();
+    }
    using LA = LinearAllocs;
    for (const auto alloc_type :
         {LA::hipMalloc, LA::hipHostMalloc, LA::hipMallocManaged, LA::mallocAndRegister}) {
      params.alloc_type = alloc_type;
      DYNAMIC_SECTION("Allocation type: " << to_string(alloc_type)) {
-        AtomicExch<TestType, false, scope>().run(params);
+        AtomicExch<TestType, false, scope, memory_scope>().run(params);
      }
    }
  }
@@ -305,7 +344,7 @@ void AtomicExchSingleDeviceSingleKernelTest(const unsigned int width, const unsi
  SECTION("Shared memory") {
    params.blocks = dim3(1);
    params.alloc_type = LinearAllocs::hipMalloc;
-    AtomicExch<TestType, true, scope>().run(params);
+    AtomicExch<TestType, true, scope, memory_scope>().run(params);
  }
 }