Add additional build job for sm90

.github/workflows/dispatch-build-and-test.yml

@@ -28,9 +28,9 @@ jobs:
         include: ${{ fromJSON(inputs.per_cuda_compiler_matrix) }}
     with:
       cpu: ${{ matrix.cpu }}
-      test_name: ${{matrix.cpu}}/${{matrix.compiler.name}}${{matrix.compiler.version}}/C++${{matrix.std}}
+      test_name: ${{matrix.cpu}}/${{matrix.compiler.name}}${{matrix.compiler.version}}/C++${{matrix.std}} ${{matrix.extra_build_args}}
       build_script: './ci/build_${{ inputs.project_name }}.sh -cxx ${{matrix.compiler.exe}} -std ${{matrix.std}} "${{matrix.extra_build_args}}"'
-      test_script:  './ci/test_${{ inputs.project_name }}.sh -cxx ${{matrix.compiler.exe}} -std ${{matrix.std}}'
+      test_script:  './ci/test_${{ inputs.project_name }}.sh -cxx ${{matrix.compiler.exe}} -std ${{matrix.std}} "${{matrix.extra_build_args}}"'
       container_image: rapidsai/devcontainers:${{inputs.devcontainer_version}}-cpp-${{matrix.compiler.name}}${{matrix.compiler.version}}-cuda${{matrix.cuda}}-${{matrix.os}}
       run_tests: ${{ contains(matrix.jobs, 'test') && !contains(github.event.head_commit.message, 'skip-tests') && matrix.os != 'windows-2022' }}
 

ci/matrix.yaml

@@ -67,6 +67,7 @@ pull_request:
     - {cuda: *cuda_newest, os: 'ubuntu20.04', cpu: 'amd64', compiler: *gcc9,     std: [11, 14, 17],     jobs: ['build']}
     - {cuda: *cuda_newest, os: 'ubuntu20.04', cpu: 'amd64', compiler: *gcc10,    std: [11, 14, 17, 20], jobs: ['build']}
     - {cuda: *cuda_newest, os: 'ubuntu22.04', cpu: 'amd64', compiler: *gcc11,    std: [11, 14, 17, 20], jobs: ['build']}
+    - {cuda: *cuda_newest, os: 'ubuntu22.04', cpu: 'amd64', compiler: *gcc12,    std: [11, 14, 17, 20], jobs: ['build'], extra_build_args: '-cmake-options -DCMAKE_CUDA_ARCHITECTURES=90a'}
     - {cuda: *cuda_newest, os: 'ubuntu22.04', cpu: 'amd64', compiler: *gcc12,    std: [11, 14, 17, 20], jobs: ['build', 'test']}
     - {cuda: *cuda_newest, os: 'ubuntu22.04', cpu: 'arm64', compiler: *gcc12,    std: [11, 14, 17, 20], jobs: ['build']}
     - {cuda: *cuda_newest, os: 'ubuntu20.04', cpu: 'amd64', compiler: *llvm9,    std: [11, 14, 17],     jobs: ['build']}

libcudacxx/include/cuda/std/detail/libcxx/include/__cuda/barrier.h

@@ -774,7 +774,6 @@ _CUDA_VSTD::uint64_t * __try_get_barrier_handle<::cuda::thread_scope_block, _CUD
 // The user is still responsible for arriving and waiting on (or otherwise
 // synchronizing with) the barrier or pipeline barrier to see the results of
 // copies from other threads participating in the synchronization object.
-extern "C" _LIBCUDACXX_HOST_DEVICE void __cuda_ptx_mbarrier_complete_tx_is_not_supported_before_SM_90__();
 struct __memcpy_completion_impl {
 
     template<typename _Group>
@@ -815,8 +814,6 @@ struct __memcpy_completion_impl {
                     if (__group.thread_rank() == 0) {
                         ::cuda::device::barrier_expect_tx(__barrier, __size);
                     }
-                ),(
-                    __cuda_ptx_mbarrier_complete_tx_is_not_supported_before_SM_90__();
                 ));
 #endif // __cccl_ptx_isa >= 800
                 return async_contract_fulfillment::async;

libcudacxx/libcxx/test/std/containers/views/mdspan/mdspan.layout.util/layout_util.hpp

@@ -8,6 +8,8 @@
 
 #include <tuple>
 
+#include "test_macros.h"
+
 template <class> struct TestLayoutCtors;
 template <class Mapping, size_t... DynamicSizes>
 struct TestLayoutCtors<std::tuple<
@@ -162,7 +164,3 @@ struct is_stride_avail< T
 
 template< class T, class RankType >
 constexpr bool is_stride_avail_v = is_stride_avail< T, RankType >::value;
-
-// Workaround for variables that are only used in static_assert's
-template< typename T >
-constexpr bool unused( T && ) { return true; }

libcudacxx/test/libcudacxx/cuda/annotated_ptr/utils.h

@@ -32,9 +32,6 @@ void assert_rt_wrap(cudaError_t code, const char *file, int line) {
 }
 #define assert_rt(ret) { assert_rt_wrap((ret), __FILE__, __LINE__); }
 
-template <typename ... T>
-__host__ __device__ constexpr bool unused(T...) {return true;}
-
 template<typename T, int N>
 __device__ __host__ __noinline__
 T* alloc(bool shared = false) {

libcudacxx/test/libcudacxx/cuda/barrier/arrive_tx_cluster.runfail.cpp

@@ -11,6 +11,7 @@
 // UNSUPPORTED: libcpp-has-no-threads
 // UNSUPPORTED: pre-sm-90
 // UNSUPPORTED: nvcc-11
+// UNSUPPORTED: no_execute
 
 // <cuda/barrier>
 

libcudacxx/test/libcudacxx/cuda/barrier/arrive_tx_device.runfail.cpp

@@ -11,6 +11,7 @@
 // UNSUPPORTED: libcpp-has-no-threads
 // UNSUPPORTED: pre-sm-90
 // UNSUPPORTED: nvcc-11
+// UNSUPPORTED: no_execute
 
 // <cuda/barrier>
 

libcudacxx/test/libcudacxx/cuda/barrier/expect_tx_device.runfail.cpp

@@ -11,6 +11,7 @@
 // UNSUPPORTED: libcpp-has-no-threads
 // UNSUPPORTED: pre-sm-90
 // UNSUPPORTED: nvcc-11
+// UNSUPPORTED: no_execute
 
 // <cuda/barrier>