Regressions observed on HIP/CUDA path after integrating CUB-HIP as a sub-module

[sriharikarnam]

Facing compilation errors from CUB while exercising Thrust API's on HIP/CUDA path. The branch currently under use is hip_port_1.7.3

Below is the list of issues and fix which works for Thrust API's, these need to be checked-in to the relevant cub-hip branch.

Issue 1: ../thrust/system/cuda/detail/cub/device/dispatch/dispatch_reduce.cuh(479): error: expected a ">>>"
../thrust/system/cuda/detail/cub/device/dispatch/dispatch_reduce.cuh(479): error: expected a "("
There exists a syntax error in the above mentioned file.

Solution: In file thrust/system/cuda/detail/cub-hip/cub/device/dispatch/dispatch_reduce.cuh ,line no:479
- hipLaunchKernelGGL(single_tile_kernel, 1, 256 0, stream,
+ hipLaunchKernelGGL(single_tile_kernel, 1, 256, 0, stream,
This resolves the issue.

Issue 2 : error: invalid conversion from ‘void ()(const int, int*, int, int, int, thrust::system::cuda::detail::cub_::GridEvenShare)’ to ‘const void*’ [-========= fpermissive]
return hipOccupancyMaxActiveBlocksPerMultiprocessor (

Solution: In file thrust/system/cuda/detail/cub/util_device.cuh, Line no:276
- kernel_ptr,
+ (const void *)kernel_ptr,

This resolves the issue.

Steps to reproduce:
$ git clone https://github.com/ROCmSoftwarePlatform/Thrust.git
$ cd Thrust
$ export HIP_PLATFORM=nvcc (For NVCC Platform )
$ cd examples
$ cp histogram.cu histogram.cpp
$ /opt/rocm/bin/hipcc histogram.cpp -I../ -o histogram.cpp.out

[scchan]

Could you create a pull request for this please? Thanks!

[sriharikarnam]

@scchan
pull request has been submitted for regression issues on HIP/CUDA.
#7

[sriharikarnam]

@scchan
one more PR raised for CUB-HIP on HIP/CUDA path:
#8, please merge it to cub-hip(1.7.3) branch after review