[Clang] PR https://github.com/llvm/llvm-project/pull/80801 cause CUDA 9.2 thrust library compile error

[yihanwg]

#80801 cause CUDA 9.2 thrust library compile error.
This problem exists from 9.2 to trunk. https://github.com/NVIDIA/cccl/blob/main/cub/cub/thread/thread_store.cuh .

cuda-9.2/include/thrust/system/cuda/detail/cub/block/../iterator/../thread/thread_store.cuh:351:56: error: a template argument list is expected after a name prefixed by the template keyword [-Wmissing-template-arg-list-after-template-kw]
  351 |     IterateThreadStore<0, VOLATILE_MULTIPLE>::template Dereference(
      |                                                        ^

/**
 * ThreadStore definition for STORE_VOLATILE modifier on non-primitive pointer types
 */
template <typename T>
__device__ __forceinline__ void ThreadStoreVolatilePtr(
    T                           *ptr,
    T                           val,
    Int2Type<false>             /*is_primitive*/)
{
    // Create a temporary using shuffle-words, then store using volatile-words
    typedef typename UnitWord<T>::VolatileWord  VolatileWord;  
    typedef typename UnitWord<T>::ShuffleWord   ShuffleWord;

    const int VOLATILE_MULTIPLE = sizeof(T) / sizeof(VolatileWord);
    const int SHUFFLE_MULTIPLE  = sizeof(T) / sizeof(ShuffleWord);
    
    VolatileWord words[VOLATILE_MULTIPLE];

    #pragma unroll
    for (int i = 0; i < SHUFFLE_MULTIPLE; ++i)
        reinterpret_cast<ShuffleWord*>(words)[i] = reinterpret_cast<ShuffleWord*>(&val)[i];

    IterateThreadStore<0, VOLATILE_MULTIPLE>::template Dereference(
        reinterpret_cast<volatile VolatileWord*>(ptr),
        words);
}

[yihanwg]

AFAIK, a lot of existing code were broken by CWG 96.

[llvmbot]

@llvm/issue-subscribers-clang-frontend

Author: Yihan Wang (yihanwg)

https://github.com//pull/80801 cause CUDA 9.2 thrust library compile error. This problem exists from 9.2 to trunk. https://github.com/NVIDIA/cccl/blob/main/cub/cub/thread/thread_store.cuh .

cuda-9.2/include/thrust/system/cuda/detail/cub/block/../iterator/../thread/thread_store.cuh:351:56: error: a template argument list is expected after a name prefixed by the template keyword [-Wmissing-template-arg-list-after-template-kw]
  351 |     IterateThreadStore<0, VOLATILE_MULTIPLE>::template Dereference(
      |                                                        ^

/**
 * ThreadStore definition for STORE_VOLATILE modifier on non-primitive pointer types
 */
template <typename T>
__device__ __forceinline__ void ThreadStoreVolatilePtr(
    T                           *ptr,
    T                           val,
    Int2Type<false>             /*is_primitive*/)
{
    // Create a temporary using shuffle-words, then store using volatile-words
    typedef typename UnitWord<T>::VolatileWord  VolatileWord;  
    typedef typename UnitWord<T>::ShuffleWord   ShuffleWord;

    const int VOLATILE_MULTIPLE = sizeof(T) / sizeof(VolatileWord);
    const int SHUFFLE_MULTIPLE  = sizeof(T) / sizeof(ShuffleWord);
    
    VolatileWord words[VOLATILE_MULTIPLE];

    #pragma unroll
    for (int i = 0; i < SHUFFLE_MULTIPLE; ++i)
        reinterpret_cast<ShuffleWord*>(words)[i] = reinterpret_cast<ShuffleWord*>(&val)[i];

    IterateThreadStore<0, VOLATILE_MULTIPLE>::template Dereference(
        reinterpret_cast<volatile VolatileWord*>(ptr),
        words);
}

[Artem-B]

This sounds like something to bring to the attention of CCCL folks. They do run their tests with clang, but may not be using the most recent builds yet.

@jrhemstad -- FYI.