From 550def052611eb514124f25b582e12a6015bf0ee Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Nicol=C3=B2=20Scipione?= <nicolo.scipione@codeplay.com>
Date: Tue, 26 Mar 2024 17:42:55 +0100
Subject: [PATCH] Apply suggestions from code review

Co-authored-by: HJA Bird <hja.bird@gmail.com>
Co-authored-by: pgorlani <92453485+pgorlani@users.noreply.github.com>
---
 src/interface/blas1/backend/amd_gpu.hpp    | 14 ++++++-------
 src/operations/blas1/WGAtomicReduction.hpp | 23 ++++++++--------------
 2 files changed, 15 insertions(+), 22 deletions(-)

diff --git a/src/interface/blas1/backend/amd_gpu.hpp b/src/interface/blas1/backend/amd_gpu.hpp
index e6d8e44ed..3a1409b8d 100644
--- a/src/interface/blas1/backend/amd_gpu.hpp
+++ b/src/interface/blas1/backend/amd_gpu.hpp
@@ -35,13 +35,13 @@ typename sb_handle_t::event_t _asum(
     sb_handle_t& sb_handle, index_t _N, container_0_t _vx, increment_t _incx,
     container_1_t _rs, const typename sb_handle_t::event_t& _dependencies) {
   /**
-   * This compile time check is absolutely necessary for AMD gpu.
-   * AMD atomic operations required a specific combination of hardware that we
-   *cannot check neither enforce to users. Since reduction operators kernel
-   *implementation useses atomic operation without that particular combination
-   *the operator may fail silently. This check enforce a different atomic
-   *address space causing a big performance degradation, but making the kernel
-   *behaves correctly also with managed memory (aka malloc_shared allocation).
+   * This compile time check is absolutely necessary for AMD GPUs.
+   * AMD's atomic operations require a specific combination of hardware that
+   * cannot be checked nor enforced. Since the reduction operator kernel
+   * implementation uses atomic operations, without that particular hardware combination
+   * the reduction may silently fail. This check enforces a different atomic
+   * address space causing a big performance degradation, but also making the kernel
+   * behave correctly with managed memory (aka malloc_shared allocation).
    **/
 #ifdef SB_ENABLE_USM
   bool usm_managed_mem{false};
diff --git a/src/operations/blas1/WGAtomicReduction.hpp b/src/operations/blas1/WGAtomicReduction.hpp
index 21756c7c0..ca46b8269 100644
--- a/src/operations/blas1/WGAtomicReduction.hpp
+++ b/src/operations/blas1/WGAtomicReduction.hpp
@@ -126,21 +126,14 @@ PORTBLAS_INLINE
                                       cl::sycl::plus<value_t>());
   }
   if (ndItem.get_local_id()[0] == 0) {
-    if constexpr (!usmManagedMem) {
-      auto atomic_res =
-          cl::sycl::atomic_ref<value_t, cl::sycl::memory_order::relaxed,
-                               cl::sycl::memory_scope::device,
-                               cl::sycl::access::address_space::global_space>(
-              lhs_.get_data()[0]);
-      atomic_res += val;
-    } else {
-      auto atomic_res =
-          cl::sycl::atomic_ref<value_t, cl::sycl::memory_order::relaxed,
-                               cl::sycl::memory_scope::device,
-                               cl::sycl::access::address_space::generic_space>(
-              lhs_.get_data()[0]);
-      atomic_res += val;
-    }
+    constexpr cl::sycl::access::address_space addr_sp =
+        usmManagedMem ? cl::sycl::access::address_space::generic_space
+                      : cl::sycl::access::address_space::global_space;
+    auto atomic_res =
+        cl::sycl::atomic_ref<value_t, cl::sycl::memory_order::relaxed,
+                             cl::sycl::memory_scope::device, addr_sp>(
+            lhs_.get_data()[0]);
+    atomic_res += val;
   }
 
   return {};