Arguments out of order in built-in rocmPRIM DeviceScan::ExclusiveScan

[skyreflectedinmirrors]

edit: removed my previous description, as I've figure out it was a simple issue of the arguments being out of order:

The current call of:

        return ExclusiveScan(
            d_temp_storage, temp_storage_bytes,
            d_in, d_out, T(0), num_items, ::hipcub::Sum(),
            stream, debug_synchronous
        );

in hipcub/rocprim/device/device_scan.hpp is incorrect, as template method with the ScanOpT expects the scan operation immediately after d_out.

This results in the rather confusing error:

/opt/rocm/bin/hipcc -std=c++11 -O2 -march=native -ffast-math -DNAMD_CUDA  -D__HIP_PLATFORM_HCC__=   -I/opt/rocm/include -I/opt/rocm/hcc/include -I/opt/rocm/include/ -o obj/CudaTileListKernel.o -c `echo src/`CudaTileListKernel.cu
In file included from src/CudaTileListKernel.cu:1:
In file included from src/HIPWrapperCUB.h:42:
In file included from /opt/rocm/include/hipcub/hipcub.hpp:34:
In file included from /opt/rocm/include/hipcub/rocprim/hipcub.hpp:65:
/opt/rocm/include/hipcub/rocprim/device/device_scan.hpp:99:16: error: use of undeclared identifier 'ExclusiveScan'
        return ExclusiveScan(
               ^
src/CudaTileListKernel.cu:1246:38: note: in instantiation of function template specialization 'hipcub::DeviceScan::ExclusiveSum<int *, int *>' requested here
          cudaCheck(cub::DeviceScan::ExclusiveSum(NULL, size,
                                     ^
/opt/rocm/include/hipcub/rocprim/device/device_scan.hpp:113:16: note: must qualify identifier to find this declaration in dependent base class
    hipError_t ExclusiveScan(void *d_temp_storage,
               ^
/opt/rocm/include/hipcub/rocprim/device/device_scan.hpp:99:16: error: no matching function for call to 'ExclusiveScan'
        return ExclusiveScan(
               ^~~~~~~~~~~~~
/opt/rocm/include/hipcub/rocprim/device/device_scan.hpp:113:16: note: candidate function template not viable: no known conversion from '::hipcub::Sum' to 'int' for 7th argument
    hipError_t ExclusiveScan(void *d_temp_storage,
               ^
/opt/rocm/include/hipcub/rocprim/device/device_scan.hpp:99:16: error: no matching function for call to 'ExclusiveScan'
        return ExclusiveScan(
               ^~~~~~~~~~~~~
/opt/rocm/include/hipcub/rocprim/device/device_scan.hpp:113:16: note: candidate function template not viable: no known conversion from '::hipcub::Sum' to 'int' for 7th argument
    hipError_t ExclusiveScan(void *d_temp_storage,
               ^
3 errors generated.

Submitting a PR shortly

[ex-rzr]

Hi!

It will be automatically downloaded and built by CMake script.

Unfortunately NAMD doesn't use CMake so all rocPRIM have to be installed manually (from deb or using make install).

This does appear to be true, but there also seems to be some files / directories missing in the resulting /opt/roc/hipCUB/include/hipcub/rocprim folder. For instance, all of the detail folders (e.g., rocprim/warp/detail) are missing, and all files there-in. More importantly, using rocPRIM as supplied by hipCUB seems to not install the main rocprim.hpp header file... which results in an error as soon as you try to include hipcub.hpp

hipcub/rocprim is a wrapper of rocPRIM, i.e. rocPRIM backend.

Now, if you build and install rocPRIM separately resolves this issue (as the code can now find <rocprim/rocprim.hpp> under /opt/rocm

So this is a right way.

There is a wiki page I created long ago: https://github.com/AMDComputeLibraries/NAMD/wiki

There are instructions how to build NAMD but they need to be updated and rechecked. Many issues in HIP/rocFFT... I reported/fixed are merged now. Also rocPRIM was separated into two libraries with own repositories: rocPRIM and hipCUB.

I'll recheck them (probably) tomorrow.

The issues with ExclusiveScan/ExclusiveSum is fixed in develop branch, please use it (we merge it into master later).

[skyreflectedinmirrors]

@ex-rzr ok, thank you!