diff --git a/test/rocprim/test_hip_block_exchange.cpp b/test/rocprim/test_hip_block_exchange.cpp index 4314e1f22..eb4ef6e76 100644 --- a/test/rocprim/test_hip_block_exchange.cpp +++ b/test/rocprim/test_hip_block_exchange.cpp @@ -28,8 +28,6 @@ // Google Test #include -// HC API -#include // rocPRIM API #include diff --git a/test/rocprim/test_hip_block_load_store.cpp b/test/rocprim/test_hip_block_load_store.cpp index 6739a79f7..6dd1c5e3b 100644 --- a/test/rocprim/test_hip_block_load_store.cpp +++ b/test/rocprim/test_hip_block_load_store.cpp @@ -26,9 +26,9 @@ // Google Test #include -// HC API -#include -#include +// HIP API +#include +#include // rocPRIM API #include @@ -221,26 +221,26 @@ typedef ::testing::Types< params, params, - params, - params, - params, - params, - params, - params, - - params, - params, - params, - params, - params, - params, - - params, - params, - params, - params, - params, - params + params, + params, + params, + params, + params, + params, + + params, + params, + params, + params, + params, + params, + + params, + params, + params, + params, + params, + params > Params; TYPED_TEST_CASE(RocprimBlockLoadStoreClassTests, ClassParams); diff --git a/test/rocprim/test_hip_transform_iterator.cpp b/test/rocprim/test_hip_transform_iterator.cpp index 55d64a81e..1190f5f1a 100644 --- a/test/rocprim/test_hip_transform_iterator.cpp +++ b/test/rocprim/test_hip_transform_iterator.cpp @@ -39,6 +39,7 @@ template struct times_two { + ROCPRIM_HOST_DEVICE T operator()(const T& value) const { return 2 * value; @@ -48,6 +49,7 @@ struct times_two template struct plus_ten { + ROCPRIM_HOST_DEVICE T operator()(const T& value) const { return value + 10; diff --git a/test/rocprim/test_hip_warp_reduce.cpp b/test/rocprim/test_hip_warp_reduce.cpp index e4fde26d8..0b144d79b 100644 --- a/test/rocprim/test_hip_warp_reduce.cpp +++ b/test/rocprim/test_hip_warp_reduce.cpp @@ -97,7 +97,7 @@ void warp_reduce_sum_kernel(T* device_input, T* device_output) T value = device_input[index]; using wreduce_t = rp::warp_reduce; - tile_static typename wreduce_t::storage_type storage[warps_no]; + __shared__ typename wreduce_t::storage_type storage[warps_no]; wreduce_t().reduce(value, value, storage[warp_id]); if(hipThreadIdx_x%LogicalWarpSize == 0) @@ -204,7 +204,7 @@ void warp_allreduce_sum_kernel(T* device_input, T* device_output) T value = device_input[index]; using wreduce_t = rp::warp_reduce; - tile_static typename wreduce_t::storage_type storage[warps_no]; + __shared__ typename wreduce_t::storage_type storage[warps_no]; wreduce_t().reduce(value, value, storage[warp_id]); device_output[index] = value; @@ -312,7 +312,7 @@ void warp_reduce_sum_kernel(T* device_input, T* device_output, size_t valid) T value = device_input[index]; using wreduce_t = rp::warp_reduce; - tile_static typename wreduce_t::storage_type storage[warps_no]; + __shared__ typename wreduce_t::storage_type storage[warps_no]; wreduce_t().reduce(value, value, valid, storage[warp_id]); if(hipThreadIdx_x%LogicalWarpSize == 0) @@ -420,7 +420,7 @@ void warp_allreduce_sum_kernel(T* device_input, T* device_output, size_t valid) T value = device_input[index]; using wreduce_t = rp::warp_reduce; - tile_static typename wreduce_t::storage_type storage[warps_no]; + __shared__ typename wreduce_t::storage_type storage[warps_no]; wreduce_t().reduce(value, value, valid, storage[warp_id]); device_output[index] = value; @@ -623,7 +623,7 @@ void head_segmented_warp_reduce_kernel(T* input, Flag* flags, T* output) auto flag = flags[index]; using wreduce_t = rp::warp_reduce; - tile_static typename wreduce_t::storage_type storage[warps_no]; + __shared__ typename wreduce_t::storage_type storage[warps_no]; wreduce_t().head_segmented_reduce(value, value, flag, storage[warp_id]); output[index] = value; @@ -750,7 +750,7 @@ void tail_segmented_warp_reduce_kernel(T* input, Flag* flags, T* output) auto flag = flags[index]; using wreduce_t = rp::warp_reduce; - tile_static typename wreduce_t::storage_type storage[warps_no]; + __shared__ typename wreduce_t::storage_type storage[warps_no]; wreduce_t().tail_segmented_reduce(value, value, flag, storage[warp_id]); output[index] = value; diff --git a/test/rocprim/test_hip_warp_sort.cpp b/test/rocprim/test_hip_warp_sort.cpp index 9f896cb92..03484e4f2 100644 --- a/test/rocprim/test_hip_warp_sort.cpp +++ b/test/rocprim/test_hip_warp_sort.cpp @@ -52,7 +52,7 @@ class RocprimWarpSortShuffleBasedTests : public ::testing::Test { }; template -bool test(const T& a, const T& b) [[hc]] +bool test(const T& a, const T& b) { return a < b; }