migration errors in cub and thrust

[tims111]

while migrating a pytorch file, following errors are observed, any idea? Thanks!

'/usr/local/cuda-11.6/targets/x86_64-linux/include/thrust/system/cuda/detail/core/util.h:534:9: error: no template named 'CacheModifiedInputIterator' in namespace 'thrust::cub'; did you mean 'at_cuda_detail::cub::CacheModifiedInputIterator'?
        cub::CacheModifiedInputIterator<PtxPlan::LOAD_MODIFIER,
        ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
        at_cuda_detail::cub::CacheModifiedInputIterator
/usr/local/cuda-11.6/targets/x86_64-linux/include/cub/block/../iterator/cache_modified_input_iterator.cuh:103:7: note: 'at_cuda_detail::cub::CacheModifiedInputIterator' declared here
class CacheModifiedInputIterator
      ^
/usr/local/cuda-11.6/targets/x86_64-linux/include/thrust/system/cuda/detail/core/util.h:576:18: error: no template named 'BlockLoad' in namespace 'thrust::cub'; did you mean 'at_cuda_detail::cub::BlockLoad'?
    using type = cub::BlockLoad<T,
                 ^~~~~~~~~~~~~~
                 at_cuda_detail::cub::BlockLoad
/usr/local/cuda-11.6/targets/x86_64-linux/include/cub/block/block_load.cuh:642:7: note: 'at_cuda_detail::cub::BlockLoad' declared here
class BlockLoad

/usr/local/cuda-11.6/targets/x86_64-linux/include/thrust/system/cuda/detail/core/util.h:593:18: error: no template named 'BlockStore' in namespace 'thrust::cub'; did you mean 'at_cuda_detail::cub::BlockStore'?
    using type = cub::BlockStore<T,
                 ^~~~~~~~~~~~~~~
                 at_cuda_detail::cub::BlockStore
/usr/local/cuda-11.6/targets/x86_64-linux/include/cub/block/block_store.cuh:532:7: note: 'at_cuda_detail::cub::BlockStore' declared here
class BlockStore

/usr/local/cuda-11.6/targets/x86_64-linux/include/thrust/system/cuda/detail/core/util.h:635:31: error: no member named 'PtxVersion' in namespace 'thrust::cub'
    cudaError_t status = cub::PtxVersion(ptx_version);
                         ~~~~~^
/usr/local/cuda-11.6/targets/x86_64-linux/include/thrust/system/cuda/detail/core/util.h:642:12: error: no member named 'SyncStream' in namespace 'thrust::cub'; did you mean 'at_cuda_detail::cub::SyncStream'?
    return cub::SyncStream(stream);
           ^~~~~~~~~~~~~~~
           at_cuda_detail::cub::SyncStream
/usr/local/cuda-11.6/targets/x86_64-linux/include/cub/block/../iterator/../util_device.cuh:557:41: note: 'at_cuda_detail::cub::SyncStream' declared here
CUB_RUNTIME_FUNCTION inline cudaError_t SyncStream(cudaStream_t stream)

/usr/local/cuda-11.6/targets/x86_64-linux/include/thrust/system/cuda/detail/core/util.h:647:5: error: no member named 'CTA_SYNC' in namespace 'thrust::cub'; did you mean 'at_cuda_detail::cub::CTA_SYNC'?
    cub::CTA_SYNC();
    ^~~~~~~~~~~~~
    at_cuda_detail::cub::CTA_SYNC
/usr/local/cuda-11.6/targets/x86_64-linux/include/cub/block/../util_ptx.cuh:227:34: note: 'at_cuda_detail::cub::CTA_SYNC' declared here
__device__  __forceinline__ void CTA_SYNC()

/usr/local/cuda-11.6/targets/x86_64-linux/include/thrust/system/cuda/detail/core/util.h:663:22: error: no template named 'UnitWord' in namespace 'thrust::cub'; did you mean 'at_cuda_detail::cub::UnitWord'?
    typedef typename cub::UnitWord<T>::DeviceWord DeviceWord;
                     ^~~~~~~~~~~~~
                     at_cuda_detail::cub::UnitWord
/usr/local/cuda-11.6/targets/x86_64-linux/include/cub/util_type.cuh:425:8: note: 'at_cuda_detail::cub::UnitWord' declared here
struct UnitWord

/usr/local/cuda-11.6/targets/x86_64-linux/include/thrust/system/cuda/detail/core/util.h:755:12: error: no member named 'AliasTemporaries' in namespace 'thrust::cub'; did you mean 'at_cuda_detail::cub::AliasTemporaries'?
    return cub::AliasTemporaries(storage_ptr,
           ^~~~~~~~~~~~~~~~~~~~~
           at_cuda_detail::cub::AliasTemporaries
/usr/local/cuda-11.6/targets/x86_64-linux/include/cub/block/../iterator/../util_device.cuh:65:13: note: 'at_cuda_detail::cub::AliasTemporaries' declared here
cudaError_t AliasTemporaries(

[tomflinda]

Pls provide a reproducer and the steps on how to reproduce this issue, so that we can do a further investigation.

[jinz2014]

After reading the post, I tried to run the example https://github.com/NVIDIA/cub/blob/main/examples/block/example_block_radix_sort.cu.

I observe the following message, but the migration does not produce any errors.

warning: DPCT1007:5: Migration of cub::StoreDirectStriped is not supported
warning: DPCT1007:6: Migration of cub::BlockLoad.Load is not supported
warning: DPCT1007:7: Migration of cub::BlockRadixSort.SortBlockedToStriped is not supported
warning: DPCT1007:8: Migration of cub::Debug is not supported

[tomflinda]

@tims111

1. currently these cub APIs(cub::StoreDirectStriped, cub::BlockLoad.Load, cub::BlockRadixSort.SortBlockedToStriped, cub::Debug) have no mappings in SYCL, after the mappings are available, we will plan to support them;
2. as the parsing errors, we cannot reproduce them, could you provide the migration commands and the host system info you used?

[jinz2014]

Another example is the CUB's run length encoding. It seems that DPCT does not convert the function.

https://github.com/zjin-lcf/HeCBench/blob/master/rle-cuda/main.cu

[tomflinda]

@jinz2014 reproduced for DeviceRunLengthEncode migration issue, we plan to fix it, and #1142 is filed to track this issue you reported, and this issue will be closed.

[tims111]

@tomflinda
I took another look. I think the parsing error is related to missing defination in the compilation command. The missing defination is used to control the name space to use. I found even nvcc has that, so I didn't notice I dropped it. I guess if the missing cub API is fixed, it shall be in a better shape.

[tomflinda]

@tims111 Okay, thanks for your feedback, I will close this issue, if you find any problems in future, pls feel free to open new issues to us.

[jinz2014]

@tomflinda Ok. Thank you.