From 33ac26c9b8375fa028bc891d4e18250f0cc76cb2 Mon Sep 17 00:00:00 2001
From: Eiden Yoshida <47196116+eidenyoshida@users.noreply.github.com>
Date: Wed, 8 Jul 2020 15:52:41 -0600
Subject: [PATCH] Add additional check for gfx908 asic revision (#171)

---
 rocprim/include/rocprim/device/device_partition.hpp | 10 ++++++++--
 rocprim/include/rocprim/device/device_scan.hpp      | 10 ++++++++--
 2 files changed, 16 insertions(+), 4 deletions(-)

diff --git a/rocprim/include/rocprim/device/device_partition.hpp b/rocprim/include/rocprim/device/device_partition.hpp
index fda8bf1c2..9e79c40ff 100644
--- a/rocprim/include/rocprim/device/device_partition.hpp
+++ b/rocprim/include/rocprim/device/device_partition.hpp
@@ -199,7 +199,13 @@ hipError_t partition_impl(void * temporary_storage,
     static_cast<void>(hipGetDevice(&deviceId));
     static_cast<void>(hipGetDeviceProperties(&prop, deviceId));
 
-    if (prop.gcnArch == 908)
+#if HIP_VERSION >= 307
+    int asicRevision = prop.asicRevision;
+#else
+    int asicRevision = 0;
+#endif
+
+    if (prop.gcnArch == 908 && asicRevision < 2)
     {
         hipLaunchKernelGGL(
             HIP_KERNEL_NAME(init_offset_scan_state_kernel<offset_scan_state_with_sleep_type>),
@@ -220,7 +226,7 @@ hipError_t partition_impl(void * temporary_storage,
 
     if(debug_synchronous) start = std::chrono::high_resolution_clock::now();
     grid_size = number_of_blocks;
-    if (prop.gcnArch == 908)
+    if (prop.gcnArch == 908 && asicRevision < 2)
     {
         hipLaunchKernelGGL(
             HIP_KERNEL_NAME(partition_kernel<
diff --git a/rocprim/include/rocprim/device/device_scan.hpp b/rocprim/include/rocprim/device/device_scan.hpp
index a6ef93a7d..8ebe07be2 100644
--- a/rocprim/include/rocprim/device/device_scan.hpp
+++ b/rocprim/include/rocprim/device/device_scan.hpp
@@ -380,7 +380,13 @@ auto scan_impl(void * temporary_storage,
         if(debug_synchronous) start = std::chrono::high_resolution_clock::now();
         auto grid_size = (number_of_blocks + block_size - 1)/block_size;
 
-        if (prop.gcnArch == 908)
+#if HIP_VERSION >= 307
+        int asicRevision = prop.asicRevision;
+#else
+        int asicRevision = 0;
+#endif
+
+        if (prop.gcnArch == 908 && asicRevision < 2)
         {
             hipLaunchKernelGGL(
                 HIP_KERNEL_NAME(init_lookback_scan_state_kernel<scan_state_with_sleep_type>),
@@ -399,7 +405,7 @@ auto scan_impl(void * temporary_storage,
 
         if(debug_synchronous) start = std::chrono::high_resolution_clock::now();
         grid_size = number_of_blocks;
-        if (prop.gcnArch == 908)
+        if (prop.gcnArch == 908 && asicRevision < 2)
         {
             hipLaunchKernelGGL(
                 HIP_KERNEL_NAME(lookback_scan_kernel<