Add __launch_bounds__
MODIFY
test/test_binary_search.cpp
test/test_binary_search.cpp
@@ -36,7 +36,7 @@ struct custom_less
template <class T>
__global__
-THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
+THRUST_HIP_LAUNCH_BOUNDS_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;
@@ -88,7 +88,7 @@ TYPED_TEST(BinarySearchTestsInKernel, TestLowerBound)
template <class T>
__global__
-THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
+THRUST_HIP_LAUNCH_BOUNDS_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;
@@ -140,7 +140,7 @@ TYPED_TEST(BinarySearchTestsInKernel, TestUpperBound)
template <class T>
__global__
-THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
void binary_search_kernel(size_t n, T* input, bool* output)
{
output[0] = thrust::binary_search(thrust::device, input, input + n, T(0));

MODIFY
test/test_binary_search_vector.cpp
test/test_binary_search_vector.cpp
@@ -37,7 +37,7 @@ struct custom_less
template <class T>
__global__
-THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
void lower_bound_kernel(size_t n, T* input, ptrdiff_t* output)
{
thrust::counting_iterator<T> values(0);
@@ -82,7 +82,7 @@ TYPED_TEST(BinarySearchVectorTestsInKernel, TestLowerBound)
template <class T>
__global__
-THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
void upper_bound_kernel(size_t n, T* input, ptrdiff_t* output)
{
thrust::counting_iterator<T> values(0);
@@ -126,7 +126,7 @@ TYPED_TEST(BinarySearchVectorTestsInKernel, TestUpperBound)
template <class T>
__global__
-THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
void binary_search_kernel(size_t n, T* input, bool* output)
{
thrust::counting_iterator<T> values(0);

MODIFY
test/test_dereference.cpp
test/test_dereference.cpp
template <typename Iterator1, typename Iterator2>
__global__
-THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
void simple_copy_on_device(Iterator1 first1, Iterator1 last1, Iterator2 first2)
{
while(first1 != last1)

MODIFY
test/test_for_each.cpp
test/test_for_each.cpp
@@ -90,7 +90,7 @@ TYPED_TEST(ForEachTests, HostPathSimpleTest)
template <class F>
__global__
-THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
void simple_test_kernel(F func, int size)
{
// (void) func; (void) size;

MODIFY
test/test_sort.cpp
test/test_sort.cpp
@@ -432,7 +432,7 @@ TEST(SortTests, TestSortBoolDescending)
//TODO: refactor this test into a different set of tests
__global__
-THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
+THRUST_HIP_LAUNCH_BOUNDS_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
@@ -243,7 +243,7 @@ TEST(SortByKeyTests, TestSortByKeyBoolDescending)
//TODO: refactor this test into a different set of tests
__global__
-THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
void SortByKeyKernel(int const N, int* keys, short* values)
{
if(threadIdx.x == 0)

MODIFY
testing/dereference.cu
testing/dereference.cu
@@ -13,7 +13,7 @@ __THRUST_DISABLE_MSVC_POSSIBLE_LOSS_OF_DATA_WARNING_BEGIN
template <typename Iterator1, typename Iterator2>
#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA || THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_HIP
__global__
-THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
#endif
void simple_copy_on_device(Iterator1 first1, Iterator1 last1, Iterator2 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
+// TODO: These paremeters should be tuned for NAVI.
+#ifndef THRUST_HIP_DEFAULT_MAX_BLOCK_SIZE
+ #define THRUST_HIP_DEFAULT_MAX_BLOCK_SIZE 256
+#endif
+#ifndef THRUST_HIP_DEFAULT_MIN_WARPS_PER_EU
+ #define THRUST_HIP_DEFAULT_MIN_WARPS_PER_EU 1
+#endif
+#define THRUST_HIP_LAUNCH_BOUNDS(BlockSize) __launch_bounds__(BlockSize, THRUST_HIP_DEFAULT_MIN_WARPS_PER_EU)
+#define THRUST_HIP_LAUNCH_BOUNDS_DEFAULT THRUST_HIP_LAUNCH_BOUNDS(THRUST_HIP_DEFAULT_MAX_BLOCK_SIZE)
+
#ifdef __HIP_DEVICE_COMPILE__
#define THRUST_HIP_DEVICE_CODE
#endif

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

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

MODIFY
thrust/system/hip/detail/set_operations.h
thrust/system/hip/detail/set_operations.h
@@ -784,7 +784,7 @@ namespace __set_operations
class SetOp,
class LookBackScanState>
__global__
- THRUST_HIP_ATTRIBUTE_WORK_GROUP_SIZE_RANGE_DEFAULT
+ THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
void lookback_set_op_kernel(KeysIt1 keys1,
KeysIt2 keys2,
ValuesIt1 values1,

Add __launch_bounds__