Add additional check for gfx908 asic revision (#171)

ROCm · Jul 8, 2020 · 33ac26c · 33ac26c
1 parent f2d6ae0
commit 33ac26c
Show file tree

Hide file tree

Showing 2 changed files with 16 additions and 4 deletions.
diff --git 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
@@ -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<