[FEA] Make cuco compilable using clang

[MatthiasKohl]

Is your feature request related to a problem? Please describe.
Any project dependent on cuco cannot be compiled using clang if the headers in cuco cannot be compiled using clang.

Describe the solution you'd like
It would be great if cuco has a check in place verifying that commits can be compiled using clang.
I'm currently setting something like this up for RAFT (see rapidsai/raft#424) so I'd be happy to provide something similar for cuco.

Additional context
In order to use something like clang-tidy, plus to get better compiler diagnostics, it's always good that code can be compiled using clang.

A few specific things I ran into while applying the above mentioned PR to RAFT (all applies to git tag f0eecb203590f1f4ac4a9f1700229f4434ac64dc of this repo which RAFT depends on at the moment):

1. include/cuco/detail/static_map.inl:336:20, include/cuco/detail/static_map.inl:252:20, include/cuco/detail/static_map.inl:216:21, include/cuco/detail/static_map.inl:324:23: should use this->[[next/initial]]_slot
2. include/cuco/detail/static_map.inl:119:12: requires typename keyword before static_map

[PointKernel]

@MatthiasKohl Can you please specify how you build cuco with clang? Any particular version of clang you want to test with?

[MatthiasKohl]

@PointKernel Basically, I'd like to be able to build RAFT (rapids primitives) with clang, including the device code. This is to support tools like clang-tidy in RAFT, as well as getting better diagnostics in dependent projects.
I can imagine other projects dependent on cuCollections may have similar requirements, either right now or in the future.

For RAPIDS, it would be great to support clang 11+, but if that's difficult for some reason, clang 13+ is good enough.

EDIT: didn't mention this before: RAFT depends on cuCollections for some of its device code, although it's a small-ish part of RAFT.

[sleeepyjack]

I can take care of this issue since I am also using clang tools in my dev pipeline. -> please assign me

[PointKernel]

@sleeepyjack I started this adventure a while back but then got distracted.

IIRC, the main blocking point was build failures in dependencies like thrust and nvbench. There are issues in cuco as well but I think there is not much we can do before resolving the former.

[sleeepyjack]

Okay, maybe I was a bit too eager then 😆 However, I can take care of tracking the associated issues in the external projects and pinging the people in charge. I really don't want to miss out on clang-tidy et al. while developing on cucu. Maybe we can have a meeting where you can share your findings with me?

[PointKernel]

@sleeepyjack Will ping you once I recall what I've done.

[MatthiasKohl]

@sleeepyjack in case you're interested, we have a script in Raft which turns nvcc/g++ compiler commands into clang++ commands, based on a cmake compilation command database.
So if you're using cmake, you can use this to test all the compiler commands

[AustinSchuh]

FYI, you quickly hit NVIDIA/cccl#1020 .

[AustinSchuh]

I'm now down to 1 error (at least when I try to include cuco/static_map.cuh, I haven't tried anything else yet)

