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(hipGetDevice(&deviceId)); static_cast(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), @@ -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), @@ -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<