NVIDIA · jrhemstad · Oct 11, 2023 · Sep 26, 2023 · Sep 26, 2023 · Sep 26, 2023
@@ -12,6 +12,10 @@ explode_std_versions() {
   jq -cr 'map(. as $o | {std: $o.std[]} + del($o.std))'
 }
 
+explode_libs() {
+  jq -cr 'map(. as $o | {lib: $o.lib[]} + del($o.lib))'
+}
+
 extract_matrix() {
   local file="$1"
   local type="$2"
@@ -23,6 +27,8 @@ extract_matrix() {
   write_output "HOST_COMPILERS" "$(echo "$nvcc_full_matrix" | jq -cr '[.[] | .compiler.name] | unique')"
   write_output "PER_CUDA_COMPILER_MATRIX" "$(echo "$nvcc_full_matrix" | jq -cr ' group_by(.cuda + .compiler.name) | map({(.[0].cuda + "-" + .[0].compiler.name): .}) | add')"
   write_output "NVRTC_MATRIX" "$(echo "$matrix" | jq '.nvrtc' | explode_std_versions)"
+  local clang_cuda_matrix="$(echo "$matrix" | jq -cr '.["clang-cuda"]' | explode_std_versions | explode_libs)"
+  write_output "CLANG_CUDA_MATRIX" "$clang_cuda_matrix"
   write_output "EXAMPLES_MATRIX" "$(echo "$matrix" | jq -cr '.examples' )"
 }
 

@@ -41,6 +41,7 @@ jobs:
       HOST_COMPILERS: ${{steps.set-outputs.outputs.HOST_COMPILERS}}
       PER_CUDA_COMPILER_MATRIX: ${{steps.set-outputs.outputs.PER_CUDA_COMPILER_MATRIX}}
       NVRTC_MATRIX: ${{steps.set-outputs.outputs.NVRTC_MATRIX}}
+      CLANG_CUDA_MATRIX: ${{steps.set-outputs.outputs.CLANG_CUDA_MATRIX}}
       EXAMPLES_MATRIX: ${{steps.set-outputs.outputs.EXAMPLES_MATRIX}}
     steps:
       - name: Checkout repo
@@ -111,9 +112,25 @@ jobs:
       devcontainer_version: ${{ needs.compute-matrix.outputs.DEVCONTAINER_VERSION }}
       is_windows: ${{ contains(matrix.compiler, 'cl') }}
 
+  clang-cuda:
+    name: ${{matrix.lib}} ${{matrix.cpu}}/CTK${{matrix.cuda}}/clang-cuda
+    needs: compute-matrix
+    strategy:
+      fail-fast: false
+      matrix:
+        include: ${{ fromJSON(needs.compute-matrix.outputs.CLANG_CUDA_MATRIX) }}
+    uses: ./.github/workflows/run-as-coder.yml
+    with:
+      name: ${{matrix.lib}} CTK${{matrix.cuda}} clang-cuda${{matrix.compiler.version}}/${{matrix.std}}
+      runner: linux-${{matrix.cpu}}-cpu16
+      image: rapidsai/devcontainers:${{needs.compute-matrix.outputs.DEVCONTAINER_VERSION}}-cpp-${{matrix.compiler.name}}${{matrix.compiler.version}}-cuda${{matrix.cuda}}-${{matrix.os}}
+      command: |
+        CMAKE_CUDA_COMPILER="${{matrix.compiler.exe}}" ./ci/build_${{matrix.lib}}.sh ${{matrix.compiler.exe}} ${{matrix.std}} ${{matrix.gpu_build_archs}}
+
   examples:
     name: CCCL Examples
     needs: compute-matrix
+    if: ${{ !contains(github.event.head_commit.message, 'skip-tests') }}
     strategy:
       fail-fast: false
       matrix:
@@ -137,6 +154,7 @@ jobs:
     runs-on: ubuntu-latest
     name: CI
     needs:
+      - clang-cuda
       - cub
       - libcudacxx
       - nvrtc

@@ -6,12 +6,13 @@ set -eo pipefail
 cd "$( cd "$( dirname "${BASH_SOURCE[0]}" )" && pwd )";
 
 # Script defaults
-CUDA_COMPILER=nvcc
+CUDA_COMPILER=${CMAKE_CUDA_COMPILER:-nvcc}
 
 # Check if the correct number of arguments has been provided
 function usage {
     echo "Usage: $0 [OPTIONS] <HOST_COMPILER> <CXX_STANDARD> <GPU_ARCHS>"
     echo "The PARALLEL_LEVEL environment variable controls the amount of build parallelism. Default is the number of cores."
+    echo "The CMAKE_CUDA_COMPILER environment variable can be used to control the CUDA compiler. The -nvcc flag takes precedence."
     echo "Example: PARALLEL_LEVEL=8 $0 g++-8 14 \"70\" "
     echo "Example: $0 clang++-8 17 \"70;75;80-virtual\" "
     echo "Possible options: "
@@ -54,9 +55,7 @@ readonly CXX_STANDARD=$2
 
 # Replace spaces, commas and semicolons with semicolons for CMake list
 readonly GPU_ARCHS=$(echo $3 | tr ' ,' ';')
-
 readonly PARALLEL_LEVEL=${PARALLEL_LEVEL:=$(nproc)}
-readonly NVCC_VERSION=$($CUDA_COMPILER --version | grep release | awk '{print $6}' | cut -c2-)
 
 if [ -z ${DEVCONTAINER_NAME+x} ]; then
     BUILD_DIR=../build/local
@@ -83,7 +82,7 @@ COMMON_CMAKE_OPTIONS="
 echo "========================================"
 echo "Begin build"
 echo "pwd=$(pwd)"
-echo "NVCC_VERSION=$NVCC_VERSION"
+echo "CUDA_COMPILER=$CUDA_COMPILER"
 echo "HOST_COMPILER=$HOST_COMPILER"
 echo "CXX_STANDARD=$CXX_STANDARD"
 echo "GPU_ARCHS=$GPU_ARCHS"

@@ -2,7 +2,6 @@
 
 source "$(dirname "$0")/build_common.sh"
 
-
 # CUB benchmarks require at least CUDA nvcc 11.5 for int128
 # Returns "true" if the first version is greater than or equal to the second
 version_compare() {
@@ -12,12 +11,20 @@ version_compare() {
         echo "false"
     fi
 }
-readonly ENABLE_CUB_BENCHMARKS=${ENABLE_CUB_BENCHMARKS:=$(version_compare $NVCC_VERSION 11.5)}
 
-if [[ $ENABLE_CUB_BENCHMARKS == "true" ]]; then
-    echo "CUDA version is $NVCC_VERSION. Building CUB benchmarks."
+ENABLE_CUB_BENCHMARKS="false"
+ENABLE_CUB_RDC="false"
+if [[ "$CUDA_COMPILER" == *nvcc* ]]; then
+    ENABLE_CUB_RDC="true"
+    NVCC_VERSION=$($CUDA_COMPILER --version | grep release | awk '{print $6}' | cut -c2-)
+    if [[ $(version_compare $NVCC_VERSION 11.5) == "true" ]]; then
+        ENABLE_CUB_BENCHMARKS="true"
+        echo "nvcc version is $NVCC_VERSION. Building CUB benchmarks."
+    else
+        echo "nvcc version is $NVCC_VERSION. Not building CUB benchmarks because nvcc version is less than 11.5."
+    fi
 else
-    echo "CUDA version is $NVCC_VERSION. Not building CUB benchmarks because CUDA version is less than 11.5."
+    echo "nvcc version is not determined (likely using a non-NVCC compiler). Not building CUB benchmarks."
 fi
 
 CMAKE_OPTIONS="
@@ -32,6 +39,7 @@ CMAKE_OPTIONS="
     -DTHRUST_IGNORE_DEPRECATED_CPP_DIALECT=ON \
     -DCUB_IGNORE_DEPRECATED_CPP_DIALECT=ON \
     -DCUB_ENABLE_BENCHMARKS="$ENABLE_CUB_BENCHMARKS"\
+    -DCUB_ENABLE_RDC_TESTS="$ENABLE_CUB_RDC" \
 "
 
 configure_and_build "CUB" "$CMAKE_OPTIONS"
@@ -51,6 +51,8 @@ pull_request:
     - {cuda: *cuda_newest, os: 'windows2022', cpu: 'amd64', compiler: {name: 'cl',   version: '14.36', exe: 'cl++'}, gpu_build_archs: '70', std: [14, 17, 20],     jobs: ['build']}
   nvrtc:
     - {cuda: *cuda_newest, os: 'ubuntu22.04', cpu: 'amd64', gpu_build_archs: '70', std: [11, 14, 17, 20]}
+  clang-cuda:
+    - {lib: ['thrust', 'cub', 'libcudacxx'], cuda: *cuda_newest, os: 'ubuntu22.04', cpu: 'amd64', compiler: {name: 'llvm', version: '16', exe: 'clang++'}, gpu_build_archs: '70', std: [17, 20]}
   examples:
     # Strategy: Oldest CUDA + Oldest Host compiler && Newest CUDA + Newest Host Compiler
     - {cuda: *cuda_oldest, os: 'ubuntu18.04', cpu: 'amd64', compiler: {name: 'gcc', version: '6', exe: 'g++'}}

@@ -32,10 +32,10 @@ else() # NOT LIBCUDACXX_TEST_WITH_NVRTC
   set(LIBCUDACXX_CUDA_COMPILER "${CMAKE_CUDA_COMPILER}")
 endif()
 
-if (NOT MSVC)
+if (NOT MSVC AND NOT ${CMAKE_CUDA_COMPILER_ID} STREQUAL "Clang")
   set(LIBCUDACXX_WARNING_LEVEL
     "--compiler-options=-Wall \
-     --compiler-options=-Wextra")
+    --compiler-options=-Wextra")
 endif()
 
 # sccache cannot handle the -Fd option generationg pdb files
@@ -50,6 +50,18 @@ if (${CMAKE_CXX_COMPILER_ID} STREQUAL "IntelLLVM")
     --compiler-options=-fno-fast-math")
 endif()
 
+if (${CMAKE_CUDA_COMPILER_ID} STREQUAL "Clang")
+  string(APPEND LIBCUDACXX_TEST_COMPILER_FLAGS
+     " -Xclang -fcuda-allow-variadic-functions"
+     " -Xclang -Wno-unused-parameter"
+     " -Wno-unknown-cuda-version")
+
+  find_package(CUDAToolkit)
+
+  string(APPEND LIBCUDACXX_TEST_LINKER_FLAGS
+    " -L${CUDAToolkit_LIBRARY_DIR} -lcuda -lcudart")
+endif()
+
 if (${CMAKE_CUDA_COMPILER_ID} STREQUAL "NVIDIA")
   set(LIBCUDACXX_TEST_COMPILER_FLAGS
     "${LIBCUDACXX_TEST_COMPILER_FLAGS} \
@@ -67,6 +79,14 @@ if (${CMAKE_CUDA_COMPILER_ID} STREQUAL "NVHPC")
     -stdpar")
 endif()
 
+if (${CMAKE_CUDA_COMPILER_ID} STREQUAL "Clang")
+  set(LIBCUDACXX_TEST_COMPILER_FLAGS
+    "${LIBCUDACXX_TEST_COMPILER_FLAGS} \
+    ${LIBCUDACXX_FORCE_INCLUDE} \
+    -I${libcudacxx_SOURCE_DIR}/include \
+    ${LIBCUDACXX_WARNING_LEVEL}")
+endif()
+
 set(LIBCUDACXX_COMPUTE_ARCHS_STRING
   "${CMAKE_CUDA_ARCHITECTURES}")
 

@@ -21,17 +21,18 @@ TEST_NV_DIAG_SUPPRESS(set_but_not_used)
 __device__
 void test()
 {
-    __shared__ cuda::barrier<cuda::thread_scope_block> b;
-    init(&b, 2);
+    __shared__ cuda::barrier<cuda::thread_scope_block>* b;
+    shared_memory_selector<cuda::barrier<cuda::thread_scope_block>, constructor_initializer> sel;
+    b = sel.construct(2);
 
     uint64_t token;
     asm volatile ("mbarrier.arrive.b64 %0, [%1];"
         : "=l"(token)
-        : "l"(cuda::device::barrier_native_handle(b))
+        : "l"(cuda::device::barrier_native_handle(*b))
         : "memory");
     (void)token;
 
-    b.arrive_and_wait();
+    b->arrive_and_wait();
 }
 
 int main(int argc, char ** argv)

@@ -47,7 +47,7 @@ struct storage
     __host__ __device__
     friend bool operator==(const storage & lhs, const T & rhs) {
         for (cuda::std::size_t i = 0; i < size; ++i) {
-            if (lhs.data[i] != rhs + i) {
+            if (lhs.data[i] != static_cast<T>(rhs + i)) {
                 return false;
             }
         }

@@ -9,6 +9,7 @@
 //===----------------------------------------------------------------------===//
 
 // UNSUPPORTED: pre-sm-70
+// UNSUPPORTED: clang && (!nvcc)
 
 #define _LIBCUDACXX_CUDA_ABI_VERSION 2
 

@@ -20,6 +20,7 @@
 #include <stdlib.h>
 
 #include "test_macros.h"
+#include "cuda_space_selector.h"
 
 TEST_NV_DIAG_SUPPRESS(static_var_with_dynamic_init)
 TEST_NV_DIAG_SUPPRESS(186) // pointless comparison of unsigned integer with zero
@@ -30,7 +31,7 @@ constexpr size_t stages_count = 2; // Pipeline with two stages
 // Simply copy shared memory to global out
 __device__ __forceinline__ void compute(int* global_out, int const* shared_in){
     auto block = cooperative_groups::this_thread_block();
-    for (int i = 0; i < block.size(); ++i) {
+    for (int i = 0; i < static_cast<int>(block.size()); ++i) {
         global_out[i] = shared_in[i];
     }
 }
@@ -46,8 +47,11 @@ __global__ void with_staging(int* global_out, int const* global_in, size_t size,
     size_t shared_offset[stages_count] = { 0, block.size() }; // Offsets to each batch
 
     // Allocate shared storage for a two-stage cuda::pipeline:
-    __shared__ cuda::pipeline_shared_state<cuda::thread_scope::thread_scope_block, stages_count> shared_state;
-    auto pipeline = cuda::make_pipeline(block, &shared_state);
+    using pipeline_state = cuda::pipeline_shared_state<cuda::thread_scope::thread_scope_block, stages_count>;
+    __shared__ pipeline_state* shared_state;
+    shared_memory_selector<pipeline_state, constructor_initializer> sel;
+    shared_state = sel.construct();
+    auto pipeline = cuda::make_pipeline(block, shared_state);
 
     // Each thread processes `batch_sz` elements.
     // Compute offset of the batch `batch` of this thread block in global memory:

@@ -21,11 +21,13 @@
 #  define TEST_NVCC
 #elif defined(__NVCOMPILER)
 #  define TEST_NVCXX
+#elif defined(__CUDA__)
+#  define TEST_CLANG_CUDA
 #else
 #  define TEST_HOST
 #endif
 
-#if defined(TEST_NVCC)
+#if defined(TEST_NVCC) || defined(TEST_CLANG_CUDA)
 
 __host__ __device__ void test() {
 #if defined(__CUDA_ARCH__)

@@ -93,6 +93,6 @@ int main(int argc, char** argv)
     return ret;
 }
 
-#define main fake_main
+#define main __host__ __device__ fake_main
 
 #endif
@@ -7,6 +7,7 @@
 //===----------------------------------------------------------------------===//
 
 // UNSUPPORTED: nvrtc, pre-sm-70
+// XFAIL: clang && (!nvcc)
 
 // uncomment for a really verbose output detailing what test steps are being launched
 // #define DEBUG_TESTERS

@@ -7,6 +7,7 @@
 //===----------------------------------------------------------------------===//
 
 // UNSUPPORTED: nvrtc, pre-sm-70
+// XFAIL: clang && (!nvcc)
 
 // uncomment for a really verbose output detailing what test steps are being launched
 // #define DEBUG_TESTERS

@@ -80,7 +80,7 @@ using aw_aw_pw1 = performer_list<
     barrier_arrive_and_wait,
     barrier_arrive_and_wait,
     async_tester_fence,
-    clear_token,
+    clear_token
 >;
 
 using aw_aw_pw2 = performer_list<

@@ -166,7 +166,7 @@ using arithmetic_atomic_testers = extend_tester_list<
     fetch_sub_tester<30, 21, 9>,
     fetch_min_tester<9, 5, 5>,
     fetch_max_tester<5, 9, 9>,
-    fetch_sub_tester<9, 17, -8>,
+    fetch_sub_tester<9, 17, -8>
 >;
 
 using bitwise_atomic_testers = extend_tester_list<

@@ -11,6 +11,7 @@
 
 #include <cuda/std/type_traits>
 
+#include <cstdio>
 #include <new>
 #include <thread>
 #include <vector>

@@ -16,6 +16,8 @@
 #include <cuda/std/concepts>
 #include <cuda/std/type_traits>
 
+#include "test_macros.h"
+
 using cuda::std::same_as;
 
 struct S1 {};
@@ -35,6 +37,9 @@ struct S5 {
   int* p;
 };
 
+#ifdef TEST_COMPILER_CLANG_CUDA
+#pragma clang diagnostic ignored "-Wunused-private-field"
+#endif // TEST_COMPILER_CLANG_CUDA
 class C1 {};
 class C2 {
   /* [[maybe_unused]] */ int i;

@@ -55,9 +55,6 @@ int main(int, char**)
 
     // TEST(TestLayoutLeftListInitialization, test_layout_left_extent_initialization)
     {
-        typedef int    data_t ;
-        typedef size_t index_t;
-
         cuda::std::layout_left::mapping<cuda::std::extents<size_t,dyn, dyn>> m{cuda::std::dextents<size_t,2>{16, 32}};
 
         static_assert( m.is_exhaustive()          == true, "" );

@@ -35,7 +35,6 @@ int main(int, char**)
     typed_test_compare_left();
 
     using index_t = size_t;
-    using ext1d_t = cuda::std::extents<index_t,dyn>;
     using ext2d_t = cuda::std::extents<index_t,dyn,dyn>;
 
     {

@@ -35,7 +35,6 @@ int main(int, char**)
     typed_test_compare_right();
 
     using index_t = size_t;
-    using ext1d_t = cuda::std::extents<index_t,dyn>;
     using ext2d_t = cuda::std::extents<index_t,dyn,dyn>;
 
     {