In file included from frc971/orin/cuda.cc:3:
In file included from external/com_github_nvidia_cuco/include/cuco/static_map.cuh:24:
In file included from external/com_github_nvidia_cuco/include/cuco/pair.cuh:147:
In file included from external/com_github_nvidia_cuco/include/cuco/detail/pair/pair.inl:54:
external/com_github_nvidia_cuco/include/cuco/detail/pair/tuple_helpers.inl:18:8: error: 
      cannot specialize a dependent template
   18 | struct tuple_size<typename cuco::pair<T1, T2>> : integral_constant<size_t, 2> {
      |        ^
external/com_github_nvidia_cuco/include/cuco/detail/pair/tuple_helpers.inl:22:8: error: 
      cannot specialize a dependent template
   22 | struct tuple_size<const cuco::pair<T1, T2>> : tuple_size<cuco::pair<T1, T2>> {
      |        ^
external/com_github_nvidia_cuco/include/cuco/detail/pair/tuple_helpers.inl:26:8: error: 
      cannot specialize a dependent template
   26 | struct tuple_size<volatile cuco::pair<T1, T2>> : tuple_size<cuco::pair<T1, T2>> {
      |        ^
external/com_github_nvidia_cuco/include/cuco/detail/pair/tuple_helpers.inl:30:8: error: 
      cannot specialize a dependent template
   30 | struct tuple_size<const volatile cuco::pair<T1, T2>> : tuple_size<cuco::pair<T1, T2>> {
      |        ^
external/com_github_nvidia_cuco/include/cuco/detail/pair/tuple_helpers.inl:34:8: error: 
      cannot specialize a dependent template
   34 | struct tuple_element<I, cuco::pair<T1, T2>> {
      |        ^
external/com_github_nvidia_cuco/include/cuco/detail/pair/tuple_helpers.inl:39:8: error: 
      cannot specialize a dependent template
   39 | struct tuple_element<0, cuco::pair<T1, T2>> {
      |        ^
external/com_github_nvidia_cuco/include/cuco/detail/pair/tuple_helpers.inl:44:8: error: 
      cannot specialize a dependent template
   44 | struct tuple_element<1, cuco::pair<T1, T2>> {
      |        ^
external/com_github_nvidia_cuco/include/cuco/detail/pair/tuple_helpers.inl:49:8: error: 
      cannot specialize a dependent template
   49 | struct tuple_element<0, const cuco::pair<T1, T2>> : tuple_element<0, cuco::pair<T1, T2>> {
      |        ^
external/com_github_nvidia_cuco/include/cuco/detail/pair/tuple_helpers.inl:53:8: error: 
      cannot specialize a dependent template
   53 | struct tuple_element<1, const cuco::pair<T1, T2>> : tuple_element<1, cuco::pair<T1, T2>> {
      |        ^
external/com_github_nvidia_cuco/include/cuco/detail/pair/tuple_helpers.inl:57:8: error: 
      cannot specialize a dependent template
   57 | struct tuple_element<0, volatile cuco::pair<T1, T2>> : tuple_element<0, cuco::pair<T1, T2>> {
      |        ^
external/com_github_nvidia_cuco/include/cuco/detail/pair/tuple_helpers.inl:61:8: error: 
      cannot specialize a dependent template
   61 | struct tuple_element<1, volatile cuco::pair<T1, T2>> : tuple_element<1, cuco::pair<T1, T2>> {
      |        ^
external/com_github_nvidia_cuco/include/cuco/detail/pair/tuple_helpers.inl:65:8: error: 
      cannot specialize a dependent template
   65 | struct tuple_element<0, const volatile cuco::pair<T1, T2>> : tuple_element<0, cuco::pair<T1, T2>> {
      |        ^
external/com_github_nvidia_cuco/include/cuco/detail/pair/tuple_helpers.inl:69:8: error: 
      cannot specialize a dependent template
   69 | struct tuple_element<1, const volatile cuco::pair<T1, T2>> : tuple_element<1, cuco::pair<T1, T2>> {
      |        ^
13 errors generated when compiling for sm_80.

I'm having trouble convincing myself if the following is going in the right direction, but it does fix the compiler error...

diff --git a/include/cuco/detail/pair/pair.inl b/include/cuco/detail/pair/pair.inl
index 3279a91..496c969 100644
--- a/include/cuco/detail/pair/pair.inl
+++ b/include/cuco/detail/pair/pair.inl
@@ -50,9 +50,10 @@ __host__ __device__ constexpr bool operator==(cuco::pair<T1, T2> const& lhs,
 
 }  // namespace cuco
 
-namespace thrust {
+/*namespace thrust {
 #include <cuco/detail/pair/tuple_helpers.inl>
 }  // namespace thrust
+*/
 
 namespace cuda::std {
 #include <cuco/detail/pair/tuple_helpers.inl>

In thrust/pair.h, we see:

template <class T>
using tuple_size = ::cuda::std::tuple_size<T>;

Which leads me to believe that the stanza in include/cuco/detail/pair/pair.inl covers it:

namespace cuda::std {
#include <cuco/detail/pair/tuple_helpers.inl>
}  // namespace cuda::std

I can write a small kernel which hits https://github.com/NVIDIA/cuCollections/blob/dev/include/cuco/pair.cuh#L103 but I can't hit https://github.com/NVIDIA/cuCollections/blob/dev/include/cuco/pair.cuh#L89

cuco::pair<int, int> f(5, 3); 
const cuda::std::pair<int, int> spair(5, 3); 
cuco::pair<int, int> g(spair);
cuco::pair<int, int> e(f);
if (e  == g && e == f) {
  printf("success\n");
} else {
  printf("faillure\n");
}