Building issue for rocPRIM exclusive_scan

[sriharikarnam]

Issue
Build issue in mxnet related to exclusive_scan of rocPRIM when integrating rocPRIM in place of cub-hip
I am running into this issue. I needed some help

/opt/rocm/hipcub/include/hipcub/rocprim/block/block_scan.hpp:193:20: error: no matching member function for call to 'exclusive_scan'
base_type::exclusive_scan(input, output, temp_storage_);
~~~~~~~~~~~^~~~~~~~~~~~~~
src/operator/tensor/./cast_storage-inl.cuh:445:33: note: in instantiation of member function 'hipcub::BlockScan<long, 256, hipcub::BLOCK_SCAN_RAKING, 1, 1, 1>::ExclusiveSum' requested here
BlockScan(temp_storage).ExclusiveSum(nnz, nnz);

The error was no matching member function call in the block_scan.hpp for exclusive_scan in https://github.com/ROCmSoftwarePlatform/rocPRIM/blob/master/hipcub/include/hipcub/rocprim/block/block_scan.hpp#L193

[ex-rzr]

This issue is fixed here: c11dda7

But #33 is not merged yet.

[sriharikarnam]

@ex-rzr
Thanks for the inputs.
With the given inputs the issue solved for exclusive_scan.
But i see another issue as below
/opt/rocm/hipcub/include/hipcub/cub/device/device_radix_sort.hpp(61): error: DoubleBuffer is not a template

/opt/rocm/hipcub/include/hipcub/cub/device/device_radix_sort.hpp(62): error: DoubleBuffer is not a template

/opt/rocm/hipcub/include/hipcub/cub/device/device_radix_sort.hpp(109): error: DoubleBuffer is not a template

/opt/rocm/hipcub/include/hipcub/cub/device/device_radix_sort.hpp(110): error: DoubleBuffer is not a template

/opt/rocm/hipcub/include/hipcub/cub/device/device_radix_sort.hpp(153): error: DoubleBuffer is not a template

/opt/rocm/hipcub/include/hipcub/cub/device/device_radix_sort.hpp(196): error: DoubleBuffer is not a template

[ex-rzr]

This is very strange.

I've checked hipCUB tests on NVCC: all tests can be built without any errors or warnings.

How are you using hipCUB? Can you show the code?

[sriharikarnam]

I've tried to integrate rocprim in place of cub-hip in mxnet source as bellow for headers and namespace

-#include <cub/cub.cuh>
+#include <hipcub/hipcub.hpp>

-#include <cub/device/device_run_length_encode.cuh>
+#include <cub/device/device_run_length_encode.hpp>

-#include <cub/device/device_scan.cuh>
+#include <cub/device/device_scan.hpp>

and replaced the cub namespace with hipcub

• cub::DeviceSelect::Unique(NULL, unique_temp_bytes, null_dptr, null_dptr,

• hipcub::DeviceSelect::Unique(NULL, unique_temp_bytes, null_dptr, null_dptr,

and also added the CUB directory for NVCC compilation.

The code works for hcc compilation ,but for nvcc the issue is as below /opt/rocm/hipcub/include/hipcub/cub/device/device_radix_sort.hpp(61): error: DoubleBuffer is not a template

/opt/rocm/hipcub/include/hipcub/cub/device/device_radix_sort.hpp(62): error: DoubleBuffer is not a template

/opt/rocm/hipcub/include/hipcub/cub/device/device_radix_sort.hpp(109): error: DoubleBuffer is not a template

/opt/rocm/hipcub/include/hipcub/cub/device/device_radix_sort.hpp(110): error: DoubleBuffer is not a template

/opt/rocm/hipcub/include/hipcub/cub/device/device_radix_sort.hpp(153): error: DoubleBuffer is not a template

/opt/rocm/hipcub/include/hipcub/cub/device/device_radix_sort.hpp(196): error: DoubleBuffer is not a template

[ex-rzr]

Only #include <hipcub/hipcub.hpp> is required here. In the original code 2 other includes do nothing because cub/cub.cuh already includes them, but in the new code this may cause some conflicts.

So my suggestion is to remove all CUB includes and keep only hipcub/hipcub.hpp.

[sriharikarnam]

The suggestion does fix this issue.

thank you for the quick help!

[VincentSC]

#33 is merged in develop