Add amdgpu_flat_work_group_size
MODIFY
test/test_binary_search.cpp
test/test_binary_search.cpp
@@ -35,7 +35,9 @@ struct custom_less
};
template <class T>
-__global__ void lower_bound_kernel(size_t n, T* input, ptrdiff_t* output)
+__global__
+THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
+void lower_bound_kernel(size_t n, T* input, ptrdiff_t* output)
{
output[0] = thrust::lower_bound(thrust::device, input, input + n, T(0), custom_less()) - input;
output[1] = thrust::lower_bound(thrust::device, input, input + n, T(1)) - input;
@@ -85,7 +87,9 @@ TYPED_TEST(BinarySearchTestsInKernel, TestLowerBound)
}
template <class T>
-__global__ void upper_bound_kernel(size_t n, T* input, ptrdiff_t* output)
+__global__
+THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
+void upper_bound_kernel(size_t n, T* input, ptrdiff_t* output)
{
output[0] = thrust::upper_bound(thrust::device, input, input + n, T(0)) - input;
output[1] = thrust::upper_bound(thrust::device, input, input + n, T(1)) - input;
@@ -135,7 +139,9 @@ TYPED_TEST(BinarySearchTestsInKernel, TestUpperBound)
}
template <class T>
-__global__ void binary_search_kernel(size_t n, T* input, bool* output)
+__global__
+THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
+void binary_search_kernel(size_t n, T* input, bool* output)
{
output[0] = thrust::binary_search(thrust::device, input, input + n, T(0));
output[1] = thrust::binary_search(thrust::device, input, input + n, T(1));

MODIFY
test/test_binary_search_vector.cpp
test/test_binary_search_vector.cpp
@@ -36,7 +36,9 @@ struct custom_less
};
template <class T>
-__global__ void lower_bound_kernel(size_t n, T* input, ptrdiff_t* output)
+__global__
+THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
+void lower_bound_kernel(size_t n, T* input, ptrdiff_t* output)
{
thrust::counting_iterator<T> values(0);
thrust::lower_bound(
@@ -79,7 +81,9 @@ TYPED_TEST(BinarySearchVectorTestsInKernel, TestLowerBound)
}
template <class T>
-__global__ void upper_bound_kernel(size_t n, T* input, ptrdiff_t* output)
+__global__
+THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
+void upper_bound_kernel(size_t n, T* input, ptrdiff_t* output)
{
thrust::counting_iterator<T> values(0);
thrust::upper_bound(thrust::device, input, input + n, values, values + 10, output);
@@ -121,7 +125,9 @@ TYPED_TEST(BinarySearchVectorTestsInKernel, TestUpperBound)
}
template <class T>
-__global__ void binary_search_kernel(size_t n, T* input, bool* output)
+__global__
+THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
+void binary_search_kernel(size_t n, T* input, bool* output)
{
thrust::counting_iterator<T> values(0);
thrust::binary_search(thrust::device, input, input + n, values, values + 10, output);

MODIFY
test/test_dereference.cpp
test/test_dereference.cpp
#include "test_header.hpp"
template <typename Iterator1, typename Iterator2>
-__global__ void simple_copy_on_device(Iterator1 first1, Iterator1 last1, Iterator2 first2)
+__global__
+THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
+void simple_copy_on_device(Iterator1 first1, Iterator1 last1, Iterator2 first2)
{
while(first1 != last1)
*(first2++) = *(first1++);

MODIFY
test/test_for_each.cpp
test/test_for_each.cpp
@@ -89,7 +89,9 @@ TYPED_TEST(ForEachTests, HostPathSimpleTest)
}
template <class F>
-__global__ void simple_test_kernel(F func, int size)
+__global__
+THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
+void simple_test_kernel(F func, int size)
{
// (void) func; (void) size;
thrust::for_each(thrust::seq,

MODIFY
test/test_sort.cpp
test/test_sort.cpp
@@ -431,7 +431,9 @@ TEST(SortTests, TestSortBoolDescending)
}
//TODO: refactor this test into a different set of tests
-__global__ void SortKernel(int const N, int* array)
+__global__
+THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
+void SortKernel(int const N, int* array)
{
if(threadIdx.x == 0)
{

MODIFY
test/test_sort_by_key.cpp
test/test_sort_by_key.cpp
@@ -242,7 +242,9 @@ TEST(SortByKeyTests, TestSortByKeyBoolDescending)
}
//TODO: refactor this test into a different set of tests
-__global__ void SortByKeyKernel(int const N, int* keys, short* values)
+__global__
+THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
+void SortByKeyKernel(int const N, int* keys, short* values)
{
if(threadIdx.x == 0)
{

MODIFY
testing/dereference.cu
testing/dereference.cu
@@ -11,8 +11,9 @@ __THRUST_DISABLE_MSVC_POSSIBLE_LOSS_OF_DATA_WARNING_BEGIN
template <typename Iterator1, typename Iterator2>
-#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
+#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA || THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_HIP
__global__
+THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
#endif
void simple_copy_on_device(Iterator1 first1, Iterator1 last1, Iterator2 first2)
{
@@ -23,7 +24,7 @@ void simple_copy_on_device(Iterator1 first1, Iterator1 last1, Iterator2 first2)
template <typename Iterator1, typename Iterator2>
void simple_copy(Iterator1 first1, Iterator1 last1, Iterator2 first2)
{
-#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
+#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA || THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_HIP
simple_copy_on_device<<<1,1>>>(first1, last1, first2);
#else
simple_copy_on_device(first1, last1, first2);

MODIFY
thrust/system/hip/config.h
thrust/system/hip/config.h
#define __THRUST_HAS_HIPRT__ 0
#define THRUST_HIP_RUNTIME_FUNCTION __host__ __forceinline__
#endif
+ #define THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE(BlockSize) \
+ __attribute__((amdgpu_flat_work_group_size(BlockSize,BlockSize)))
+ #define THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE(MinBlockSize,MaxBlockSize) \
+ __attribute__((amdgpu_flat_work_group_size(MinBlockSize,MaxBlockSize)))
+ #define THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT \
+ THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE(128,256)
#else
#define __THRUST_HAS_HIPRT__ 0
#define THRUST_HIP_RUNTIME_FUNCTION __host__ __forceinline__
+ #define THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE(BlockSize)
+ #define THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE(MinBlockSize,MaxBlockSize)
+ #define THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
#endif
#ifdef __HIP_DEVICE_COMPILE__

MODIFY
thrust/system/hip/detail/adjacent_difference.h
thrust/system/hip/detail/adjacent_difference.h
@@ -102,7 +102,9 @@ namespace __adjacent_difference
class InputIt,
class HeadsIt,
class Size>
- __global__ void block_heads_fill(InputIt input, HeadsIt block_heads, Size input_size)
+ __global__
+ THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE(BlockSize)
+ void block_heads_fill(InputIt input, HeadsIt block_heads, Size input_size)
{
constexpr auto items_per_block = BlockSize * ItemsPerThread;
Size tile_base = blockIdx.x * items_per_block;
@@ -140,7 +142,9 @@ namespace __adjacent_difference
class HeadsIt,
class OutputIt,
class BinaryOp>
- __global__ void adjacent_difference_kernel(InputIt input,
+ __global__
+ THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE(BlockSize)
+ void adjacent_difference_kernel(InputIt input,
HeadsIt block_heads,
const size_t input_size,
OutputIt output,

MODIFY
thrust/system/hip/detail/parallel_for.h
thrust/system/hip/detail/parallel_for.h
@@ -47,7 +47,9 @@ namespace __parallel_for
};
template <unsigned int BlockSize, unsigned int ItemsPerThread, class F, class Size>
- __global__ void kernel(F f, Size num_items)
+ __global__
+ THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE(BlockSize)
+ void kernel(F f, Size num_items)
{
constexpr auto items_per_block = BlockSize * ItemsPerThread;
Size tile_base = blockIdx.x * items_per_block;

MODIFY
thrust/system/hip/detail/set_operations.h
thrust/system/hip/detail/set_operations.h
@@ -783,8 +783,9 @@ namespace __set_operations
class CompareOp,
class SetOp,
class LookBackScanState>
- __global__ void
- lookback_set_op_kernel(KeysIt1 keys1,
+ __global__
+ THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
+ void lookback_set_op_kernel(KeysIt1 keys1,
KeysIt2 keys2,
ValuesIt1 values1,
ValuesIt2 values2,

Add amdgpu_flat_work_group_size