diff --git a/.github/actions/workflow-build/build-workflow.py b/.github/actions/workflow-build/build-workflow.py index fa79da282a67..3ff859ef0f96 100755 --- a/.github/actions/workflow-build/build-workflow.py +++ b/.github/actions/workflow-build/build-workflow.py @@ -334,6 +334,19 @@ def get_job_type_info(job): return result +@memoize_result +def get_codegen_target(codegen_target): + if codegen_target not in matrix_yaml["codegen_targets"]: + raise Exception( + f"Unknown codegen target '{codegen_target}'. Valid options are: " + + ", ".join(matrix_yaml["codegen_targets"].keys()) + ) + + result = matrix_yaml["codegen_targets"][codegen_target] + result["id"] = codegen_target + return result + + @memoize_result def get_tag_info(tag): if tag not in matrix_yaml["tags"].keys(): @@ -364,6 +377,7 @@ def get_all_matrix_job_tags_sorted(): sorted_important_tags = [ "project", "jobs", + "codegen_target", "cudacxx", "cxx", "ctk", @@ -432,6 +446,10 @@ def generate_dispatch_group_name(matrix_job): def generate_dispatch_job_name(matrix_job, job_type): job_info = get_job_type_info(job_type) + job_name = job_info["name"] + if "codegen_target" in matrix_job: + codegen_target = get_codegen_target(matrix_job["codegen_target"]) + job_name += f" {codegen_target['name']}" cpu_str = matrix_job["cpu"] if job_info["gpu"]: gpu = get_gpu(matrix_job["gpu"]) @@ -470,7 +488,7 @@ def generate_dispatch_job_name(matrix_job, job_type): else "" ) - return f"[{config_tag}] {job_info['name']}({cpu_str}{gpu_str}){extra_info}" + return f"[{config_tag}] {job_name}({cpu_str}{gpu_str}){extra_info}" def generate_dispatch_job_runner(matrix_job, job_type): @@ -556,6 +574,13 @@ def generate_dispatch_job_command(matrix_job, job_type): command += f' -py-version "{py_version}"' if py_ctk_mode: command += f' -ctk-mode "{py_ctk_mode}"' + if "codegen_target" in matrix_job: + codegen_target = get_codegen_target(matrix_job["codegen_target"]) + command += f' -target "{codegen_target["cmake_target"]}"' + command += ( + " -cmake-options " + f'"-DLIBCUDACXX_CODEGEN_FILECHECK_TESTS={codegen_target["id"]}"' + ) if extra_args: command += f" {extra_args}" @@ -599,6 +624,11 @@ def generate_dispatch_job_origin(matrix_job, job_type): if "args" in origin_job and not origin_job["args"]: del origin_job["args"] + if "codegen_target" in origin_job: + origin_job["codegen_target"] = get_codegen_target(origin_job["codegen_target"])[ + "name" + ] + origin["matrix_job"] = origin_job return origin @@ -1054,6 +1084,33 @@ def validate_tags(matrix_job, ignore_required=False): error_message_with_matrix_job(matrix_job, f"Unknown tag '{tag}'") ) + jobs = matrix_job.get("jobs", []) + jobs = jobs if isinstance(jobs, list) else [jobs] + has_codegen_job = "codegen_filecheck" in jobs + if has_codegen_job and "codegen_target" not in matrix_job: + raise Exception( + error_message_with_matrix_job( + matrix_job, + "The codegen_filecheck job requires a codegen_target tag.", + ) + ) + if "codegen_target" in matrix_job and any( + job != "codegen_filecheck" for job in jobs + ): + raise Exception( + error_message_with_matrix_job( + matrix_job, + "The codegen_target tag is only valid for codegen_filecheck jobs.", + ) + ) + if "codegen_target" in matrix_job: + codegen_targets = matrix_job["codegen_target"] + codegen_targets = ( + codegen_targets if isinstance(codegen_targets, list) else [codegen_targets] + ) + for codegen_target in codegen_targets: + get_codegen_target(codegen_target) + if "gpu" in matrix_job: gpus = ( matrix_job["gpu"] diff --git a/CMakePresets.json b/CMakePresets.json index 6865c3434fae..60e60bb9059b 100644 --- a/CMakePresets.json +++ b/CMakePresets.json @@ -145,6 +145,16 @@ "LIBCUDACXX_ENABLE_LIBCUDACXX_TESTS": true } }, + { + "name": "libcudacxx-codegen-filecheck", + "displayName": "libcu++: Codegen FileCheck", + "inherits": "libcudacxx", + "cacheVariables": { + "CMAKE_CUDA_ARCHITECTURES": "80", + "LIBCUDACXX_CODEGEN_FILECHECK_TESTS": "all", + "LIBCUDACXX_REQUIRE_CODEGEN_TEST_TOOLS": true + } + }, { "name": "libcudacxx-cpp17", "displayName": "libcu++: C++17", @@ -530,12 +540,20 @@ "libcudacxx.test.public_headers_host_only", "libcudacxx.test.lit.precompile", "libcudacxx.test.nvtarget", - "libcudacxx.test.atomics.ptx", - "libcudacxx.test.simd.ptx", "libcudacxx.test.c2h_all", "libcudacxx.test.debugging" ] }, + { + "name": "libcudacxx-codegen-filecheck", + "configurePreset": "libcudacxx-codegen-filecheck", + "targets": [ + "libcudacxx.test.atomics.ptx", + "libcudacxx.test.atomics.sass", + "libcudacxx.test.simd.ptx", + "libcudacxx.test.simd.sass" + ] + }, { "name": "libcudacxx-cpp17", "configurePreset": "libcudacxx-cpp17", diff --git a/ci/build_libcudacxx.sh b/ci/build_libcudacxx.sh index 29a8708ac768..522bf051ed39 100755 --- a/ci/build_libcudacxx.sh +++ b/ci/build_libcudacxx.sh @@ -4,14 +4,52 @@ set -euo pipefail ci_dir=$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd) +build_target="" +codegen_tests=false +common_args=() +while [[ $# -ne 0 ]]; do + case "$1" in + -codegen-tests) + codegen_tests=true + shift + ;; + -target) + if [[ $# -lt 2 ]]; then + echo "Error: -target requires a value" >&2 + exit 1 + fi + build_target="$2" + shift 2 + ;; + *) + common_args+=("$1") + shift + ;; + esac +done +set -- "${common_args[@]}" + # shellcheck source=ci/build_common.sh source "${ci_dir}/build_common.sh" print_environment_details PRESET="libcudacxx" +if $codegen_tests; then + PRESET="libcudacxx-codegen-filecheck" +fi + CMAKE_OPTIONS=("-DCMAKE_CXX_STANDARD=${CXX_STANDARD}" "-DCMAKE_CUDA_STANDARD=${CXX_STANDARD}") +if [[ -n "$build_target" ]]; then + configure_preset "$PRESET" "$PRESET" "${CMAKE_OPTIONS[@]}" + if ! $CONFIGURE_ONLY; then + build_preset "$PRESET" "$PRESET" --target "$build_target" + fi + print_time_summary + exit 0 +fi + upload_test_artifacts=false if [[ -n "${GITHUB_ACTIONS:-}" ]] && "${ci_dir}/util/workflow/has_consumers.sh"; then upload_test_artifacts=true diff --git a/ci/matrix.yaml b/ci/matrix.yaml index caa02f183525..cdb2549f5f88 100644 --- a/ci/matrix.yaml +++ b/ci/matrix.yaml @@ -110,6 +110,18 @@ workflows: - {jobs: ['test_gpu'], project: 'thrust', cmake_options: '-DTHRUST_DISPATCH_TYPE=Force32bit', gpu: 'rtx4090'} - {jobs: ['nvrtc'], project: 'libcudacxx', std: 'all', gpu: 'rtx2080', sm: 'gpu'} - {jobs: ['verify_codegen'], project: 'libcudacxx'} + # libcu++ Codegen FileCheck: PRs use one GCC host compiler per CTK. + # Suite-specific architectures are added separately. + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', ctk: '12.0', cxx: 'gcc12', codegen_target: ['atomics-ptx', 'atomics-sass'], sm: [75, 80, 90]} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', ctk: '12.X', cxx: 'gcc14', codegen_target: ['atomics-ptx', 'atomics-sass', 'simd-ptx', 'simd-sass'], sm: [80, 90, 100, 120]} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', ctk: '12.X', cxx: 'gcc14', codegen_target: ['atomics-ptx', 'atomics-sass'], sm: 75} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', ctk: '12.X', cxx: 'gcc14', codegen_target: 'simd-sass', sm: [103, '120f']} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', ctk: '13.0', cxx: 'gcc15', codegen_target: ['atomics-ptx', 'atomics-sass', 'simd-ptx', 'simd-sass'], sm: [80, 90, 100, 120]} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', ctk: '13.0', cxx: 'gcc15', codegen_target: ['atomics-ptx', 'atomics-sass'], sm: 75} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', ctk: '13.0', cxx: 'gcc15', codegen_target: 'simd-sass', sm: [103, '120f']} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', cxx: 'gcc15', codegen_target: ['atomics-ptx', 'atomics-sass', 'simd-ptx', 'simd-sass'], sm: [80, 90, 100, 120]} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', cxx: 'gcc15', codegen_target: ['atomics-ptx', 'atomics-sass'], sm: 75} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', cxx: 'gcc15', codegen_target: 'simd-sass', sm: [103, '120f']} # c.parallel -- pinned to gcc13 / msvc2022 to match python - {jobs: ['test'], project: 'cccl_c_parallel', ctk: '12.X', cxx: ['gcc13', 'msvc2022'], gpu: ['t4']} - {jobs: ['test'], project: 'cccl_c_parallel', ctk: '13.X', cxx: ['gcc13', 'msvc2022'], gpu: ['rtx2080', 'l4', 'h100']} @@ -284,6 +296,18 @@ workflows: # NVRTC tests don't currently support 12.0: - {jobs: ['nvrtc'], project: 'libcudacxx', ctk: [ '12.X', '13.0', '13.X'], cxx: 'gcc12', std: 'all', gpu: 'rtx2080', sm: 'gpu'} - {jobs: ['verify_codegen'], project: 'libcudacxx'} + # libcu++ Codegen FileCheck: nightly covers the GCC and Clang host compilers + # supported by each CTK. Suite-specific architectures are added separately. + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', ctk: '12.0', cxx: ['gcc12', 'clang14'], codegen_target: ['atomics-ptx', 'atomics-sass'], sm: [75, 80, 90]} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', ctk: '12.X', cxx: ['gcc14', 'clang19'], codegen_target: ['atomics-ptx', 'atomics-sass', 'simd-ptx', 'simd-sass'], sm: [80, 90, 100, 120]} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', ctk: '12.X', cxx: ['gcc14', 'clang19'], codegen_target: ['atomics-ptx', 'atomics-sass'], sm: 75} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', ctk: '12.X', cxx: ['gcc14', 'clang19'], codegen_target: 'simd-sass', sm: [103, '120f']} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', ctk: '13.0', cxx: ['gcc15', 'clang20'], codegen_target: ['atomics-ptx', 'atomics-sass', 'simd-ptx', 'simd-sass'], sm: [80, 90, 100, 120]} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', ctk: '13.0', cxx: ['gcc15', 'clang20'], codegen_target: ['atomics-ptx', 'atomics-sass'], sm: 75} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', ctk: '13.0', cxx: ['gcc15', 'clang20'], codegen_target: 'simd-sass', sm: [103, '120f']} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', cxx: ['gcc15', 'clang21'], codegen_target: ['atomics-ptx', 'atomics-sass', 'simd-ptx', 'simd-sass'], sm: [80, 90, 100, 120]} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', cxx: ['gcc15', 'clang21'], codegen_target: ['atomics-ptx', 'atomics-sass'], sm: 75} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', cxx: ['gcc15', 'clang21'], codegen_target: 'simd-sass', sm: [103, '120f']} # c.parallel -- pinned to gcc13 / msvc2022 to match python - {jobs: ['test'], project: ['cccl_c_parallel'], ctk: '12.X', cxx: ['gcc13', 'msvc2022'], gpu: ['t4']} - {jobs: ['test'], project: ['cccl_c_parallel'], ctk: '13.X', cxx: ['gcc13', 'msvc2022'], gpu: ['rtx2080', 'l4', 'h100']} @@ -393,6 +417,18 @@ workflows: # NVRTC tests don't currently support 12.0: - {jobs: ['nvrtc'], project: 'libcudacxx', ctk: [ '12.X', '13.0', '13.X'], cxx: 'gcc12', std: 'all', gpu: 'rtx2080', sm: 'gpu'} - {jobs: ['verify_codegen'], project: 'libcudacxx'} + # libcu++ Codegen FileCheck: weekly covers the GCC and Clang host compilers + # supported by each CTK. Suite-specific architectures are added separately. + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', ctk: '12.0', cxx: ['gcc12', 'clang14'], codegen_target: ['atomics-ptx', 'atomics-sass'], sm: [75, 80, 90]} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', ctk: '12.X', cxx: ['gcc14', 'clang19'], codegen_target: ['atomics-ptx', 'atomics-sass', 'simd-ptx', 'simd-sass'], sm: [80, 90, 100, 120]} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', ctk: '12.X', cxx: ['gcc14', 'clang19'], codegen_target: ['atomics-ptx', 'atomics-sass'], sm: 75} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', ctk: '12.X', cxx: ['gcc14', 'clang19'], codegen_target: 'simd-sass', sm: [103, '120f']} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', ctk: '13.0', cxx: ['gcc15', 'clang20'], codegen_target: ['atomics-ptx', 'atomics-sass', 'simd-ptx', 'simd-sass'], sm: [80, 90, 100, 120]} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', ctk: '13.0', cxx: ['gcc15', 'clang20'], codegen_target: ['atomics-ptx', 'atomics-sass'], sm: 75} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', ctk: '13.0', cxx: ['gcc15', 'clang20'], codegen_target: 'simd-sass', sm: [103, '120f']} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', cxx: ['gcc15', 'clang21'], codegen_target: ['atomics-ptx', 'atomics-sass', 'simd-ptx', 'simd-sass'], sm: [80, 90, 100, 120]} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', cxx: ['gcc15', 'clang21'], codegen_target: ['atomics-ptx', 'atomics-sass'], sm: 75} + - {jobs: ['codegen_filecheck'], project: 'libcudacxx', std: 'max', cxx: ['gcc15', 'clang21'], codegen_target: 'simd-sass', sm: [103, '120f']} # c.parallel -- pinned to gcc13 / msvc2022 to match python - {jobs: ['test'], project: ['cccl_c_parallel'], ctk: '12.X', cxx: ['gcc13', 'msvc2022'], gpu: ['t4']} - {jobs: ['test'], project: ['cccl_c_parallel'], ctk: '13.X', cxx: ['gcc13', 'msvc2022'], gpu: ['rtx2080', 'l4', 'h100']} @@ -607,6 +643,7 @@ jobs: # libcudacxx: nvrtc: { gpu: true, name: 'NVRTC' } verify_codegen: { gpu: false, name: 'VerifyCodegen' } + codegen_filecheck: { gpu: false, name: 'libcu++ Codegen FileCheck', invoke: { prefix: 'build', args: '-codegen-tests' } } # CUB: build_nolid: { name: 'BuildNoLaunch', gpu: false, invoke: { prefix: 'build', args: '-no-lid'} } @@ -656,6 +693,12 @@ jobs: dc: { gpu: false } dc_ext: { gpu: false, cuda_ext: true } +codegen_targets: + atomics-ptx: { name: 'Atomics PTX', cmake_target: 'libcudacxx.test.atomics.ptx' } + atomics-sass: { name: 'Atomics SASS', cmake_target: 'libcudacxx.test.atomics.sass' } + simd-ptx: { name: 'SIMD PTX', cmake_target: 'libcudacxx.test.simd.ptx' } + simd-sass: { name: 'SIMD SASS', cmake_target: 'libcudacxx.test.simd.sass' } + # Projects have the following properties: # # Keys are project subdirectories names. These will also be used in script names. @@ -783,9 +826,11 @@ gpus: # - required: Whether the tag is required. Default is false. # - default: The default value for the tag. Default is null. tags: - # An array of jobs (e.g. 'build', 'test', 'nvrtc', 'infra', 'verify_codegen', ...) + # An array of jobs (e.g. 'build', 'test', 'nvrtc', 'infra', 'verify_codegen', 'codegen_filecheck', ...) # See the `jobs` map. jobs: { required: true } + # FileCheck codegen suite selected by a codegen_filecheck job. + codegen_target: { required: false } # CUDA ToolKit version # See the `ctks` map. ctk: { default: '13.X' } diff --git a/libcudacxx/include/cuda/__atomic/atomic.h b/libcudacxx/include/cuda/__atomic/atomic.h index 5ef6fcecec7f..f79ce1dc632e 100644 --- a/libcudacxx/include/cuda/__atomic/atomic.h +++ b/libcudacxx/include/cuda/__atomic/atomic.h @@ -22,6 +22,7 @@ #endif // no system header #include +#include #include #include @@ -80,7 +81,7 @@ struct atomic : public ::cuda::std::__atomic_impl<_Tp, _Sco> template struct atomic_ref : public ::cuda::std::__atomic_ref_impl<_Tp, _Sco> { - using value_type = _Tp; + using value_type = ::cuda::std::remove_cv_t<_Tp>; static constexpr size_t required_alignment = sizeof(_Tp); @@ -90,7 +91,7 @@ struct atomic_ref : public ::cuda::std::__atomic_ref_impl<_Tp, _Sco> : ::cuda::std::__atomic_ref_impl<_Tp, _Sco>(__ref) {} - _CCCL_HOST_DEVICE_API inline _Tp operator=(_Tp __v) const noexcept + _CCCL_HOST_DEVICE_API inline value_type operator=(value_type __v) const noexcept { this->store(__v); return __v; @@ -105,12 +106,14 @@ struct atomic_ref : public ::cuda::std::__atomic_ref_impl<_Tp, _Sco> atomic_ref& operator=(const atomic_ref&) = delete; atomic_ref& operator=(const atomic_ref&) const = delete; - _CCCL_HOST_DEVICE_API inline _Tp fetch_max(const _Tp& __op, memory_order __m = memory_order_seq_cst) const noexcept + _CCCL_HOST_DEVICE_API inline value_type + fetch_max(const value_type& __op, memory_order __m = memory_order_seq_cst) const noexcept { return ::cuda::std::__atomic_fetch_max_dispatch(&this->__a, __op, __m, ::cuda::std::__scope_to_tag<_Sco>{}); } - _CCCL_HOST_DEVICE_API inline _Tp fetch_min(const _Tp& __op, memory_order __m = memory_order_seq_cst) const noexcept + _CCCL_HOST_DEVICE_API inline value_type + fetch_min(const value_type& __op, memory_order __m = memory_order_seq_cst) const noexcept { return ::cuda::std::__atomic_fetch_min_dispatch(&this->__a, __op, __m, ::cuda::std::__scope_to_tag<_Sco>{}); } diff --git a/libcudacxx/include/cuda/std/__atomic/api/common.h b/libcudacxx/include/cuda/std/__atomic/api/common.h index 04a3135bc488..3150a68fa1f9 100644 --- a/libcudacxx/include/cuda/std/__atomic/api/common.h +++ b/libcudacxx/include/cuda/std/__atomic/api/common.h @@ -22,43 +22,44 @@ #endif // no system header #include +#include // API definitions for the base atomic implementation -#define _LIBCUDACXX_ATOMIC_COMMON_IMPL(_CONST, _VOLATILE) \ +#define _LIBCUDACXX_ATOMIC_COMMON_IMPL(_CONST, _VOLATILE, _VT) \ _CCCL_HOST_DEVICE_API inline bool is_lock_free() const _VOLATILE noexcept \ { \ return _LIBCUDACXX_ATOMIC_IS_LOCK_FREE(sizeof(_Tp)); \ } \ - _CCCL_HOST_DEVICE_API inline void store(_Tp __d, memory_order __m = memory_order_seq_cst) \ + _CCCL_HOST_DEVICE_API inline void store(_VT __d, memory_order __m = memory_order_seq_cst) \ _CONST _VOLATILE noexcept _LIBCUDACXX_CHECK_STORE_MEMORY_ORDER(__m) \ { \ __atomic_store_dispatch(&__a, __d, __m, _Sco{}); \ } \ - _CCCL_HOST_DEVICE_API inline _Tp load(memory_order __m = memory_order_seq_cst) \ + _CCCL_HOST_DEVICE_API inline _VT load(memory_order __m = memory_order_seq_cst) \ const _VOLATILE noexcept _LIBCUDACXX_CHECK_LOAD_MEMORY_ORDER(__m) \ { \ return __atomic_load_dispatch(&__a, __m, _Sco{}); \ } \ - _CCCL_HOST_DEVICE_API inline operator _Tp() const _VOLATILE noexcept \ + _CCCL_HOST_DEVICE_API inline operator _VT() const _VOLATILE noexcept \ { \ return load(); \ } \ - _CCCL_HOST_DEVICE_API inline _Tp exchange(_Tp __d, memory_order __m = memory_order_seq_cst) \ + _CCCL_HOST_DEVICE_API inline _VT exchange(_VT __d, memory_order __m = memory_order_seq_cst) \ _CONST _VOLATILE noexcept \ { \ return __atomic_exchange_dispatch(&__a, __d, __m, _Sco{}); \ } \ - _CCCL_HOST_DEVICE_API inline bool compare_exchange_weak(_Tp& __e, _Tp __d, memory_order __s, memory_order __f) \ + _CCCL_HOST_DEVICE_API inline bool compare_exchange_weak(_VT& __e, _VT __d, memory_order __s, memory_order __f) \ _CONST _VOLATILE noexcept _LIBCUDACXX_CHECK_EXCHANGE_MEMORY_ORDER(__s, __f) \ { \ return __atomic_compare_exchange_weak_dispatch(&__a, &__e, __d, __s, __f, _Sco{}); \ } \ - _CCCL_HOST_DEVICE_API inline bool compare_exchange_strong(_Tp& __e, _Tp __d, memory_order __s, memory_order __f) \ + _CCCL_HOST_DEVICE_API inline bool compare_exchange_strong(_VT& __e, _VT __d, memory_order __s, memory_order __f) \ _CONST _VOLATILE noexcept _LIBCUDACXX_CHECK_EXCHANGE_MEMORY_ORDER(__s, __f) \ { \ return __atomic_compare_exchange_strong_dispatch(&__a, &__e, __d, __s, __f, _Sco{}); \ } \ - _CCCL_HOST_DEVICE_API inline bool compare_exchange_weak(_Tp& __e, _Tp __d, memory_order __m = memory_order_seq_cst) \ + _CCCL_HOST_DEVICE_API inline bool compare_exchange_weak(_VT& __e, _VT __d, memory_order __m = memory_order_seq_cst) \ _CONST _VOLATILE noexcept \ { \ if (memory_order_acq_rel == __m) \ @@ -69,7 +70,7 @@ return __atomic_compare_exchange_weak_dispatch(&__a, &__e, __d, __m, __m, _Sco{}); \ } \ _CCCL_HOST_DEVICE_API inline bool compare_exchange_strong( \ - _Tp& __e, _Tp __d, memory_order __m = memory_order_seq_cst) _CONST _VOLATILE noexcept \ + _VT& __e, _VT __d, memory_order __m = memory_order_seq_cst) _CONST _VOLATILE noexcept \ { \ if (memory_order_acq_rel == __m) \ return __atomic_compare_exchange_strong_dispatch(&__a, &__e, __d, __m, memory_order_acquire, _Sco{}); \ @@ -78,7 +79,7 @@ else \ return __atomic_compare_exchange_strong_dispatch(&__a, &__e, __d, __m, __m, _Sco{}); \ } \ - _CCCL_HOST_DEVICE_API inline void wait(_Tp __v, memory_order __m = memory_order_seq_cst) const _VOLATILE noexcept \ + _CCCL_HOST_DEVICE_API inline void wait(_VT __v, memory_order __m = memory_order_seq_cst) const _VOLATILE noexcept \ { \ __atomic_wait(&__a, __v, __m, _Sco{}); \ } \ @@ -92,105 +93,105 @@ } // API definitions for arithmetic atomics -#define _LIBCUDACXX_ATOMIC_ARITHMETIC_IMPL(_CONST, _VOLATILE) \ - _CCCL_HOST_DEVICE_API inline _Tp fetch_add(_Tp __op, memory_order __m = memory_order_seq_cst) \ +#define _LIBCUDACXX_ATOMIC_ARITHMETIC_IMPL(_CONST, _VOLATILE, _VT) \ + _CCCL_HOST_DEVICE_API inline _VT fetch_add(_VT __op, memory_order __m = memory_order_seq_cst) \ _CONST _VOLATILE noexcept \ { \ return __atomic_fetch_add_dispatch(&__a, __op, __m, _Sco{}); \ } \ - _CCCL_HOST_DEVICE_API inline _Tp fetch_sub(_Tp __op, memory_order __m = memory_order_seq_cst) \ + _CCCL_HOST_DEVICE_API inline _VT fetch_sub(_VT __op, memory_order __m = memory_order_seq_cst) \ _CONST _VOLATILE noexcept \ { \ return __atomic_fetch_sub_dispatch(&__a, __op, __m, _Sco{}); \ } \ - _CCCL_HOST_DEVICE_API inline _Tp operator++(int) _CONST _VOLATILE noexcept \ + _CCCL_HOST_DEVICE_API inline _VT operator++(int) _CONST _VOLATILE noexcept \ { \ - return fetch_add(_Tp(1)); \ + return fetch_add(_VT(1)); \ } \ - _CCCL_HOST_DEVICE_API inline _Tp operator--(int) _CONST _VOLATILE noexcept \ + _CCCL_HOST_DEVICE_API inline _VT operator--(int) _CONST _VOLATILE noexcept \ { \ - return fetch_sub(_Tp(1)); \ + return fetch_sub(_VT(1)); \ } \ - _CCCL_HOST_DEVICE_API inline _Tp operator++() _CONST _VOLATILE noexcept \ + _CCCL_HOST_DEVICE_API inline _VT operator++() _CONST _VOLATILE noexcept \ { \ - return fetch_add(_Tp(1)) + _Tp(1); \ + return fetch_add(_VT(1)) + _VT(1); \ } \ - _CCCL_HOST_DEVICE_API inline _Tp operator--() _CONST _VOLATILE noexcept \ + _CCCL_HOST_DEVICE_API inline _VT operator--() _CONST _VOLATILE noexcept \ { \ - return fetch_sub(_Tp(1)) - _Tp(1); \ + return fetch_sub(_VT(1)) - _VT(1); \ } \ - _CCCL_HOST_DEVICE_API inline _Tp operator+=(_Tp __op) _CONST _VOLATILE noexcept \ + _CCCL_HOST_DEVICE_API inline _VT operator+=(_VT __op) _CONST _VOLATILE noexcept \ { \ return fetch_add(__op) + __op; \ } \ - _CCCL_HOST_DEVICE_API inline _Tp operator-=(_Tp __op) _CONST _VOLATILE noexcept \ + _CCCL_HOST_DEVICE_API inline _VT operator-=(_VT __op) _CONST _VOLATILE noexcept \ { \ return fetch_sub(__op) - __op; \ } // API definitions for bitwise atomics -#define _LIBCUDACXX_ATOMIC_BITWISE_IMPL(_CONST, _VOLATILE) \ - _CCCL_HOST_DEVICE_API inline _Tp fetch_and(_Tp __op, memory_order __m = memory_order_seq_cst) \ +#define _LIBCUDACXX_ATOMIC_BITWISE_IMPL(_CONST, _VOLATILE, _VT) \ + _CCCL_HOST_DEVICE_API inline _VT fetch_and(_VT __op, memory_order __m = memory_order_seq_cst) \ _CONST _VOLATILE noexcept \ { \ return __atomic_fetch_and_dispatch(&__a, __op, __m, _Sco{}); \ } \ - _CCCL_HOST_DEVICE_API inline _Tp fetch_or(_Tp __op, memory_order __m = memory_order_seq_cst) \ + _CCCL_HOST_DEVICE_API inline _VT fetch_or(_VT __op, memory_order __m = memory_order_seq_cst) \ _CONST _VOLATILE noexcept \ { \ return __atomic_fetch_or_dispatch(&__a, __op, __m, _Sco{}); \ } \ - _CCCL_HOST_DEVICE_API inline _Tp fetch_xor(_Tp __op, memory_order __m = memory_order_seq_cst) \ + _CCCL_HOST_DEVICE_API inline _VT fetch_xor(_VT __op, memory_order __m = memory_order_seq_cst) \ _CONST _VOLATILE noexcept \ { \ return __atomic_fetch_xor_dispatch(&__a, __op, __m, _Sco{}); \ } \ - _CCCL_HOST_DEVICE_API inline _Tp operator&=(_Tp __op) _CONST _VOLATILE noexcept \ + _CCCL_HOST_DEVICE_API inline _VT operator&=(_VT __op) _CONST _VOLATILE noexcept \ { \ return fetch_and(__op) & __op; \ } \ - _CCCL_HOST_DEVICE_API inline _Tp operator|=(_Tp __op) _CONST _VOLATILE noexcept \ + _CCCL_HOST_DEVICE_API inline _VT operator|=(_VT __op) _CONST _VOLATILE noexcept \ { \ return fetch_or(__op) | __op; \ } \ - _CCCL_HOST_DEVICE_API inline _Tp operator^=(_Tp __op) _CONST _VOLATILE noexcept \ + _CCCL_HOST_DEVICE_API inline _VT operator^=(_VT __op) _CONST _VOLATILE noexcept \ { \ return fetch_xor(__op) ^ __op; \ } // API definitions for atomics with pointers -#define _LIBCUDACXX_ATOMIC_POINTER_IMPL(_CONST, _VOLATILE) \ - _CCCL_HOST_DEVICE_API inline _Tp fetch_add(ptrdiff_t __op, memory_order __m = memory_order_seq_cst) \ +#define _LIBCUDACXX_ATOMIC_POINTER_IMPL(_CONST, _VOLATILE, _VT) \ + _CCCL_HOST_DEVICE_API inline _VT fetch_add(ptrdiff_t __op, memory_order __m = memory_order_seq_cst) \ _CONST _VOLATILE noexcept \ { \ return __atomic_fetch_add_dispatch(&__a, __op, __m, __thread_scope_system_tag{}); \ } \ - _CCCL_HOST_DEVICE_API inline _Tp fetch_sub(ptrdiff_t __op, memory_order __m = memory_order_seq_cst) \ + _CCCL_HOST_DEVICE_API inline _VT fetch_sub(ptrdiff_t __op, memory_order __m = memory_order_seq_cst) \ _CONST _VOLATILE noexcept \ { \ return __atomic_fetch_sub_dispatch(&__a, __op, __m, __thread_scope_system_tag{}); \ } \ - _CCCL_HOST_DEVICE_API inline _Tp operator++(int) _CONST _VOLATILE noexcept \ + _CCCL_HOST_DEVICE_API inline _VT operator++(int) _CONST _VOLATILE noexcept \ { \ return fetch_add(1); \ } \ - _CCCL_HOST_DEVICE_API inline _Tp operator--(int) _CONST _VOLATILE noexcept \ + _CCCL_HOST_DEVICE_API inline _VT operator--(int) _CONST _VOLATILE noexcept \ { \ return fetch_sub(1); \ } \ - _CCCL_HOST_DEVICE_API inline _Tp operator++() _CONST _VOLATILE noexcept \ + _CCCL_HOST_DEVICE_API inline _VT operator++() _CONST _VOLATILE noexcept \ { \ return fetch_add(1) + 1; \ } \ - _CCCL_HOST_DEVICE_API inline _Tp operator--() _CONST _VOLATILE noexcept \ + _CCCL_HOST_DEVICE_API inline _VT operator--() _CONST _VOLATILE noexcept \ { \ return fetch_sub(1) - 1; \ } \ - _CCCL_HOST_DEVICE_API inline _Tp operator+=(ptrdiff_t __op) _CONST _VOLATILE noexcept \ + _CCCL_HOST_DEVICE_API inline _VT operator+=(ptrdiff_t __op) _CONST _VOLATILE noexcept \ { \ return fetch_add(__op) + __op; \ } \ - _CCCL_HOST_DEVICE_API inline _Tp operator-=(ptrdiff_t __op) _CONST _VOLATILE noexcept \ + _CCCL_HOST_DEVICE_API inline _VT operator-=(ptrdiff_t __op) _CONST _VOLATILE noexcept \ { \ return fetch_sub(__op) - __op; \ } diff --git a/libcudacxx/include/cuda/std/__atomic/api/owned.h b/libcudacxx/include/cuda/std/__atomic/api/owned.h index 4089d4c658ef..2ddd84cbb5c1 100644 --- a/libcudacxx/include/cuda/std/__atomic/api/owned.h +++ b/libcudacxx/include/cuda/std/__atomic/api/owned.h @@ -48,8 +48,8 @@ struct __atomic_common static constexpr bool is_always_lock_free = _CCCL_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), nullptr); #endif // defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE) - _LIBCUDACXX_ATOMIC_COMMON_IMPL(, ) - _LIBCUDACXX_ATOMIC_COMMON_IMPL(, volatile) + _LIBCUDACXX_ATOMIC_COMMON_IMPL(, , _Tp) + _LIBCUDACXX_ATOMIC_COMMON_IMPL(, volatile, _Tp) }; template @@ -67,11 +67,11 @@ struct __atomic_arithmetic static constexpr bool is_always_lock_free = _CCCL_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), nullptr); #endif // defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE) - _LIBCUDACXX_ATOMIC_COMMON_IMPL(, ) - _LIBCUDACXX_ATOMIC_COMMON_IMPL(, volatile) + _LIBCUDACXX_ATOMIC_COMMON_IMPL(, , _Tp) + _LIBCUDACXX_ATOMIC_COMMON_IMPL(, volatile, _Tp) - _LIBCUDACXX_ATOMIC_ARITHMETIC_IMPL(, ) - _LIBCUDACXX_ATOMIC_ARITHMETIC_IMPL(, volatile) + _LIBCUDACXX_ATOMIC_ARITHMETIC_IMPL(, , _Tp) + _LIBCUDACXX_ATOMIC_ARITHMETIC_IMPL(, volatile, _Tp) }; template @@ -89,14 +89,14 @@ struct __atomic_bitwise static constexpr bool is_always_lock_free = _CCCL_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), nullptr); #endif // defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE) - _LIBCUDACXX_ATOMIC_COMMON_IMPL(, ) - _LIBCUDACXX_ATOMIC_COMMON_IMPL(, volatile) + _LIBCUDACXX_ATOMIC_COMMON_IMPL(, , _Tp) + _LIBCUDACXX_ATOMIC_COMMON_IMPL(, volatile, _Tp) - _LIBCUDACXX_ATOMIC_ARITHMETIC_IMPL(, ) - _LIBCUDACXX_ATOMIC_ARITHMETIC_IMPL(, volatile) + _LIBCUDACXX_ATOMIC_ARITHMETIC_IMPL(, , _Tp) + _LIBCUDACXX_ATOMIC_ARITHMETIC_IMPL(, volatile, _Tp) - _LIBCUDACXX_ATOMIC_BITWISE_IMPL(, ) - _LIBCUDACXX_ATOMIC_BITWISE_IMPL(, volatile) + _LIBCUDACXX_ATOMIC_BITWISE_IMPL(, , _Tp) + _LIBCUDACXX_ATOMIC_BITWISE_IMPL(, volatile, _Tp) }; template @@ -114,11 +114,11 @@ struct __atomic_pointer static constexpr bool is_always_lock_free = _CCCL_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), nullptr); #endif // defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE) - _LIBCUDACXX_ATOMIC_COMMON_IMPL(, ) - _LIBCUDACXX_ATOMIC_COMMON_IMPL(, volatile) + _LIBCUDACXX_ATOMIC_COMMON_IMPL(, , _Tp) + _LIBCUDACXX_ATOMIC_COMMON_IMPL(, volatile, _Tp) - _LIBCUDACXX_ATOMIC_POINTER_IMPL(, ) - _LIBCUDACXX_ATOMIC_POINTER_IMPL(, volatile) + _LIBCUDACXX_ATOMIC_POINTER_IMPL(, , _Tp) + _LIBCUDACXX_ATOMIC_POINTER_IMPL(, volatile, _Tp) }; template diff --git a/libcudacxx/include/cuda/std/__atomic/api/reference.h b/libcudacxx/include/cuda/std/__atomic/api/reference.h index fa4e74326afa..446df2f26c9d 100644 --- a/libcudacxx/include/cuda/std/__atomic/api/reference.h +++ b/libcudacxx/include/cuda/std/__atomic/api/reference.h @@ -46,7 +46,7 @@ struct __atomic_ref_common static constexpr bool is_always_lock_free = _CCCL_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), nullptr); #endif // defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE) - _LIBCUDACXX_ATOMIC_COMMON_IMPL(const, ) + _LIBCUDACXX_ATOMIC_COMMON_IMPL(const, , remove_cv_t<_Tp>) }; template @@ -62,8 +62,8 @@ struct __atomic_ref_arithmetic static constexpr bool is_always_lock_free = _CCCL_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), nullptr); #endif // defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE) - _LIBCUDACXX_ATOMIC_COMMON_IMPL(const, ) - _LIBCUDACXX_ATOMIC_ARITHMETIC_IMPL(const, ) + _LIBCUDACXX_ATOMIC_COMMON_IMPL(const, , remove_cv_t<_Tp>) + _LIBCUDACXX_ATOMIC_ARITHMETIC_IMPL(const, , remove_cv_t<_Tp>) }; template @@ -79,9 +79,9 @@ struct __atomic_ref_bitwise static constexpr bool is_always_lock_free = _CCCL_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), nullptr); #endif // defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE) - _LIBCUDACXX_ATOMIC_COMMON_IMPL(const, ) - _LIBCUDACXX_ATOMIC_ARITHMETIC_IMPL(const, ) - _LIBCUDACXX_ATOMIC_BITWISE_IMPL(const, ) + _LIBCUDACXX_ATOMIC_COMMON_IMPL(const, , remove_cv_t<_Tp>) + _LIBCUDACXX_ATOMIC_ARITHMETIC_IMPL(const, , remove_cv_t<_Tp>) + _LIBCUDACXX_ATOMIC_BITWISE_IMPL(const, , remove_cv_t<_Tp>) }; template @@ -97,8 +97,8 @@ struct __atomic_ref_pointer static constexpr bool is_always_lock_free = _CCCL_ATOMIC_ALWAYS_LOCK_FREE(sizeof(_Tp), nullptr); #endif // defined(_CCCL_ATOMIC_ALWAYS_LOCK_FREE) - _LIBCUDACXX_ATOMIC_COMMON_IMPL(const, ) - _LIBCUDACXX_ATOMIC_POINTER_IMPL(const, ) + _LIBCUDACXX_ATOMIC_COMMON_IMPL(const, , remove_cv_t<_Tp>) + _LIBCUDACXX_ATOMIC_POINTER_IMPL(const, , remove_cv_t<_Tp>) }; template diff --git a/libcudacxx/include/cuda/std/__atomic/types/base.h b/libcudacxx/include/cuda/std/__atomic/types/base.h index 10f840647434..957f5e79a72c 100644 --- a/libcudacxx/include/cuda/std/__atomic/types/base.h +++ b/libcudacxx/include/cuda/std/__atomic/types/base.h @@ -100,7 +100,7 @@ _CCCL_HOST_DEVICE_API void __atomic_store_dispatch(_Sto* __a, _Up __val, memory_ template = 0> _CCCL_HOST_DEVICE_API auto __atomic_load_dispatch(const _Sto* __a, memory_order __order, _Sco = {}) - -> __atomic_underlying_t<_Sto> + -> __atomic_underlying_remove_cv_t<_Sto> { NV_DISPATCH_TARGET( NV_IS_DEVICE, @@ -111,7 +111,7 @@ _CCCL_HOST_DEVICE_API auto __atomic_load_dispatch(const _Sto* __a, memory_order template = 0> _CCCL_HOST_DEVICE_API auto __atomic_exchange_dispatch(_Sto* __a, _Up __value, memory_order __order, _Sco = {}) - -> __atomic_underlying_t<_Sto> + -> __atomic_underlying_remove_cv_t<_Sto> { NV_DISPATCH_TARGET( NV_IS_DEVICE, @@ -162,7 +162,7 @@ _CCCL_HOST_DEVICE_API bool __atomic_compare_exchange_weak_dispatch( template = 0> _CCCL_HOST_DEVICE_API auto __atomic_fetch_add_dispatch(_Sto* __a, _Up __delta, memory_order __order, _Sco = {}) - -> __atomic_underlying_t<_Sto> + -> __atomic_underlying_remove_cv_t<_Sto> { NV_DISPATCH_TARGET( NV_IS_DEVICE, @@ -173,7 +173,7 @@ _CCCL_HOST_DEVICE_API auto __atomic_fetch_add_dispatch(_Sto* __a, _Up __delta, m template = 0> _CCCL_HOST_DEVICE_API auto __atomic_fetch_sub_dispatch(_Sto* __a, _Up __delta, memory_order __order, _Sco = {}) - -> __atomic_underlying_t<_Sto> + -> __atomic_underlying_remove_cv_t<_Sto> { NV_DISPATCH_TARGET( NV_IS_DEVICE, @@ -184,7 +184,7 @@ _CCCL_HOST_DEVICE_API auto __atomic_fetch_sub_dispatch(_Sto* __a, _Up __delta, m template = 0> _CCCL_HOST_DEVICE_API auto __atomic_fetch_and_dispatch(_Sto* __a, _Up __pattern, memory_order __order, _Sco = {}) - -> __atomic_underlying_t<_Sto> + -> __atomic_underlying_remove_cv_t<_Sto> { NV_DISPATCH_TARGET( NV_IS_DEVICE, @@ -195,7 +195,7 @@ _CCCL_HOST_DEVICE_API auto __atomic_fetch_and_dispatch(_Sto* __a, _Up __pattern, template = 0> _CCCL_HOST_DEVICE_API auto __atomic_fetch_or_dispatch(_Sto* __a, _Up __pattern, memory_order __order, _Sco = {}) - -> __atomic_underlying_t<_Sto> + -> __atomic_underlying_remove_cv_t<_Sto> { NV_DISPATCH_TARGET( NV_IS_DEVICE, @@ -206,7 +206,7 @@ _CCCL_HOST_DEVICE_API auto __atomic_fetch_or_dispatch(_Sto* __a, _Up __pattern, template = 0> _CCCL_HOST_DEVICE_API auto __atomic_fetch_xor_dispatch(_Sto* __a, _Up __pattern, memory_order __order, _Sco = {}) - -> __atomic_underlying_t<_Sto> + -> __atomic_underlying_remove_cv_t<_Sto> { NV_DISPATCH_TARGET( NV_IS_DEVICE, @@ -217,7 +217,7 @@ _CCCL_HOST_DEVICE_API auto __atomic_fetch_xor_dispatch(_Sto* __a, _Up __pattern, template = 0> _CCCL_HOST_DEVICE_API auto __atomic_fetch_max_dispatch(_Sto* __a, _Up __val, memory_order __order, _Sco = {}) - -> __atomic_underlying_t<_Sto> + -> __atomic_underlying_remove_cv_t<_Sto> { NV_IF_TARGET( NV_IS_DEVICE, @@ -227,7 +227,7 @@ _CCCL_HOST_DEVICE_API auto __atomic_fetch_max_dispatch(_Sto* __a, _Up __val, mem template = 0> _CCCL_HOST_DEVICE_API auto __atomic_fetch_min_dispatch(_Sto* __a, _Up __val, memory_order __order, _Sco = {}) - -> __atomic_underlying_t<_Sto> + -> __atomic_underlying_remove_cv_t<_Sto> { NV_IF_TARGET( NV_IS_DEVICE, diff --git a/libcudacxx/include/cuda/std/__atomic/types/reference.h b/libcudacxx/include/cuda/std/__atomic/types/reference.h index 6a0e8c29044f..1b1b0e67c64d 100644 --- a/libcudacxx/include/cuda/std/__atomic/types/reference.h +++ b/libcudacxx/include/cuda/std/__atomic/types/reference.h @@ -23,6 +23,7 @@ #include #include +#include #include @@ -36,7 +37,9 @@ struct __atomic_ref_storage static constexpr __atomic_tag __tag = __atomic_tag::__atomic_base_tag; #if !_CCCL_COMPILER(GCC) || _CCCL_COMPILER(GCC, >=, 5) - static_assert(is_trivially_copyable_v<_Tp>, "std::atomic_ref requires that 'Tp' be a trivially copyable type"); + // GCC 7 cannot correctly evaluate is_trivially_copyable_v for volatile-qualified types. + static_assert(is_trivially_copyable_v>, + "std::atomic_ref requires that 'Tp' be a trivially copyable type"); #endif static_assert(sizeof(__underlying_t) <= 16, "cuda::std::atomic_ref only supports sizeof(Tp) <= 16"); diff --git a/libcudacxx/include/cuda/std/atomic b/libcudacxx/include/cuda/std/atomic index 81d9771850e3..d51135fdb486 100644 --- a/libcudacxx/include/cuda/std/atomic +++ b/libcudacxx/include/cuda/std/atomic @@ -45,6 +45,7 @@ #include #include #include +#include // clang-format on @@ -92,7 +93,7 @@ struct atomic : public __atomic_impl<_Tp> template struct atomic_ref : public __atomic_ref_impl<_Tp> { - using value_type = _Tp; + using value_type = remove_cv_t<_Tp>; static constexpr size_t required_alignment = sizeof(_Tp); @@ -102,7 +103,7 @@ struct atomic_ref : public __atomic_ref_impl<_Tp> : __atomic_ref_impl<_Tp>(__ref) {} - _CCCL_HOST_DEVICE_API inline _Tp operator=(_Tp __v) const noexcept + _CCCL_HOST_DEVICE_API inline value_type operator=(value_type __v) const noexcept { this->store(__v); return __v; diff --git a/libcudacxx/test/CMakeLists.txt b/libcudacxx/test/CMakeLists.txt index f1fb3340ea99..c3718f6310e2 100644 --- a/libcudacxx/test/CMakeLists.txt +++ b/libcudacxx/test/CMakeLists.txt @@ -97,7 +97,69 @@ if (LIBCUDACXX_TEST_WITH_NVRTC) endif() add_subdirectory(nvtarget) -include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/CodegenTest.cmake") -add_subdirectory(atomic_codegen) -add_subdirectory(simd_codegen) + +set( + LIBCUDACXX_CODEGEN_FILECHECK_TESTS + "" + CACHE STRING + "Semicolon-separated libcu++ FileCheck codegen suites to enable (all, atomics-ptx, atomics-sass, simd-ptx, simd-sass)." +) +set( + libcudacxx_codegen_filecheck_valid_suites + all + atomics-ptx + atomics-sass + simd-ptx + simd-sass +) +foreach (suite IN LISTS LIBCUDACXX_CODEGEN_FILECHECK_TESTS) + if (NOT suite IN_LIST libcudacxx_codegen_filecheck_valid_suites) + message( + FATAL_ERROR + "Unknown libcu++ FileCheck codegen suite '${suite}'. Expected one of: ${libcudacxx_codegen_filecheck_valid_suites}" + ) + endif() +endforeach() + +if (LIBCUDACXX_CODEGEN_FILECHECK_TESTS) + include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/CodegenTest.cmake") + + if ( + all IN_LIST LIBCUDACXX_CODEGEN_FILECHECK_TESTS + OR atomics-ptx IN_LIST LIBCUDACXX_CODEGEN_FILECHECK_TESTS + ) + set(libcudacxx_codegen_filecheck_atomics_ptx ON) + endif() + if ( + all IN_LIST LIBCUDACXX_CODEGEN_FILECHECK_TESTS + OR atomics-sass IN_LIST LIBCUDACXX_CODEGEN_FILECHECK_TESTS + ) + set(libcudacxx_codegen_filecheck_atomics_sass ON) + endif() + if ( + all IN_LIST LIBCUDACXX_CODEGEN_FILECHECK_TESTS + OR simd-ptx IN_LIST LIBCUDACXX_CODEGEN_FILECHECK_TESTS + ) + set(libcudacxx_codegen_filecheck_simd_ptx ON) + endif() + if ( + all IN_LIST LIBCUDACXX_CODEGEN_FILECHECK_TESTS + OR simd-sass IN_LIST LIBCUDACXX_CODEGEN_FILECHECK_TESTS + ) + set(libcudacxx_codegen_filecheck_simd_sass ON) + endif() + + if ( + libcudacxx_codegen_filecheck_atomics_ptx + OR libcudacxx_codegen_filecheck_atomics_sass + ) + add_subdirectory(atomic_codegen) + endif() + if ( + libcudacxx_codegen_filecheck_simd_ptx + OR libcudacxx_codegen_filecheck_simd_sass + ) + add_subdirectory(simd_codegen) + endif() +endif() add_subdirectory(debugging) diff --git a/libcudacxx/test/atomic_codegen/CMakeLists.txt b/libcudacxx/test/atomic_codegen/CMakeLists.txt index 5ca26699b39a..9229a20a9795 100644 --- a/libcudacxx/test/atomic_codegen/CMakeLists.txt +++ b/libcudacxx/test/atomic_codegen/CMakeLists.txt @@ -1,26 +1,29 @@ -add_custom_target(libcudacxx.test.atomics.ptx) +if (libcudacxx_codegen_filecheck_atomics_ptx) + add_custom_target(libcudacxx.test.atomics.ptx) -set(atomic_codegen_cuda_arch 80) + set(libcudacxx_atomic_codegen_tests) + if (NOT "NVHPC" STREQUAL "${CMAKE_CXX_COMPILER_ID}") + file(GLOB libcudacxx_atomic_codegen_tests CONFIGURE_DEPENDS "*.cu") + endif() -set(libcudacxx_atomic_codegen_tests) -if (NOT "NVHPC" STREQUAL "${CMAKE_CXX_COMPILER_ID}") - file(GLOB libcudacxx_atomic_codegen_tests "*.cu") + if (libcudacxx_atomic_codegen_tests) + libcudacxx_codegen_check_tools(libcudacxx_codegen_tests_enabled) + if (libcudacxx_codegen_tests_enabled) + libcudacxx_codegen_get_cuda_architectures(atomic_codegen_cuda_archs) + foreach (arch IN LISTS atomic_codegen_cuda_archs) + libcudacxx_codegen_add_ptx_tests( + AGGREGATE_TARGET libcudacxx.test.atomics.ptx + TARGET_PREFIX atomic_codegen + ARCH "${arch}" + CHECK_PREFIXES SMXX + TESTS ${libcudacxx_atomic_codegen_tests} + COMPILE_DEFINITIONS _CCCL_ATOMIC_UNSAFE_AUTOMATIC_STORAGE=1 + ) + endforeach() + endif() + endif() endif() -if (NOT libcudacxx_atomic_codegen_tests) - return() +if (libcudacxx_codegen_filecheck_atomics_sass) + add_subdirectory(sass) endif() - -libcudacxx_codegen_check_tools(libcudacxx_codegen_tests_enabled) -if (NOT libcudacxx_codegen_tests_enabled) - return() -endif() - -libcudacxx_codegen_add_ptx_tests( - AGGREGATE_TARGET libcudacxx.test.atomics.ptx - TARGET_PREFIX atomic_codegen - ARCH "${atomic_codegen_cuda_arch}" - CHECK_PREFIXES SM8X - TESTS ${libcudacxx_atomic_codegen_tests} - COMPILE_DEFINITIONS _CCCL_ATOMIC_UNSAFE_AUTOMATIC_STORAGE=1 -) diff --git a/libcudacxx/test/atomic_codegen/atomic_add_non_volatile.cu b/libcudacxx/test/atomic_codegen/atomic_add_non_volatile.cu index 4489e8942680..226a9fc94914 100644 --- a/libcudacxx/test/atomic_codegen/atomic_add_non_volatile.cu +++ b/libcudacxx/test/atomic_codegen/atomic_add_non_volatile.cu @@ -8,14 +8,14 @@ __global__ void add_relaxed_device_non_volatile(int* data, int* out, int n) /* -; SM8X-LABEL: .target sm_80 -; SM8X: .visible .entry [[FUNCTION:_.*add_relaxed_device_non_volatile.*]]( -; SM8X-DAG: ld.param.{{b|u}}64 %rd[[#ATOM:]], {{.*}}[[FUNCTION]]_param_0{{.*}} -; SM8X-DAG: ld.param.{{b|u}}64 %rd[[#RESULT:]], {{.*}}[[FUNCTION]]_param_1{{.*}} -; SM8X-DAG: ld.param.{{b|u}}32 %r[[#INPUT:]], {{.*}}[[FUNCTION]]_param_2{{.*}} -; SM8X-DAG: cvta.to.global.u64 %rd[[#GOUT:]], %rd[[#RESULT]]; -; SM8X-NEXT: {{/*[[:space:]] *}}atom.add.relaxed.gpu.s32 %r[[#DEST:]],[%rd[[#ATOM]]],%r[[#INPUT]];{{[[:space:]]/*}} -; SM8X-NEXT: st.global.{{b|u}}32 [%rd[[#GOUT]]], %r[[#DEST]]; -; SM8X-NEXT: ret; +; SMXX-LABEL: .target sm_{{[0-9]+}} +; SMXX: .visible .entry [[FUNCTION:_.*add_relaxed_device_non_volatile.*]]( +; SMXX-DAG: ld.param.{{b|u}}64 %rd[[#ATOM:]], {{.*}}[[FUNCTION]]_param_0{{.*}} +; SMXX-DAG: ld.param.{{b|u}}64 %rd[[#RESULT:]], {{.*}}[[FUNCTION]]_param_1{{.*}} +; SMXX-DAG: ld.param.{{b|u}}32 %r[[#INPUT:]], {{.*}}[[FUNCTION]]_param_2{{.*}} +; SMXX-DAG: cvta.to.global.u64 %rd[[#GOUT:]], %rd[[#RESULT]]; +; SMXX-DAG: {{/*[[:space:]] *}}atom.add.relaxed.gpu.s32 %r[[#DEST:]],[%rd[[#ATOM]]],%r[[#INPUT]];{{[[:space:]]/*}} +; SMXX: st.global.{{b|u}}32 [%rd[[#GOUT]]], %r[[#DEST]]; +; SMXX-NEXT: ret; */ diff --git a/libcudacxx/test/atomic_codegen/atomic_cas_non_volatile.cu b/libcudacxx/test/atomic_codegen/atomic_cas_non_volatile.cu index cb753870d71a..dbed366ae495 100644 --- a/libcudacxx/test/atomic_codegen/atomic_cas_non_volatile.cu +++ b/libcudacxx/test/atomic_codegen/atomic_cas_non_volatile.cu @@ -9,15 +9,15 @@ __global__ void cas_device_relaxed_non_volatile(int* data, int* out, int n) // clang-format off /* -; SM8X-LABEL: .target sm_80 -; SM8X: .visible .entry [[FUNCTION:_.*cas_device_relaxed_non_volatile.*]]( -; SM8X-DAG: ld.param.{{b|u}}64 %rd[[#ATOM:]], {{.*}}[[FUNCTION]]_param_0{{.*}} -; SM8X-DAG: ld.param.{{b|u}}64 %rd[[#EXPECTED:]], {{.*}}[[FUNCTION]]_param_1{{.*}} -; SM8X-DAG: ld.param.{{b|u}}32 %r[[#INPUT:]], {{.*}}[[FUNCTION]]_param_2{{.*}} -; SM8X-DAG: cvta.to.global.u64 %rd[[#GOUT:]], %rd[[#EXPECTED]]; -; SM8X-DAG: ld.global.{{b|u}}32 %r[[#LOCALEXP:]], [%rd[[#INPUT]]]; -; SM8X-NEXT: {{/*[[:space:]] *}}atom.cas.relaxed.gpu.b32 %r[[#DEST:]],[%rd[[#ATOM]]],%r[[#LOCALEXP]],%r[[#INPUT]];{{[[:space:]]/*}} -; SM8X-NEXT: st.global.{{b|u}}32 [%rd[[#GOUT]]], %r[[#DEST]]; -; SM8X-NEXT: ret; +; SMXX-LABEL: .target sm_{{[0-9]+}} +; SMXX: .visible .entry [[FUNCTION:_.*cas_device_relaxed_non_volatile.*]]( +; SMXX-DAG: ld.param.{{b|u}}64 %rd[[#ATOM:]], {{.*}}[[FUNCTION]]_param_0{{.*}} +; SMXX-DAG: ld.param.{{b|u}}64 %rd[[#EXPECTED:]], {{.*}}[[FUNCTION]]_param_1{{.*}} +; SMXX-DAG: ld.param.{{b|u}}32 %r[[#INPUT:]], {{.*}}[[FUNCTION]]_param_2{{.*}} +; SMXX-DAG: cvta.to.global.u64 %rd[[#GOUT:]], %rd[[#EXPECTED]]; +; SMXX-DAG: ld.global.{{b|u}}32 %r[[#LOCALEXP:]], [%rd[[#INPUT]]]; +; SMXX-NEXT: {{/*[[:space:]] *}}atom.cas.relaxed.gpu.b32 %r[[#DEST:]],[%rd[[#ATOM]]],%r[[#LOCALEXP]],%r[[#INPUT]];{{[[:space:]]/*}} +; SMXX-NEXT: st.global.{{b|u}}32 [%rd[[#GOUT]]], %r[[#DEST]]; +; SMXX-NEXT: ret; */ diff --git a/libcudacxx/test/atomic_codegen/atomic_exch_non_volatile.cu b/libcudacxx/test/atomic_codegen/atomic_exch_non_volatile.cu index 83da211a9080..a3ea58e53644 100644 --- a/libcudacxx/test/atomic_codegen/atomic_exch_non_volatile.cu +++ b/libcudacxx/test/atomic_codegen/atomic_exch_non_volatile.cu @@ -8,14 +8,14 @@ __global__ void exch_device_relaxed_non_volatile(int* data, int* out, int n) /* -; SM8X-LABEL: .target sm_80 -; SM8X: .visible .entry [[FUNCTION:_.*exch_device_relaxed_non_volatile.*]]( -; SM8X-DAG: ld.param.{{b|u}}64 %rd[[#ATOM:]], {{.*}}[[FUNCTION]]_param_0{{.*}} -; SM8X-DAG: ld.param.{{b|u}}64 %rd[[#EXPECTED:]], {{.*}}[[FUNCTION]]_param_1{{.*}} -; SM8X-DAG: ld.param.{{b|u}}32 %r[[#INPUT:]], {{.*}}[[FUNCTION]]_param_2{{.*}} -; SM8X-DAG: cvta.to.global.u64 %rd[[#GOUT:]], %rd[[#EXPECTED]]; -; SM8X-NEXT: {{/*[[:space:]] *}}atom.exch.relaxed.gpu.b32 %r[[#DEST:]],[%rd[[#ATOM]]],%r[[#INPUT]];{{[[:space:]]/*}} -; SM8X-NEXT: st.global.{{b|u}}32 [%rd[[#GOUT]]], %r[[#DEST]]; -; SM8X-NEXT: ret; +; SMXX-LABEL: .target sm_{{[0-9]+}} +; SMXX: .visible .entry [[FUNCTION:_.*exch_device_relaxed_non_volatile.*]]( +; SMXX-DAG: ld.param.{{b|u}}64 %rd[[#ATOM:]], {{.*}}[[FUNCTION]]_param_0{{.*}} +; SMXX-DAG: ld.param.{{b|u}}64 %rd[[#EXPECTED:]], {{.*}}[[FUNCTION]]_param_1{{.*}} +; SMXX-DAG: ld.param.{{b|u}}32 %r[[#INPUT:]], {{.*}}[[FUNCTION]]_param_2{{.*}} +; SMXX-DAG: cvta.to.global.u64 %rd[[#GOUT:]], %rd[[#EXPECTED]]; +; SMXX-DAG: {{/*[[:space:]] *}}atom.exch.relaxed.gpu.b32 %r[[#DEST:]],[%rd[[#ATOM]]],%r[[#INPUT]];{{[[:space:]]/*}} +; SMXX: st.global.{{b|u}}32 [%rd[[#GOUT]]], %r[[#DEST]]; +; SMXX-NEXT: ret; */ diff --git a/libcudacxx/test/atomic_codegen/atomic_load_non_volatile.cu b/libcudacxx/test/atomic_codegen/atomic_load_non_volatile.cu index fce3bb4d3abe..64365b1fede5 100644 --- a/libcudacxx/test/atomic_codegen/atomic_load_non_volatile.cu +++ b/libcudacxx/test/atomic_codegen/atomic_load_non_volatile.cu @@ -8,13 +8,13 @@ __global__ void load_relaxed_device_non_volatile(int* data, int* out) /* -; SM8X-LABEL: .target sm_80 -; SM8X: .visible .entry [[FUNCTION:_.*load_relaxed_device_non_volatile.*]]( -; SM8X-DAG: ld.param.{{b|u}}64 %rd[[#ATOM:]], {{.*}}[[FUNCTION]]_param_0{{.*}} -; SM8X-DAG: ld.param.{{b|u}}64 %rd[[#EXPECTED:]], {{.*}}[[FUNCTION]]_param_1{{.*}} -; SM8X-DAG: cvta.to.global.u64 %rd[[#GOUT:]], %rd[[#EXPECTED]]; -; SM8X-NEXT: {{/*[[:space:]] *}}ld.relaxed.gpu.b32 %r[[#DEST:]],[%rd[[#ATOM]]];{{[[:space:]]/*}} -; SM8X-NEXT: st.global.{{b|u}}32 [%rd[[#GOUT]]], %r[[#DEST]]; -; SM8X-NEXT: ret; +; SMXX-LABEL: .target sm_{{[0-9]+}} +; SMXX: .visible .entry [[FUNCTION:_.*load_relaxed_device_non_volatile.*]]( +; SMXX-DAG: ld.param.{{b|u}}64 %rd[[#ATOM:]], {{.*}}[[FUNCTION]]_param_0{{.*}} +; SMXX-DAG: ld.param.{{b|u}}64 %rd[[#EXPECTED:]], {{.*}}[[FUNCTION]]_param_1{{.*}} +; SMXX-DAG: cvta.to.global.u64 %rd[[#GOUT:]], %rd[[#EXPECTED]]; +; SMXX-DAG: {{/*[[:space:]] *}}ld.relaxed.gpu.b32 %r[[#DEST:]],[%rd[[#ATOM]]];{{[[:space:]]/*}} +; SMXX: st.global.{{b|u}}32 [%rd[[#GOUT]]], %r[[#DEST]]; +; SMXX-NEXT: ret; */ diff --git a/libcudacxx/test/atomic_codegen/atomic_store_non_volatile.cu b/libcudacxx/test/atomic_codegen/atomic_store_non_volatile.cu index cd77894668e3..55b8a8252cd9 100644 --- a/libcudacxx/test/atomic_codegen/atomic_store_non_volatile.cu +++ b/libcudacxx/test/atomic_codegen/atomic_store_non_volatile.cu @@ -8,11 +8,11 @@ __global__ void store_relaxed_device_non_volatile(int* data, int in) /* -; SM8X-LABEL: .target sm_80 -; SM8X: .visible .entry [[FUNCTION:_.*store_relaxed_device_non_volatile.*]]( -; SM8X-DAG: ld.param.{{b|u}}64 %rd[[#ATOM:]], {{.*}}[[FUNCTION]]_param_0{{.*}} -; SM8X-DAG: ld.param.{{b|u}}32 %r[[#INPUT:]], {{.*}}[[FUNCTION]]_param_1{{.*}} -; SM8X-NEXT: {{/*[[:space:]] *}}st.relaxed.gpu.b32 [%rd[[#ATOM]]],%r[[#INPUT]];{{[[:space:]]/*}} -; SM8X-NEXT: ret; +; SMXX-LABEL: .target sm_{{[0-9]+}} +; SMXX: .visible .entry [[FUNCTION:_.*store_relaxed_device_non_volatile.*]]( +; SMXX-DAG: ld.param.{{b|u}}64 %rd[[#ATOM:]], {{.*}}[[FUNCTION]]_param_0{{.*}} +; SMXX-DAG: ld.param.{{b|u}}32 %r[[#INPUT:]], {{.*}}[[FUNCTION]]_param_1{{.*}} +; SMXX-NEXT: {{/*[[:space:]] *}}st.relaxed.gpu.b32 [%rd[[#ATOM]]],%r[[#INPUT]];{{[[:space:]]/*}} +; SMXX-NEXT: ret; */ diff --git a/libcudacxx/test/atomic_codegen/atomic_sub_non_volatile.cu b/libcudacxx/test/atomic_codegen/atomic_sub_non_volatile.cu index 3b23ac14c9d5..ecba2f196e86 100644 --- a/libcudacxx/test/atomic_codegen/atomic_sub_non_volatile.cu +++ b/libcudacxx/test/atomic_codegen/atomic_sub_non_volatile.cu @@ -8,15 +8,15 @@ __global__ void sub_relaxed_device_non_volatile(int* data, int* out, int n) /* -; SM8X-LABEL: .target sm_80 -; SM8X: .visible .entry [[FUNCTION:_.*sub_relaxed_device_non_volatile.*]]( -; SM8X-DAG: ld.param.{{b|u}}64 %rd[[#ATOM:]], {{.*}}[[FUNCTION]]_param_0{{.*}} -; SM8X-DAG: ld.param.{{b|u}}64 %rd[[#RESULT:]], {{.*}}[[FUNCTION]]_param_1{{.*}} -; SM8X-DAG: ld.param.{{b|u}}32 %r[[#INPUT:]], {{.*}}[[FUNCTION]]_param_2{{.*}} -; SM8X-DAG: cvta.to.global.u64 %rd[[#GOUT:]], %rd[[#RESULT]]; -; SM8X-NEXT: neg.s32 %r[[#NEG:]], %r[[#INPUT]]; -; SM8X-NEXT: {{/*[[:space:]] *}}atom.add.relaxed.gpu.s32 %r[[#DEST:]],[%rd[[#ATOM]]],%r[[#NEG]];{{[[:space:]]/*}} -; SM8X-NEXT: st.global.{{b|u}}32 [%rd[[#GOUT]]], %r[[#DEST]]; -; SM8X-NEXT: ret; +; SMXX-LABEL: .target sm_{{[0-9]+}} +; SMXX: .visible .entry [[FUNCTION:_.*sub_relaxed_device_non_volatile.*]]( +; SMXX-DAG: ld.param.{{b|u}}64 %rd[[#ATOM:]], {{.*}}[[FUNCTION]]_param_0{{.*}} +; SMXX-DAG: ld.param.{{b|u}}64 %rd[[#RESULT:]], {{.*}}[[FUNCTION]]_param_1{{.*}} +; SMXX-DAG: ld.param.{{b|u}}32 %r[[#INPUT:]], {{.*}}[[FUNCTION]]_param_2{{.*}} +; SMXX-DAG: cvta.to.global.u64 %rd[[#GOUT:]], %rd[[#RESULT]]; +; SMXX-DAG: neg.s32 %r[[#NEG:]], %r[[#INPUT]]; +; SMXX-DAG: {{/*[[:space:]] *}}atom.add.relaxed.gpu.s32 %r[[#DEST:]],[%rd[[#ATOM]]],%r[[#NEG]];{{[[:space:]]/*}} +; SMXX: st.global.{{b|u}}32 [%rd[[#GOUT]]], %r[[#DEST]]; +; SMXX-NEXT: ret; */ diff --git a/libcudacxx/test/atomic_codegen/sass/CMakeLists.txt b/libcudacxx/test/atomic_codegen/sass/CMakeLists.txt new file mode 100644 index 000000000000..79a6c2ff6f61 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/CMakeLists.txt @@ -0,0 +1,87 @@ +##===----------------------------------------------------------------------===## +## +## Part of libcu++ in the CUDA C++ Core Libraries, +## under the Apache License v2.0 with LLVM Exceptions. +## See https://llvm.org/LICENSE.txt for license information. +## SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +## SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +## +##===----------------------------------------------------------------------===## + +add_custom_target(libcudacxx.test.atomics.sass) + +if (NOT CMAKE_CUDA_COMPILER_ID STREQUAL NVIDIA) + message("-- Skipping atomic SASS codegen tests without NVCC") + return() +endif() + +libcudacxx_codegen_get_cuda_architectures(atomic_codegen_sass_cuda_archs) + +if (CMAKE_CUDA_COMPILER_VERSION VERSION_LESS 12.1) + set(atomic_codegen_sass_cuda_version_prefix CUDA12-0) +else() + set(atomic_codegen_sass_cuda_version_prefix CUDA12-1-PLUS) +endif() + +set(libcudacxx_atomic_codegen_tests) +if (NOT "NVHPC" STREQUAL "${CMAKE_CXX_COMPILER_ID}") + file(GLOB libcudacxx_atomic_codegen_tests CONFIGURE_DEPENDS "*.cu") +endif() + +if (NOT libcudacxx_atomic_codegen_tests) + return() +endif() + +libcudacxx_codegen_check_tools(libcudacxx_codegen_tests_enabled) +if (NOT libcudacxx_codegen_tests_enabled) + return() +endif() + +include(CheckSourceCompiles) +include(CMakePushCheckState) +cmake_push_check_state(RESET) +list(APPEND CMAKE_REQUIRED_INCLUDES "${libcudacxx_SOURCE_DIR}/include") +check_source_compiles( + CUDA + [[ + #include + + #if !_CCCL_HAS_INT128() + # error "128-bit integers are not supported" + #endif + + int main() { return 0; } + ]] + libcudacxx_atomic_codegen_has_int128 +) +cmake_pop_check_state() + +if (NOT libcudacxx_atomic_codegen_has_int128) + message( + STATUS + "Skipping 128-bit atomic SASS codegen tests: 128-bit integers are not supported" + ) + list(FILTER libcudacxx_atomic_codegen_tests EXCLUDE REGEX "_types_128_") +else() + # Native 128-bit atomic_ref operations require the PTX ISA 8.4 path available + # with CUDA 12.4 and newer. + if (CMAKE_CUDA_COMPILER_VERSION VERSION_LESS 12.4) + list( + FILTER libcudacxx_atomic_codegen_tests + EXCLUDE + REGEX "_types_128_atomic_ref[.]cu$" + ) + endif() +endif() + +# Restrict FileCheck to the wrapper so instructions in outlined helpers cannot +# mask an inlining regression. +libcudacxx_codegen_add_sass_tests( + AGGREGATE_TARGET libcudacxx.test.atomics.sass + TARGET_PREFIX atomic_codegen + ARCHITECTURES ${atomic_codegen_sass_cuda_archs} + DUMP_FUNCTIONS atomic_codegen_test + CHECK_PREFIXES ${atomic_codegen_sass_cuda_version_prefix} + TESTS ${libcudacxx_atomic_codegen_tests} + COMPILE_DEFINITIONS _CCCL_ATOMIC_UNSAFE_AUTOMATIC_STORAGE=1 +) diff --git a/libcudacxx/test/atomic_codegen/sass/atomic_codegen_helpers.h b/libcudacxx/test/atomic_codegen/sass/atomic_codegen_helpers.h new file mode 100644 index 000000000000..8508b4044c5f --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/atomic_codegen_helpers.h @@ -0,0 +1,61 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX_TEST_ATOMIC_CODEGEN_SASS_ATOMIC_CODEGEN_HELPERS_H +#define _LIBCUDACXX_TEST_ATOMIC_CODEGEN_SASS_ATOMIC_CODEGEN_HELPERS_H + +#include + +struct __half; +struct __nv_bfloat16; + +using f16 = __half; +using bf16 = __nv_bfloat16; + +inline constexpr auto tsb = cuda::thread_scope_block; +inline constexpr auto tsd = cuda::thread_scope_device; +inline constexpr auto tss = cuda::thread_scope_system; + +inline constexpr auto mor = cuda::std::memory_order_relaxed; +inline constexpr auto moa = cuda::std::memory_order_acquire; +inline constexpr auto more = cuda::std::memory_order_release; +inline constexpr auto moar = cuda::std::memory_order_acq_rel; +inline constexpr auto mosc = cuda::std::memory_order_seq_cst; + +#if _CCCL_HAS_INT128() +using i128 = __int128_t; +using u128 = __uint128_t; +#endif // _CCCL_HAS_INT128() + +template +using ca = cuda::atomic; + +template +using car = cuda::atomic_ref; + +template +using csa = cuda::std::atomic; + +template +using csar = cuda::std::atomic_ref; + +template +using vca = volatile cuda::atomic; + +template +using vcar = cuda::atomic_ref; + +template +using vcsa = volatile cuda::std::atomic; + +template +using vcsar = cuda::std::atomic_ref; + +#endif // _LIBCUDACXX_TEST_ATOMIC_CODEGEN_SASS_ATOMIC_CODEGEN_HELPERS_H diff --git a/libcudacxx/test/atomic_codegen/sass/compare_exchange_apis.cu b/libcudacxx/test/atomic_codegen/sass/compare_exchange_apis.cu new file mode 100644 index 000000000000..372cf739cf1c --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/compare_exchange_apis.cu @@ -0,0 +1,46 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% TEMPLATE,SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE api cab=ca,tsb,CTA,block:cad=ca,tsd,GPU,non_block:cas=ca,tss,SYS,non_block:carb=car,tsb,CTA,block:card=car,tsd,GPU,non_block:cars=car,tss,SYS,non_block:csa=csa,tss,SYS,non_block:csar=csar,tss,SYS,non_block +// %PARAM% CAS cas compare_exchange_weak:compare_exchange_strong +// %PARAM% SUCCESS_ORDER,FAILURE_ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order rr=mor,mor,,non_seq_cst,no_acquire,no_membar:ar=moa,mor,,non_seq_cst,acquire,no_membar:aa=moa,moa,,non_seq_cst,acquire,no_membar:er=more,mor,ALL,non_seq_cst,no_acquire,membar:br=moar,mor,ALL,non_seq_cst,acquire,membar:ba=moar,moa,ALL,non_seq_cst,acquire,membar:sr=mosc,mor,SC,seq_cst,acquire,membar:sa=mosc,moa,SC,seq_cst,acquire,membar:ss=mosc,mosc,SC,seq_cst,acquire,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ bool atomic_codegen_test(TEMPLATE& atom, int32_t& expected, int32_t desired) +{ + return atom.CAS(expected, desired, SUCCESS_ORDER, FAILURE_ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; BLOCK: {{.*}}ATOM.E.CAS.STRONG.{{CTA|SM}}{{.*}} +; NON_BLOCK: {{.*}}ATOM.E.CAS.STRONG.[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/compare_exchange_single_order.cu b/libcudacxx/test/atomic_codegen/sass/compare_exchange_single_order.cu new file mode 100644 index 000000000000..0275522456b2 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/compare_exchange_single_order.cu @@ -0,0 +1,48 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SINGLE_ORDER,OVERLOAD_KIND,FILECHECK_PREFIX_SCOPE overload single=1,single,non_block:pair=0,pair,non_block +// %PARAM% CAS cas compare_exchange_weak:compare_exchange_strong +// %PARAM% SUCCESS_ORDER,FAILURE_ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order relaxed=mor,mor,,non_seq_cst,no_acquire,no_membar:acquire=moa,moa,,non_seq_cst,acquire,no_membar:release=more,mor,ALL,non_seq_cst,no_acquire,membar:acq_rel=moar,moa,ALL,non_seq_cst,acquire,membar:seq_cst=mosc,mosc,SC,seq_cst,acquire,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ bool +atomic_codegen_test(cuda::atomic_ref& atom, int32_t& expected, int32_t desired) +{ +#if SINGLE_ORDER + return atom.CAS(expected, desired, SUCCESS_ORDER); +#else // ^^^ SINGLE_ORDER ^^^ / vvv !SINGLE_ORDER vvv + return atom.CAS(expected, desired, SUCCESS_ORDER, FAILURE_ORDER); +#endif // !SINGLE_ORDER +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].GPU{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX: {{.*}}ATOM.E.CAS.STRONG.GPU{{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/compare_exchange_types.cu b/libcudacxx/test/atomic_codegen/sass/compare_exchange_types.cu new file mode 100644 index 000000000000..c31e60f1e67f --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/compare_exchange_types.cu @@ -0,0 +1,47 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% TEMPLATE,SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE api cab=ca,tsb,CTA,block:card=car,tsd,GPU,non_block:csa=csa,tss,SYS,non_block +// %PARAM% TYPE,SASS_SIZE type i32=int32_t,:u32=uint32_t,:i64=int64_t,.64:u64=uint64_t,.64:f32=float,:f64=double,.64:ptr1=char*,.64:ptr4=int32_t*,.64 +// %PARAM% CAS cas compare_exchange_weak:compare_exchange_strong +// %PARAM% SUCCESS_ORDER,FAILURE_ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order rr=mor,mor,,non_seq_cst,no_acquire,no_membar:ar=moa,mor,,non_seq_cst,acquire,no_membar:aa=moa,moa,,non_seq_cst,acquire,no_membar:er=more,mor,ALL,non_seq_cst,no_acquire,membar:br=moar,mor,ALL,non_seq_cst,acquire,membar:ba=moar,moa,ALL,non_seq_cst,acquire,membar:sr=mosc,mor,SC,seq_cst,acquire,membar:sa=mosc,moa,SC,seq_cst,acquire,membar:ss=mosc,mosc,SC,seq_cst,acquire,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ bool atomic_codegen_test(TEMPLATE& atom, TYPE& expected, TYPE desired) +{ + return atom.CAS(expected, desired, SUCCESS_ORDER, FAILURE_ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; BLOCK: {{.*}}ATOM.E.CAS[[SASS_SIZE]].STRONG.{{CTA|SM}}{{.*}} +; NON_BLOCK: {{.*}}ATOM.E.CAS[[SASS_SIZE]].STRONG.[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/compare_exchange_types_128_atomic.cu b/libcudacxx/test/atomic_codegen/sass/compare_exchange_types_128_atomic.cu new file mode 100644 index 000000000000..698bd25f2645 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/compare_exchange_types_128_atomic.cu @@ -0,0 +1,52 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope block=tsb,CTA,block:device=tsd,GPU,non_block:system=tss,SYS,non_block +// %PARAM% TYPE type i128:u128 +// %PARAM% CAS cas compare_exchange_weak:compare_exchange_strong +// %PARAM% SUCCESS_ORDER,FAILURE_ORDER order rr=mor,mor:ar=moa,mor:aa=moa,moa:er=more,mor:br=moar,mor:ba=moar,moa:sr=mosc,mor:sa=mosc,moa:ss=mosc,mosc +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ bool atomic_codegen_test(cuda::atomic& atom, TYPE& expected, TYPE desired) +{ + return atom.CAS(expected, desired, SUCCESS_ORDER, FAILURE_ORDER); +} + +// clang-format off +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}MEMBAR.SC.{{.*}} +; SMXX: {{.*}}BSSY [[SYNC:B[0-9]+]], {{.*}} +; BLOCK: {{.*}}ATOM.E.EXCH.STRONG.{{CTA|SM}} PT, [[LOCK_STATE:R[0-9]+]], {{.*\[}}[[BASE_ADDR:R[0-9]+]]{{(\.64)?\+0x10\].*}}, {{R[0-9]+}}{{.*}} +; NON_BLOCK: {{.*}}ATOM.E.EXCH.STRONG.[[SASS_SCOPE]] PT, [[LOCK_STATE:R[0-9]+]], {{.*\[}}[[BASE_ADDR:R[0-9]+]]{{(\.64)?\+0x10\].*}}, {{R[0-9]+}}{{.*}} +; SMXX-DAG: {{.*}}ISETP.NE{{(\.U32)?}}.AND [[LOCK_ACQUIRED:P[0-9]+]], PT, [[LOCK_STATE]], 0x1, PT{{.*}} +; NON_BLOCK-DAG: {{.*}}CCTL.IVALL{{.*}} +; SMXX: {{.*}}{{@!?}}[[LOCK_ACQUIRED]] BRA{{(\.U)?}} {{.*}} +; SMXX: {{.*}}BSYNC [[SYNC]]{{.*}} +; SMXX: {{.*}}LD.E.128{{(\.SYS)?}} {{R[0-9]+}}, {{.*\[}}[[BASE_ADDR]]{{(\.64)?\].*}} +; SMXX: {{.*}}LD.E.128{{(\.SYS)?}} {{R[0-9]+}}, {{.*}} +; SMXX: {{.*}}{{LOP3\.LUT|ISETP\.NE[^ ]*}} [[CAS_RESULT_PRED:P[0-9]+]], {{.*}} +; CUDA12-0-DAG: {{.*}}SEL [[STORE_DATA:R[0-9]+]], {{R[0-9]+}}, {{R[0-9]+}}, ![[CAS_RESULT_PRED]]{{.*}} +; CUDA12-0-DAG: {{.*}}SEL [[STORE_ADDR:R[0-9]+]], [[BASE_ADDR]], {{R[0-9]+}}, ![[CAS_RESULT_PRED]]{{.*}} +; CUDA12-0-DAG: {{.*}}SEL {{R[0-9]+}}, RZ, 0x1, [[CAS_RESULT_PRED]]{{.*}} +; CUDA12-0-DAG: {{.*}}ST.E.128{{(\.SYS)?}} {{.*\[}}[[STORE_ADDR]]{{(\.64)?\].*}}, [[STORE_DATA]]{{.*}} +; CUDA12-1-PLUS-DAG: {{.*}}@[[CAS_RESULT_PRED]] ST.E.128{{(\.SYS)?}} {{.*}} +; CUDA12-1-PLUS-DAG: {{.*}}@![[CAS_RESULT_PRED]] ST.E.128{{(\.SYS)?}} {{.*\[}}[[BASE_ADDR]]{{(\.64)?\].*}} +; SMXX-DAG: {{.*}}MEMBAR.ALL.[[SASS_SCOPE]]{{.*}} +; BLOCK-DAG: {{.*}}ST.E.STRONG.{{CTA|SM}} {{.*\[}}[[BASE_ADDR]]{{(\.64)?\+0x10\].*}}, RZ{{.*}} +; NON_BLOCK-DAG: {{.*}}ST.E.STRONG.[[SASS_SCOPE]] {{.*\[}}[[BASE_ADDR]]{{(\.64)?\+0x10\].*}}, RZ{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ +// clang-format on diff --git a/libcudacxx/test/atomic_codegen/sass/compare_exchange_types_128_atomic_ref.cu b/libcudacxx/test/atomic_codegen/sass/compare_exchange_types_128_atomic_ref.cu new file mode 100644 index 000000000000..f86e08cde999 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/compare_exchange_types_128_atomic_ref.cu @@ -0,0 +1,53 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope block=tsb,CTA,block:device=tsd,GPU,non_block:system=tss,SYS,non_block +// %PARAM% TYPE type i128:u128 +// %PARAM% CAS cas compare_exchange_weak:compare_exchange_strong +// %PARAM% SUCCESS_ORDER,FAILURE_ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order rr=mor,mor,,non_seq_cst,no_acquire,no_membar:ar=moa,mor,,non_seq_cst,acquire,no_membar:aa=moa,moa,,non_seq_cst,acquire,no_membar:er=more,mor,ALL,non_seq_cst,no_acquire,membar:br=moar,mor,ALL,non_seq_cst,acquire,membar:ba=moar,moa,ALL,non_seq_cst,acquire,membar:sr=mosc,mor,SC,seq_cst,acquire,membar:sa=mosc,moa,SC,seq_cst,acquire,membar:ss=mosc,mosc,SC,seq_cst,acquire,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ bool atomic_codegen_test(cuda::atomic_ref& atom, TYPE& expected, TYPE desired) +{ + return atom.CAS(expected, desired, SUCCESS_ORDER, FAILURE_ORDER); +} + +// clang-format off +/* + +; SM90-PLUS-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX: {{.*}}LD.E.64{{(\.SYS)?}} [[ATOM_ADDR:R[0-9]+]], {{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR-DAG: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST-DAG: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-DAG: {{.*}}LD.E.128{{(\.SYS)?}} [[EXPECTED:R[0-9]+]], {{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; BLOCK: {{.*}}ATOM.E.CAS.128.STRONG.{{CTA|SM}} {{P(T|[0-9]+)}}, [[OLD:R[0-9]+]], {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}}, [[EXPECTED]], {{R[0-9]+}}{{.*}} +; NON_BLOCK: {{.*}}ATOM.E.CAS.128.STRONG.[[SASS_SCOPE]] {{P(T|[0-9]+)}}, [[OLD:R[0-9]+]], {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}}, [[EXPECTED]], {{R[0-9]+}}{{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX: {{.*}}ST.E.128{{(\.SYS)?}} {{.*}}, [[OLD]]{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ +// clang-format on diff --git a/libcudacxx/test/atomic_codegen/sass/compare_exchange_types_8_16_atomic.cu b/libcudacxx/test/atomic_codegen/sass/compare_exchange_types_8_16_atomic.cu new file mode 100644 index 000000000000..38be8de337c0 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/compare_exchange_types_8_16_atomic.cu @@ -0,0 +1,50 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope block=tsb,CTA,block:device=tsd,GPU,non_block:system=tss,SYS,non_block +// %PARAM% TYPE type int8_t:uint8_t:int16_t:uint16_t:f16:bf16 +// %PARAM% CAS cas compare_exchange_weak:compare_exchange_strong +// %PARAM% SUCCESS_ORDER,FAILURE_ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order rr=mor,mor,,non_seq_cst,no_acquire,no_membar:ar=moa,mor,,non_seq_cst,acquire,no_membar:aa=moa,moa,,non_seq_cst,acquire,no_membar:er=more,mor,ALL,non_seq_cst,no_acquire,membar:br=moar,mor,ALL,non_seq_cst,acquire,membar:ba=moar,moa,ALL,non_seq_cst,acquire,membar:sr=mosc,mor,SC,seq_cst,acquire,membar:sa=mosc,moa,SC,seq_cst,acquire,membar:ss=mosc,mosc,SC,seq_cst,acquire,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include +#include + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ bool atomic_codegen_test(cuda::atomic& atom, TYPE& expected, TYPE desired) +{ + return atom.CAS(expected, desired, SUCCESS_ORDER, FAILURE_ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; BLOCK: {{.*}}ATOM.E.CAS.STRONG.{{CTA|SM}}{{.*}} +; NON_BLOCK: {{.*}}ATOM.E.CAS.STRONG.[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/compare_exchange_types_8_16_atomic_ref.cu b/libcudacxx/test/atomic_codegen/sass/compare_exchange_types_8_16_atomic_ref.cu new file mode 100644 index 000000000000..ebabe772d71d --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/compare_exchange_types_8_16_atomic_ref.cu @@ -0,0 +1,55 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope block=tsb,CTA,block:device=tsd,GPU,non_block:system=tss,SYS,non_block +// %PARAM% TYPE type int8_t:uint8_t:int16_t:uint16_t:f16:bf16 +// %PARAM% CAS cas compare_exchange_weak:compare_exchange_strong +// %PARAM% SUCCESS_ORDER,FAILURE_ORDER,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order rr=mor,mor,non_seq_cst,no_acquire,no_membar:ar=moa,mor,non_seq_cst,acquire,no_membar:aa=moa,moa,non_seq_cst,acquire,no_membar:er=more,mor,non_seq_cst,no_acquire,release:br=moar,mor,non_seq_cst,acquire,release:ba=moar,moa,non_seq_cst,acquire,release:sr=mosc,mor,seq_cst,acquire,seq_cst:sa=mosc,moa,seq_cst,acquire,seq_cst:ss=mosc,mosc,seq_cst,acquire,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include +#include + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ bool atomic_codegen_test(cuda::atomic_ref& atom, TYPE& expected, TYPE desired) +{ + return atom.CAS(expected, desired, SUCCESS_ORDER, FAILURE_ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX: {{.*}}LD.E.64{{(\.SYS)?}} [[ATOM_ADDR:R[0-9]+]], {{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-DAG: {{.*}}LOP3.LUT [[ALIGNED_ADDR:R[0-9]+]], [[ATOM_ADDR]]{{(\.reuse)?}}, 0xfffffffc, {{.*}} +; SEQ_CST-DAG: {{.*}}MEMBAR.SC.[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST-DAG: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK: {{.*}}LD.E.STRONG.{{CTA|SM}} {{R[0-9]+}}, {{.*\[}}[[ALIGNED_ADDR]]{{(\.64)?\].*}} +; NON_BLOCK: {{.*}}LD.E.STRONG.[[SASS_SCOPE]] {{R[0-9]+}}, {{.*\[}}[[ALIGNED_ADDR]]{{(\.64)?\].*}} +; RELEASE: {{.*}}MEMBAR.ALL.[[SASS_SCOPE]]{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; BLOCK: {{.*}}ATOM.E.CAS.STRONG.{{CTA|SM}}{{.*\[}}[[ALIGNED_ADDR]]{{\].*}} +; NON_BLOCK: {{.*}}ATOM.E.CAS.STRONG.[[SASS_SCOPE]]{{.*\[}}[[ALIGNED_ADDR]]{{\].*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/compare_exchange_volatile_apis.cu b/libcudacxx/test/atomic_codegen/sass/compare_exchange_volatile_apis.cu new file mode 100644 index 000000000000..fe598f52da27 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/compare_exchange_volatile_apis.cu @@ -0,0 +1,49 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% TEMPLATE,SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE api vcab=vca,tsb,CTA,block:vcad=vca,tsd,GPU,non_block:vcas=vca,tss,SYS,non_block:vcarb=vcar,tsb,CTA,block:vcard=vcar,tsd,GPU,non_block:vcars=vcar,tss,SYS,non_block:vcsa=vcsa,tss,SYS,non_block:vcsar=vcsar,tss,SYS,non_block +// %PARAM% CAS cas compare_exchange_weak:compare_exchange_strong +// %PARAM% SUCCESS_ORDER,FAILURE_ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order rr=mor,mor,,non_seq_cst,no_acquire,no_membar:ar=moa,mor,,non_seq_cst,acquire,no_membar:aa=moa,moa,,non_seq_cst,acquire,no_membar:er=more,mor,ALL,non_seq_cst,no_acquire,membar:br=moar,mor,ALL,non_seq_cst,acquire,membar:ba=moar,moa,ALL,non_seq_cst,acquire,membar:sr=mosc,mor,SC,seq_cst,acquire,membar:sa=mosc,moa,SC,seq_cst,acquire,membar:ss=mosc,mosc,SC,seq_cst,acquire,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ bool atomic_codegen_test(TEMPLATE& atom, int32_t& expected, int32_t desired) +{ + return atom.CAS(expected, desired, SUCCESS_ORDER, FAILURE_ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.CAS{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.CAS{{.*}} +; BLOCK: {{.*}}ATOM.E.CAS.STRONG.{{CTA|SM}}{{.*}} +; NON_BLOCK: {{.*}}ATOM.E.CAS.STRONG.[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.CAS{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/compare_exchange_volatile_types.cu b/libcudacxx/test/atomic_codegen/sass/compare_exchange_volatile_types.cu new file mode 100644 index 000000000000..80f65613db74 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/compare_exchange_volatile_types.cu @@ -0,0 +1,50 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% TEMPLATE,SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE api vcad=vca,tsd,GPU,non_block:vcard=vcar,tsd,GPU,non_block +// %PARAM% TYPE,SASS_SIZE type i32=int32_t,:u32=uint32_t,:i64=int64_t,.64:u64=uint64_t,.64:f32=float,:f64=double,.64:ptr1=char*,.64:ptr4=int32_t*,.64 +// %PARAM% CAS cas compare_exchange_weak:compare_exchange_strong +// %PARAM% SUCCESS_ORDER,FAILURE_ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order rr=mor,mor,,non_seq_cst,no_acquire,no_membar:ar=moa,mor,,non_seq_cst,acquire,no_membar:aa=moa,moa,,non_seq_cst,acquire,no_membar:er=more,mor,ALL,non_seq_cst,no_acquire,membar:br=moar,mor,ALL,non_seq_cst,acquire,membar:ba=moar,moa,ALL,non_seq_cst,acquire,membar:sr=mosc,mor,SC,seq_cst,acquire,membar:sa=mosc,moa,SC,seq_cst,acquire,membar:ss=mosc,mosc,SC,seq_cst,acquire,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ bool atomic_codegen_test(TEMPLATE& atom, TYPE& expected, TYPE desired) +{ + return atom.CAS(expected, desired, SUCCESS_ORDER, FAILURE_ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.CAS{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.CAS{{.*}} +; BLOCK: {{.*}}ATOM.E.CAS[[SASS_SIZE]].STRONG.{{CTA|SM}}{{.*}} +; NON_BLOCK: {{.*}}ATOM.E.CAS[[SASS_SIZE]].STRONG.[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.CAS{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/compare_exchange_volatile_types_128_atomic_ref.cu b/libcudacxx/test/atomic_codegen/sass/compare_exchange_volatile_types_128_atomic_ref.cu new file mode 100644 index 000000000000..2a9bf67733fe --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/compare_exchange_volatile_types_128_atomic_ref.cu @@ -0,0 +1,57 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope block=tsb,CTA,block:device=tsd,GPU,non_block:system=tss,SYS,non_block +// %PARAM% TYPE type i128:u128 +// %PARAM% CAS cas compare_exchange_weak:compare_exchange_strong +// %PARAM% SUCCESS_ORDER,FAILURE_ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order rr=mor,mor,,non_seq_cst,no_acquire,no_membar:ar=moa,mor,,non_seq_cst,acquire,no_membar:aa=moa,moa,,non_seq_cst,acquire,no_membar:er=more,mor,ALL,non_seq_cst,no_acquire,membar:br=moar,mor,ALL,non_seq_cst,acquire,membar:ba=moar,moa,ALL,non_seq_cst,acquire,membar:sr=mosc,mor,SC,seq_cst,acquire,membar:sa=mosc,moa,SC,seq_cst,acquire,membar:ss=mosc,mosc,SC,seq_cst,acquire,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ bool +atomic_codegen_test(cuda::atomic_ref& atom, TYPE& expected, TYPE desired) +{ + return atom.CAS(expected, desired, SUCCESS_ORDER, FAILURE_ORDER); +} + +// clang-format off +/* + +; SM90-PLUS-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX: {{.*}}LD.E.64{{(\.SYS)?}} [[ATOM_ADDR:R[0-9]+]], {{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR-DAG: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST-DAG: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-DAG: {{.*}}LD.E.128{{(\.SYS)?}} [[EXPECTED:R[0-9]+]], {{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.E{{.*\[}}[[ATOM_ADDR]]{{(\.64)?(\+0x[0-9a-f]+)?\].*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; BLOCK: {{.*}}ATOM.E.CAS.128.STRONG.{{CTA|SM}} {{P(T|[0-9]+)}}, [[OLD:R[0-9]+]], {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}}, [[EXPECTED]], {{R[0-9]+}}{{.*}} +; NON_BLOCK: {{.*}}ATOM.E.CAS.128.STRONG.[[SASS_SCOPE]] {{P(T|[0-9]+)}}, [[OLD:R[0-9]+]], {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}}, [[EXPECTED]], {{R[0-9]+}}{{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.E{{.*\[}}[[ATOM_ADDR]]{{(\.64)?(\+0x[0-9a-f]+)?\].*}} +; SMXX: {{.*}}ST.E.128{{(\.SYS)?}} {{.*}}, [[OLD]]{{.*}} +; SMXX-NOT: {{.*}}ATOM.E{{.*\[}}[[ATOM_ADDR]]{{(\.64)?(\+0x[0-9a-f]+)?\].*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ +// clang-format on diff --git a/libcudacxx/test/atomic_codegen/sass/compare_exchange_volatile_types_8_16_atomic.cu b/libcudacxx/test/atomic_codegen/sass/compare_exchange_volatile_types_8_16_atomic.cu new file mode 100644 index 000000000000..cf439eb681e0 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/compare_exchange_volatile_types_8_16_atomic.cu @@ -0,0 +1,54 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope device=tsd,GPU,non_block +// %PARAM% TYPE type int8_t:uint8_t:int16_t:uint16_t:f16:bf16 +// Strong compare-exchange may retry internally using weak compare-exchange; only the weak overload must contain one CAS. +// %PARAM% CAS,FILECHECK_PREFIX_SINGLE_CAS cas compare_exchange_weak=compare_exchange_weak,single_cas:compare_exchange_strong=compare_exchange_strong,smxx +// %PARAM% SUCCESS_ORDER,FAILURE_ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order rr=mor,mor,,non_seq_cst,no_acquire,no_membar:ar=moa,mor,,non_seq_cst,acquire,no_membar:aa=moa,moa,,non_seq_cst,acquire,no_membar:er=more,mor,ALL,non_seq_cst,no_acquire,membar:br=moar,mor,ALL,non_seq_cst,acquire,membar:ba=moar,moa,ALL,non_seq_cst,acquire,membar:sr=mosc,mor,SC,seq_cst,acquire,membar:sa=mosc,moa,SC,seq_cst,acquire,membar:ss=mosc,mosc,SC,seq_cst,acquire,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include +#include + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ bool atomic_codegen_test(volatile cuda::atomic& atom, TYPE& expected, TYPE desired) +{ + return atom.CAS(expected, desired, SUCCESS_ORDER, FAILURE_ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.CAS{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.CAS{{.*}} +; BLOCK: {{.*}}ATOM.E.CAS.STRONG.{{CTA|SM}}{{.*}} +; NON_BLOCK: {{.*}}ATOM.E.CAS.STRONG.[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SINGLE_CAS-NOT: {{.*}}ATOM.E.CAS{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/compare_exchange_volatile_types_8_16_atomic_ref.cu b/libcudacxx/test/atomic_codegen/sass/compare_exchange_volatile_types_8_16_atomic_ref.cu new file mode 100644 index 000000000000..a26e5a983525 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/compare_exchange_volatile_types_8_16_atomic_ref.cu @@ -0,0 +1,63 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope device=tsd,GPU,non_block +// %PARAM% TYPE type int8_t:uint8_t:int16_t:uint16_t:f16:bf16 +// Strong compare-exchange may retry internally using weak compare-exchange; only the weak overload must contain one CAS. +// %PARAM% CAS,FILECHECK_PREFIX_SINGLE_CAS cas compare_exchange_weak=compare_exchange_weak,single_cas:compare_exchange_strong=compare_exchange_strong,smxx +// %PARAM% SUCCESS_ORDER,FAILURE_ORDER,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order rr=mor,mor,non_seq_cst,no_acquire,no_membar:ar=moa,mor,non_seq_cst,acquire,no_membar:aa=moa,moa,non_seq_cst,acquire,no_membar:er=more,mor,non_seq_cst,no_acquire,release:br=moar,mor,non_seq_cst,acquire,release:ba=moar,moa,non_seq_cst,acquire,release:sr=mosc,mor,seq_cst,acquire,seq_cst:sa=mosc,moa,seq_cst,acquire,seq_cst:ss=mosc,mosc,seq_cst,acquire,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include +#include + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ bool +atomic_codegen_test(cuda::atomic_ref& atom, TYPE& expected, TYPE desired) +{ + return atom.CAS(expected, desired, SUCCESS_ORDER, FAILURE_ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX: {{.*}}LD.E.64{{(\.SYS)?}} [[ATOM_ADDR:R[0-9]+]], {{.*}} +; SMXX-NOT: {{.*}}LD.E{{.*}}.STRONG{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-DAG: {{.*}}LOP3.LUT [[ALIGNED_ADDR:R[0-9]+]], [[ATOM_ADDR]]{{(\.reuse)?}}, 0xfffffffc, {{.*}} +; SEQ_CST-DAG: {{.*}}MEMBAR.SC.[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST-DAG: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}LD.E{{.*}}.STRONG{{.*}} +; BLOCK: {{.*}}LD.E.STRONG.{{CTA|SM}} {{R[0-9]+}}, {{.*\[}}[[ALIGNED_ADDR]]{{(\.64)?\].*}} +; NON_BLOCK: {{.*}}LD.E.STRONG.[[SASS_SCOPE]] {{R[0-9]+}}, {{.*\[}}[[ALIGNED_ADDR]]{{(\.64)?\].*}} +; SMXX-NOT: {{.*}}LD.E{{.*}}.STRONG{{.*}} +; RELEASE: {{.*}}MEMBAR.ALL.[[SASS_SCOPE]]{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}LD.E{{.*}}.STRONG{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.CAS{{.*}} +; BLOCK: {{.*}}ATOM.E.CAS.STRONG.{{CTA|SM}}{{.*\[}}[[ALIGNED_ADDR]]{{\].*}} +; NON_BLOCK: {{.*}}ATOM.E.CAS.STRONG.[[SASS_SCOPE]]{{.*\[}}[[ALIGNED_ADDR]]{{\].*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SINGLE_CAS-NOT: {{.*}}ATOM.E.CAS{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/exchange_apis.cu b/libcudacxx/test/atomic_codegen/sass/exchange_apis.cu new file mode 100644 index 000000000000..e5629f537509 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/exchange_apis.cu @@ -0,0 +1,45 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% TEMPLATE,SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE api cab=ca,tsb,CTA,block:cad=ca,tsd,GPU,non_block:cas=ca,tss,SYS,non_block:carb=car,tsb,CTA,block:card=car,tsd,GPU,non_block:cars=car,tss,SYS,non_block:csa=csa,tss,SYS,non_block:csar=csar,tss,SYS,non_block +// %PARAM% ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order relaxed=mor,,non_seq_cst,no_acquire,no_membar:acquire=moa,,non_seq_cst,acquire,no_membar:release=more,ALL,non_seq_cst,no_acquire,membar:acq_rel=moar,ALL,non_seq_cst,acquire,membar:seq_cst=mosc,SC,seq_cst,acquire,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ auto atomic_codegen_test(TEMPLATE& atom, int32_t value) +{ + return atom.exchange(value, ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}}CAS{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}}CAS{{.*}} +; BLOCK: {{.*}}ATOM.E.EXCH.STRONG.{{CTA|SM}} PT, R4, {{.*}}, R6{{.*}} +; NON_BLOCK: {{.*}}ATOM.E.EXCH.STRONG.[[SASS_SCOPE]] PT, R4, {{.*}}, R6{{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}}CAS{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/exchange_types.cu b/libcudacxx/test/atomic_codegen/sass/exchange_types.cu new file mode 100644 index 000000000000..5e48a3e91b7e --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/exchange_types.cu @@ -0,0 +1,46 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% TEMPLATE,SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE api cab=ca,tsb,CTA,block:card=car,tsd,GPU,non_block:csa=csa,tss,SYS,non_block +// %PARAM% TYPE,SASS_SIZE type i32=int32_t,:u32=uint32_t,:i64=int64_t,.64:u64=uint64_t,.64:f32=float,:f64=double,.64:ptr1=char*,.64:ptr4=int32_t*,.64 +// %PARAM% ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order relaxed=mor,,non_seq_cst,no_acquire,no_membar:acquire=moa,,non_seq_cst,acquire,no_membar:release=more,ALL,non_seq_cst,no_acquire,membar:acq_rel=moar,ALL,non_seq_cst,acquire,membar:seq_cst=mosc,SC,seq_cst,acquire,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ auto atomic_codegen_test(TEMPLATE& atom, TYPE value) +{ + return atom.exchange(value, ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}}CAS{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}}CAS{{.*}} +; BLOCK: {{.*}}ATOM.E.EXCH[[SASS_SIZE]].STRONG.{{CTA|SM}} {{P(T|[0-9]+)}}, R4, {{.*}}, R6{{.*}} +; NON_BLOCK: {{.*}}ATOM.E.EXCH[[SASS_SIZE]].STRONG.[[SASS_SCOPE]] {{P(T|[0-9]+)}}, R4, {{.*}}, R6{{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}}CAS{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/exchange_types_128_atomic.cu b/libcudacxx/test/atomic_codegen/sass/exchange_types_128_atomic.cu new file mode 100644 index 000000000000..4ea7a1a9824a --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/exchange_types_128_atomic.cu @@ -0,0 +1,44 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope block=tsb,CTA,block:device=tsd,GPU,non_block:system=tss,SYS,non_block +// %PARAM% TYPE type i128:u128 +// %PARAM% ORDER order mor:moa:more:moar:mosc +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ auto atomic_codegen_test(cuda::atomic& atom, TYPE value) +{ + return atom.exchange(value, ORDER); +} + +// clang-format off +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}MEMBAR.SC.{{.*}} +; SMXX: {{.*}}BSSY [[SYNC:B[0-9]+]], {{.*}} +; BLOCK: {{.*}}ATOM.E.EXCH.STRONG.{{CTA|SM}} PT, [[LOCK_STATE:R[0-9]+]], {{.*\[}}[[BASE_ADDR:R[0-9]+]]{{(\.64)?\+0x10\].*}}, {{R[0-9]+}}{{.*}} +; NON_BLOCK: {{.*}}ATOM.E.EXCH.STRONG.[[SASS_SCOPE]] PT, [[LOCK_STATE:R[0-9]+]], {{.*\[}}[[BASE_ADDR:R[0-9]+]]{{(\.64)?\+0x10\].*}}, {{R[0-9]+}}{{.*}} +; SMXX-DAG: {{.*}}ISETP.NE{{(\.U32)?}}.AND [[LOCK_FREE:P[0-9]+]], PT, [[LOCK_STATE]], 0x1, PT{{.*}} +; NON_BLOCK-DAG: {{.*}}CCTL.IVALL{{.*}} +; SMXX: {{.*}}{{@!?}}[[LOCK_FREE]] BRA{{(\.U)?}} {{.*}} +; SMXX: {{.*}}BSYNC [[SYNC]]{{.*}} +; SMXX: {{.*}}LD.E.128{{(\.SYS)?}} R4, {{.*\[}}[[BASE_ADDR]]{{(\.64)?\].*}} +; SMXX: {{.*}}ST.E.128{{(\.SYS)?}} {{.*\[}}[[BASE_ADDR]]{{(\.64)?\].*}}, R8{{.*}} +; SMXX: {{.*}}MEMBAR.ALL.[[SASS_SCOPE]]{{.*}} +; BLOCK: {{.*}}ST.E.STRONG.{{CTA|SM}} {{.*\[}}[[BASE_ADDR]]{{(\.64)?\+0x10\].*}}, RZ{{.*}} +; NON_BLOCK: {{.*}}ST.E.STRONG.[[SASS_SCOPE]] {{.*\[}}[[BASE_ADDR]]{{(\.64)?\+0x10\].*}}, RZ{{.*}} +; SMXX-NEXT: {{.*}}RET.ABS.NODEC{{.*}} + +*/ +// clang-format on diff --git a/libcudacxx/test/atomic_codegen/sass/exchange_types_128_atomic_ref.cu b/libcudacxx/test/atomic_codegen/sass/exchange_types_128_atomic_ref.cu new file mode 100644 index 000000000000..7aa53eae6396 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/exchange_types_128_atomic_ref.cu @@ -0,0 +1,50 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope block=tsb,CTA,block:device=tsd,GPU,non_block:system=tss,SYS,non_block +// %PARAM% TYPE type i128:u128 +// %PARAM% ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order relaxed=mor,,non_seq_cst,no_acquire,no_membar:acquire=moa,,non_seq_cst,acquire,no_membar:release=more,ALL,non_seq_cst,no_acquire,membar:acq_rel=moar,ALL,non_seq_cst,acquire,membar:seq_cst=mosc,SC,seq_cst,acquire,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ auto atomic_codegen_test(cuda::atomic_ref& atom, TYPE value) +{ + return atom.exchange(value, ORDER); +} + +/* + +; SM90-PLUS-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.E.EXCH.STRONG{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}}CAS{{.*}} +; SMXX: {{.*}}LD.E.64{{(\.SYS)?}} [[ATOM_ADDR:R[0-9]+]], {{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH.STRONG{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}}CAS{{.*}} +; BLOCK: {{.*}}ATOM.E.EXCH.128.STRONG.{{CTA|SM}} PT, R4, {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}}, R8{{.*}} +; NON_BLOCK: {{.*}}ATOM.E.EXCH.128.STRONG.[[SASS_SCOPE]] PT, R4, {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}}, R8{{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH.STRONG{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}}CAS{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/exchange_types_8_16_atomic.cu b/libcudacxx/test/atomic_codegen/sass/exchange_types_8_16_atomic.cu new file mode 100644 index 000000000000..6a1374ff2c15 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/exchange_types_8_16_atomic.cu @@ -0,0 +1,49 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope block=tsb,CTA,block:device=tsd,GPU,non_block:system=tss,SYS,non_block +// %PARAM% TYPE type int8_t:uint8_t:int16_t:uint16_t:f16:bf16 +// %PARAM% ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order relaxed=mor,,non_seq_cst,no_acquire,no_membar:acquire=moa,,non_seq_cst,acquire,no_membar:release=more,ALL,non_seq_cst,no_acquire,membar:acq_rel=moar,ALL,non_seq_cst,acquire,membar:seq_cst=mosc,SC,seq_cst,acquire,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include +#include + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ auto atomic_codegen_test(cuda::atomic& atom, TYPE value) +{ + return atom.exchange(value, ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}}CAS{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}}CAS{{.*}} +; BLOCK: {{.*}}ATOM.E.EXCH.STRONG.{{CTA|SM}}{{.*}} +; NON_BLOCK: {{.*}}ATOM.E.EXCH.STRONG.[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}}CAS{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/exchange_types_8_16_atomic_ref.cu b/libcudacxx/test/atomic_codegen/sass/exchange_types_8_16_atomic_ref.cu new file mode 100644 index 000000000000..a9899fba38ec --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/exchange_types_8_16_atomic_ref.cu @@ -0,0 +1,55 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope block=tsb,CTA,block:device=tsd,GPU,non_block:system=tss,SYS,non_block +// %PARAM% TYPE type int8_t:uint8_t:int16_t:uint16_t:f16:bf16 +// %PARAM% ORDER,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order relaxed=mor,non_seq_cst,no_acquire,no_membar:acquire=moa,non_seq_cst,acquire,no_membar:release=more,non_seq_cst,no_acquire,release:acq_rel=moar,non_seq_cst,acquire,release:seq_cst=mosc,seq_cst,acquire,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include +#include + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ auto atomic_codegen_test(cuda::atomic_ref& atom, TYPE value) +{ + return atom.exchange(value, ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX: {{.*}}LD.E.64{{(\.SYS)?}} [[ATOM_ADDR:R[0-9]+]], {{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-DAG: {{.*}}LOP3.LUT [[A:R[0-9]+]], [[ATOM_ADDR]], 0xfffffffc, {{.*}} +; SEQ_CST-DAG: {{.*}}MEMBAR.SC.[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST-DAG: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK: {{.*}}LD.E.STRONG.{{CTA|SM}} [[E:R[0-9]+]], {{.*\[}}[[A]]{{(\.64)?\].*}} +; NON_BLOCK: {{.*}}LD.E.STRONG.[[SASS_SCOPE]] [[E:R[0-9]+]], {{.*\[}}[[A]]{{(\.64)?\].*}} +; RELEASE: {{.*}}MEMBAR.ALL.[[SASS_SCOPE]]{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; BLOCK: {{.*}}ATOM.E.CAS.STRONG.{{CTA|SM}} PT, [[C:R[0-9]+]], {{\[}}[[A]]{{\]}}, [[E]], [[C]]{{.*}} +; NON_BLOCK: {{.*}}ATOM.E.CAS.STRONG.[[SASS_SCOPE]] PT, [[C:R[0-9]+]], {{\[}}[[A]]{{\]}}, [[E]], [[C]]{{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX: {{.*}}ISETP.NE{{.*}} [[C]], [[E]], {{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/exchange_volatile_apis.cu b/libcudacxx/test/atomic_codegen/sass/exchange_volatile_apis.cu new file mode 100644 index 000000000000..bea1451c1e77 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/exchange_volatile_apis.cu @@ -0,0 +1,48 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% TEMPLATE,SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE api vcab=vca,tsb,CTA,block:vcad=vca,tsd,GPU,non_block:vcas=vca,tss,SYS,non_block:vcarb=vcar,tsb,CTA,block:vcard=vcar,tsd,GPU,non_block:vcars=vcar,tss,SYS,non_block:vcsa=vcsa,tss,SYS,non_block:vcsar=vcsar,tss,SYS,non_block +// %PARAM% ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order relaxed=mor,,non_seq_cst,no_acquire,no_membar:acquire=moa,,non_seq_cst,acquire,no_membar:release=more,ALL,non_seq_cst,no_acquire,membar:acq_rel=moar,ALL,non_seq_cst,acquire,membar:seq_cst=mosc,SC,seq_cst,acquire,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ auto atomic_codegen_test(TEMPLATE& atom, int32_t value) +{ + return atom.exchange(value, ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}}CAS{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}}CAS{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; BLOCK: {{.*}}ATOM.E.EXCH.STRONG.{{CTA|SM}} PT, R4, {{.*}}, R6{{.*}} +; NON_BLOCK: {{.*}}ATOM.E.EXCH.STRONG.[[SASS_SCOPE]] PT, R4, {{.*}}, R6{{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}}CAS{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/exchange_volatile_types.cu b/libcudacxx/test/atomic_codegen/sass/exchange_volatile_types.cu new file mode 100644 index 000000000000..4cf8e8da07f4 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/exchange_volatile_types.cu @@ -0,0 +1,49 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% TEMPLATE,SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE api vcad=vca,tsd,GPU,non_block:vcard=vcar,tsd,GPU,non_block +// %PARAM% TYPE,SASS_SIZE type i32=int32_t,:u32=uint32_t,:i64=int64_t,.64:u64=uint64_t,.64:f32=float,:f64=double,.64:ptr1=char*,.64:ptr4=int32_t*,.64 +// %PARAM% ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order relaxed=mor,,non_seq_cst,no_acquire,no_membar:acquire=moa,,non_seq_cst,acquire,no_membar:release=more,ALL,non_seq_cst,no_acquire,membar:acq_rel=moar,ALL,non_seq_cst,acquire,membar:seq_cst=mosc,SC,seq_cst,acquire,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ auto atomic_codegen_test(TEMPLATE& atom, TYPE value) +{ + return atom.exchange(value, ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}}CAS{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}}CAS{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; BLOCK: {{.*}}ATOM.E.EXCH[[SASS_SIZE]].STRONG.{{CTA|SM}} {{P(T|[0-9]+)}}, R4, {{.*}}, R6{{.*}} +; NON_BLOCK: {{.*}}ATOM.E.EXCH[[SASS_SIZE]].STRONG.[[SASS_SCOPE]] {{P(T|[0-9]+)}}, R4, {{.*}}, R6{{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}}CAS{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/exchange_volatile_types_128_atomic_ref.cu b/libcudacxx/test/atomic_codegen/sass/exchange_volatile_types_128_atomic_ref.cu new file mode 100644 index 000000000000..c8e1c6c137a1 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/exchange_volatile_types_128_atomic_ref.cu @@ -0,0 +1,52 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope block=tsb,CTA,block:device=tsd,GPU,non_block:system=tss,SYS,non_block +// %PARAM% TYPE type i128:u128 +// %PARAM% ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order relaxed=mor,,non_seq_cst,no_acquire,no_membar:acquire=moa,,non_seq_cst,acquire,no_membar:release=more,ALL,non_seq_cst,no_acquire,membar:acq_rel=moar,ALL,non_seq_cst,acquire,membar:seq_cst=mosc,SC,seq_cst,acquire,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ auto atomic_codegen_test(cuda::atomic_ref& atom, TYPE value) +{ + return atom.exchange(value, ORDER); +} + +/* + +; SM90-PLUS-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.E.EXCH.STRONG{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}}CAS{{.*}} +; SMXX: {{.*}}LD.E.64{{(\.SYS)?}} [[ATOM_ADDR:R[0-9]+]], {{.*}} +; SMXX-NOT: {{.*}}ATOM.E{{.*\[}}[[ATOM_ADDR]]{{(\.64)?(\+0x[0-9a-f]+)?\].*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.E{{.*\[}}[[ATOM_ADDR]]{{(\.64)?(\+0x[0-9a-f]+)?\].*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH.STRONG{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}}CAS{{.*}} +; BLOCK: {{.*}}ATOM.E.EXCH.128.STRONG.{{CTA|SM}} PT, R4, {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}}, R8{{.*}} +; NON_BLOCK: {{.*}}ATOM.E.EXCH.128.STRONG.[[SASS_SCOPE]] PT, R4, {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}}, R8{{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.E{{.*\[}}[[ATOM_ADDR]]{{(\.64)?(\+0x[0-9a-f]+)?\].*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}}CAS{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/exchange_volatile_types_8_16_atomic.cu b/libcudacxx/test/atomic_codegen/sass/exchange_volatile_types_8_16_atomic.cu new file mode 100644 index 000000000000..2dd4e3e72401 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/exchange_volatile_types_8_16_atomic.cu @@ -0,0 +1,52 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope device=tsd,GPU,non_block +// %PARAM% TYPE type int8_t:uint8_t:int16_t:uint16_t:f16:bf16 +// %PARAM% ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order relaxed=mor,,non_seq_cst,no_acquire,no_membar:acquire=moa,,non_seq_cst,acquire,no_membar:release=more,ALL,non_seq_cst,no_acquire,membar:acq_rel=moar,ALL,non_seq_cst,acquire,membar:seq_cst=mosc,SC,seq_cst,acquire,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include +#include + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ auto atomic_codegen_test(volatile cuda::atomic& atom, TYPE value) +{ + return atom.exchange(value, ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}}CAS{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}}CAS{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; BLOCK: {{.*}}ATOM.E.EXCH.STRONG.{{CTA|SM}}{{.*}} +; NON_BLOCK: {{.*}}ATOM.E.EXCH.STRONG.[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}}CAS{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/exchange_volatile_types_8_16_atomic_ref.cu b/libcudacxx/test/atomic_codegen/sass/exchange_volatile_types_8_16_atomic_ref.cu new file mode 100644 index 000000000000..ecddd2ad3fc2 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/exchange_volatile_types_8_16_atomic_ref.cu @@ -0,0 +1,60 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope device=tsd,GPU,non_block +// %PARAM% TYPE type int8_t:uint8_t:int16_t:uint16_t:f16:bf16 +// %PARAM% ORDER,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order relaxed=mor,non_seq_cst,no_acquire,no_membar:acquire=moa,non_seq_cst,acquire,no_membar:release=more,non_seq_cst,no_acquire,release:acq_rel=moar,non_seq_cst,acquire,release:seq_cst=mosc,seq_cst,acquire,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include +#include + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ auto atomic_codegen_test(cuda::atomic_ref& atom, TYPE value) +{ + return atom.exchange(value, ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX: {{.*}}LD.E.64{{(\.SYS)?}} [[ATOM_ADDR:R[0-9]+]], {{.*}} +; SMXX-NOT: {{.*}}LD.E{{.*}}.STRONG{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-DAG: {{.*}}LOP3.LUT [[A:R[0-9]+]], [[ATOM_ADDR]], 0xfffffffc, {{.*}} +; SEQ_CST-DAG: {{.*}}MEMBAR.SC.[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST-DAG: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}LD.E{{.*}}.STRONG{{.*}} +; BLOCK: {{.*}}LD.E.STRONG.{{CTA|SM}} [[E:R[0-9]+]], {{.*\[}}[[A]]{{(\.64)?\].*}} +; NON_BLOCK: {{.*}}LD.E.STRONG.[[SASS_SCOPE]] [[E:R[0-9]+]], {{.*\[}}[[A]]{{(\.64)?\].*}} +; SMXX-NOT: {{.*}}LD.E{{.*}}.STRONG{{.*}} +; RELEASE: {{.*}}MEMBAR.ALL.[[SASS_SCOPE]]{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}LD.E{{.*}}.STRONG{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; BLOCK: {{.*}}ATOM.E.CAS.STRONG.{{CTA|SM}} PT, [[C:R[0-9]+]], {{\[}}[[A]]{{\]}}, [[E]], [[C]]{{.*}} +; NON_BLOCK: {{.*}}ATOM.E.CAS.STRONG.[[SASS_SCOPE]] PT, [[C:R[0-9]+]], {{\[}}[[A]]{{\]}}, [[E]], [[C]]{{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX: {{.*}}ISETP.NE{{.*}} [[C]], [[E]], {{.*}} +; SMXX-NOT: {{.*}}ATOM.E.EXCH{{.*}} +; SMXX-NOT: {{.*}}ATOM.E.CAS{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/load_apis.cu b/libcudacxx/test/atomic_codegen/sass/load_apis.cu new file mode 100644 index 000000000000..e3cd70859b8e --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/load_apis.cu @@ -0,0 +1,45 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% TEMPLATE,SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE api cab=ca,tsb,CTA,block:cad=ca,tsd,GPU,non_block:cas=ca,tss,SYS,non_block:carb=car,tsb,CTA,block:card=car,tsd,GPU,non_block:cars=car,tss,SYS,non_block:csa=csa,tss,SYS,non_block:csar=csar,tss,SYS,non_block +// %PARAM% ORDER,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order relaxed=mor,no_acquire,non_sc:acquire=moa,acquire,non_sc:seq_cst=mosc,acquire,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ auto atomic_codegen_test(TEMPLATE& atom) +{ + return atom.load(ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}CCTL.IVALL{{.*}} +; SEQ_CST: {{.*}}MEMBAR.SC.[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; BLOCK: {{.*}}LD.E.STRONG.{{CTA|SM}} R4, {{.*}} +; NON_BLOCK: {{.*}}LD.E.STRONG.[[SASS_SCOPE]] R4, {{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/load_types.cu b/libcudacxx/test/atomic_codegen/sass/load_types.cu new file mode 100644 index 000000000000..beccc2690265 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/load_types.cu @@ -0,0 +1,46 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% TEMPLATE,SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE api cab=ca,tsb,CTA,block:card=car,tsd,GPU,non_block:csa=csa,tss,SYS,non_block +// %PARAM% TYPE,SASS_SIZE type i32=int32_t,:u32=uint32_t,:i64=int64_t,.64:u64=uint64_t,.64:f32=float,:f64=double,.64:ptr1=char*,.64:ptr4=int32_t*,.64 +// %PARAM% ORDER,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order relaxed=mor,no_acquire,non_sc:acquire=moa,acquire,non_sc:seq_cst=mosc,acquire,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ auto atomic_codegen_test(TEMPLATE& atom) +{ + return atom.load(ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}CCTL.IVALL{{.*}} +; SEQ_CST: {{.*}}MEMBAR.SC.[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; BLOCK: {{.*}}LD.E[[SASS_SIZE]].STRONG.{{CTA|SM}} R4, {{.*}} +; NON_BLOCK: {{.*}}LD.E[[SASS_SIZE]].STRONG.[[SASS_SCOPE]] R4, {{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/load_types_128_atomic.cu b/libcudacxx/test/atomic_codegen/sass/load_types_128_atomic.cu new file mode 100644 index 000000000000..aa25f73278f0 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/load_types_128_atomic.cu @@ -0,0 +1,43 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope block=tsb,CTA,block:device=tsd,GPU,non_block:system=tss,SYS,non_block +// %PARAM% TYPE type i128:u128 +// %PARAM% ORDER order mor:moa:mosc +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ auto atomic_codegen_test(cuda::atomic& atom) +{ + return atom.load(ORDER); +} + +// clang-format off +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}MEMBAR.SC.{{.*}} +; SMXX: {{.*}}BSSY [[SYNC:B[0-9]+]], {{.*}} +; BLOCK: {{.*}}ATOM.E.EXCH.STRONG.{{CTA|SM}} PT, [[LOCK_STATE:R[0-9]+]], {{.*\[}}[[BASE_ADDR:R[0-9]+]]{{(\.64)?\+0x10\].*}}, {{R[0-9]+}}{{.*}} +; NON_BLOCK: {{.*}}ATOM.E.EXCH.STRONG.[[SASS_SCOPE]] PT, [[LOCK_STATE:R[0-9]+]], {{.*\[}}[[BASE_ADDR:R[0-9]+]]{{(\.64)?\+0x10\].*}}, {{R[0-9]+}}{{.*}} +; SMXX-DAG: {{.*}}ISETP.NE{{(\.U32)?}}.AND [[LOCK_FREE:P[0-9]+]], PT, [[LOCK_STATE]], 0x1, PT{{.*}} +; NON_BLOCK-DAG: {{.*}}CCTL.IVALL{{.*}} +; SMXX: {{.*}}{{@!?}}[[LOCK_FREE]] BRA{{(\.U)?}} {{.*}} +; SMXX: {{.*}}BSYNC [[SYNC]]{{.*}} +; SMXX: {{.*}}LD.E.128{{(\.SYS)?}} R4, {{.*\[}}[[BASE_ADDR]]{{(\.64)?\].*}} +; SMXX: {{.*}}MEMBAR.ALL.[[SASS_SCOPE]]{{.*}} +; BLOCK: {{.*}}ST.E.STRONG.{{CTA|SM}} {{.*\[}}[[BASE_ADDR]]{{(\.64)?\+0x10\].*}}, RZ{{.*}} +; NON_BLOCK: {{.*}}ST.E.STRONG.[[SASS_SCOPE]] {{.*\[}}[[BASE_ADDR]]{{(\.64)?\+0x10\].*}}, RZ{{.*}} +; SMXX-NEXT: {{.*}}RET.ABS.NODEC{{.*}} + +*/ +// clang-format on diff --git a/libcudacxx/test/atomic_codegen/sass/load_types_128_atomic_ref.cu b/libcudacxx/test/atomic_codegen/sass/load_types_128_atomic_ref.cu new file mode 100644 index 000000000000..63e761a698a4 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/load_types_128_atomic_ref.cu @@ -0,0 +1,47 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope block=tsb,CTA,block:device=tsd,GPU,non_block:system=tss,SYS,non_block +// %PARAM% TYPE type i128:u128 +// %PARAM% ORDER,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order relaxed=mor,no_acquire,non_sc:acquire=moa,acquire,non_sc:seq_cst=mosc,acquire,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ auto atomic_codegen_test(cuda::atomic_ref& atom) +{ + return atom.load(ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX: {{.*}}LD.E.64{{(\.SYS)?}} [[ATOM_ADDR:R[0-9]+]], {{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}CCTL.IVALL{{.*}} +; SEQ_CST: {{.*}}MEMBAR.SC.[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}MEMBAR.SC.{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; BLOCK: {{.*}}LD.E.128.STRONG.{{CTA|SM}} R4, {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}} +; NON_BLOCK: {{.*}}LD.E.128.STRONG.[[SASS_SCOPE]] R4, {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/load_types_8_16_atomic.cu b/libcudacxx/test/atomic_codegen/sass/load_types_8_16_atomic.cu new file mode 100644 index 000000000000..e914ba0ae210 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/load_types_8_16_atomic.cu @@ -0,0 +1,49 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope block=tsb,CTA,block:device=tsd,GPU,non_block:system=tss,SYS,non_block +// %PARAM% TYPE type int8_t:uint8_t:int16_t:uint16_t:f16:bf16 +// %PARAM% ORDER,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order relaxed=mor,no_acquire,non_sc:acquire=moa,acquire,non_sc:seq_cst=mosc,acquire,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include +#include + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ auto atomic_codegen_test(cuda::atomic& atom) +{ + return atom.load(ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}CCTL.IVALL{{.*}} +; SEQ_CST: {{.*}}MEMBAR.SC.[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; BLOCK: {{.*}}LD.E.STRONG.{{CTA|SM}} {{R[0-9]+}}, {{.*}} +; NON_BLOCK: {{.*}}LD.E.STRONG.[[SASS_SCOPE]] {{R[0-9]+}}, {{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/load_types_8_16_atomic_ref.cu b/libcudacxx/test/atomic_codegen/sass/load_types_8_16_atomic_ref.cu new file mode 100644 index 000000000000..46fdb88ae682 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/load_types_8_16_atomic_ref.cu @@ -0,0 +1,50 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope block=tsb,CTA,block:device=tsd,GPU,non_block:system=tss,SYS,non_block +// %PARAM% TYPE,SASS_SIZE type i8=int8_t,.U8:u8=uint8_t,.U8:i16=int16_t,.U16:u16=uint16_t,.U16:f16=f16,.U16:bf16=bf16,.U16 +// %PARAM% ORDER,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order relaxed=mor,no_acquire,non_sc:acquire=moa,acquire,non_sc:seq_cst=mosc,acquire,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include +#include + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ auto atomic_codegen_test(cuda::atomic_ref& atom) +{ + return atom.load(ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX: {{.*}}LD.E.64{{(\.SYS)?}} [[ATOM_ADDR:R[0-9]+]], {{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}CCTL.IVALL{{.*}} +; SEQ_CST: {{.*}}MEMBAR.SC.[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; BLOCK: {{.*}}LD.E[[SASS_SIZE]].STRONG.{{CTA|SM}} {{R[0-9]+}}, {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}} +; NON_BLOCK: {{.*}}LD.E[[SASS_SIZE]].STRONG.[[SASS_SCOPE]] {{R[0-9]+}}, {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/load_volatile_apis.cu b/libcudacxx/test/atomic_codegen/sass/load_volatile_apis.cu new file mode 100644 index 000000000000..6e2e75d18f21 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/load_volatile_apis.cu @@ -0,0 +1,50 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% TEMPLATE,SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE api vcab=vca,tsb,CTA,block:vcad=vca,tsd,GPU,non_block:vcas=vca,tss,SYS,non_block:vcarb=vcar,tsb,CTA,block:vcard=vcar,tsd,GPU,non_block:vcars=vcar,tss,SYS,non_block:vcsa=vcsa,tss,SYS,non_block:vcsar=vcsar,tss,SYS,non_block +// %PARAM% ORDER,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order relaxed=mor,no_acquire,non_sc:acquire=moa,acquire,non_sc:seq_cst=mosc,acquire,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ auto atomic_codegen_test(TEMPLATE& atom) +{ + return atom.load(ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX-NOT: {{.*}}LD.E{{.*}}.STRONG{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}CCTL.IVALL{{.*}} +; SEQ_CST: {{.*}}MEMBAR.SC.[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX-NOT: {{.*}}LD.E{{.*}}.STRONG{{.*}} +; BLOCK: {{.*}}LD.E.STRONG.{{CTA|SM}} R4, {{.*}} +; NON_BLOCK: {{.*}}LD.E.STRONG.[[SASS_SCOPE]] R4, {{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX-NOT: {{.*}}LD.E{{.*}}.STRONG{{.*}} +; SMXX-NOT: {{.*}}LDL{{.*}} +; SMXX-NOT: {{.*}}STL{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/load_volatile_types.cu b/libcudacxx/test/atomic_codegen/sass/load_volatile_types.cu new file mode 100644 index 000000000000..15787062f016 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/load_volatile_types.cu @@ -0,0 +1,51 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% TEMPLATE,SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE api vcad=vca,tsd,GPU,non_block:vcard=vcar,tsd,GPU,non_block +// %PARAM% TYPE,SASS_SIZE type i32=int32_t,:u32=uint32_t,:i64=int64_t,.64:u64=uint64_t,.64:f32=float,:f64=double,.64:ptr1=char*,.64:ptr4=int32_t*,.64 +// %PARAM% ORDER,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order relaxed=mor,no_acquire,non_sc:acquire=moa,acquire,non_sc:seq_cst=mosc,acquire,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ auto atomic_codegen_test(TEMPLATE& atom) +{ + return atom.load(ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX-NOT: {{.*}}LD.E{{.*}}.STRONG{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}CCTL.IVALL{{.*}} +; SEQ_CST: {{.*}}MEMBAR.SC.[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX-NOT: {{.*}}LD.E{{.*}}.STRONG{{.*}} +; BLOCK: {{.*}}LD.E[[SASS_SIZE]].STRONG.{{CTA|SM}} R4, {{.*}} +; NON_BLOCK: {{.*}}LD.E[[SASS_SIZE]].STRONG.[[SASS_SCOPE]] R4, {{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX-NOT: {{.*}}LD.E{{.*}}.STRONG{{.*}} +; SMXX-NOT: {{.*}}LDL{{.*}} +; SMXX-NOT: {{.*}}STL{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/load_volatile_types_128_atomic_ref.cu b/libcudacxx/test/atomic_codegen/sass/load_volatile_types_128_atomic_ref.cu new file mode 100644 index 000000000000..becccabfb183 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/load_volatile_types_128_atomic_ref.cu @@ -0,0 +1,50 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope block=tsb,CTA,block:device=tsd,GPU,non_block:system=tss,SYS,non_block +// %PARAM% TYPE type i128:u128 +// %PARAM% ORDER,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order relaxed=mor,no_acquire,non_sc:acquire=moa,acquire,non_sc:seq_cst=mosc,acquire,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ auto atomic_codegen_test(cuda::atomic_ref& atom) +{ + return atom.load(ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX: {{.*}}LD.E.64{{(\.SYS)?}} [[ATOM_ADDR:R[0-9]+]], {{.*}} +; SMXX-NOT: {{.*}}LD.E{{.*\[}}[[ATOM_ADDR]]{{(\.64)?(\+0x[0-9a-f]+)?\].*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}CCTL.IVALL{{.*}} +; SEQ_CST: {{.*}}MEMBAR.SC.[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}MEMBAR.SC.{{.*}} +; SMXX-NOT: {{.*}}LD.E{{.*\[}}[[ATOM_ADDR]]{{(\.64)?(\+0x[0-9a-f]+)?\].*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; BLOCK: {{.*}}LD.E.128.STRONG.{{CTA|SM}} R4, {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}} +; NON_BLOCK: {{.*}}LD.E.128.STRONG.[[SASS_SCOPE]] R4, {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}LD.E{{.*\[}}[[ATOM_ADDR]]{{(\.64)?(\+0x[0-9a-f]+)?\].*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/load_volatile_types_8_16_atomic.cu b/libcudacxx/test/atomic_codegen/sass/load_volatile_types_8_16_atomic.cu new file mode 100644 index 000000000000..9e8cccce9afd --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/load_volatile_types_8_16_atomic.cu @@ -0,0 +1,52 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope device=tsd,GPU,non_block +// %PARAM% TYPE type int8_t:uint8_t:int16_t:uint16_t:f16:bf16 +// %PARAM% ORDER,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order relaxed=mor,no_acquire,non_sc:acquire=moa,acquire,non_sc:seq_cst=mosc,acquire,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include +#include + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ auto atomic_codegen_test(volatile cuda::atomic& atom) +{ + return atom.load(ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX-NOT: {{.*}}LD.E{{.*}}.STRONG{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}CCTL.IVALL{{.*}} +; SEQ_CST: {{.*}}MEMBAR.SC.[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX-NOT: {{.*}}LD.E{{.*}}.STRONG{{.*}} +; BLOCK: {{.*}}LD.E.STRONG.{{CTA|SM}} {{R[0-9]+}}, {{.*}} +; NON_BLOCK: {{.*}}LD.E.STRONG.[[SASS_SCOPE]] {{R[0-9]+}}, {{.*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX-NOT: {{.*}}LD.E{{.*}}.STRONG{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/load_volatile_types_8_16_atomic_ref.cu b/libcudacxx/test/atomic_codegen/sass/load_volatile_types_8_16_atomic_ref.cu new file mode 100644 index 000000000000..c91c46c2349b --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/load_volatile_types_8_16_atomic_ref.cu @@ -0,0 +1,53 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope device=tsd,GPU,non_block +// %PARAM% TYPE,SASS_SIZE type i8=int8_t,.U8:u8=uint8_t,.U8:i16=int16_t,.U16:u16=uint16_t,.U16:f16=f16,.U16:bf16=bf16,.U16 +// %PARAM% ORDER,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order relaxed=mor,no_acquire,non_sc:acquire=moa,acquire,non_sc:seq_cst=mosc,acquire,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// %FILECHECK% PREFIX_COMBINE non_block,acquire +// clang-format on + +#include +#include + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ auto atomic_codegen_test(cuda::atomic_ref& atom) +{ + return atom.load(ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX: {{.*}}LD.E.64{{(\.SYS)?}} [[ATOM_ADDR:R[0-9]+]], {{.*}} +; SMXX-NOT: {{.*}}LD.E{{.*}}.STRONG{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}CCTL.IVALL{{.*}} +; SEQ_CST: {{.*}}MEMBAR.SC.[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SC-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX-NOT: {{.*}}LD.E{{.*}}.STRONG{{.*}} +; BLOCK: {{.*}}LD.E[[SASS_SIZE]].STRONG.{{CTA|SM}} {{R[0-9]+}}, {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}} +; NON_BLOCK: {{.*}}LD.E[[SASS_SIZE]].STRONG.[[SASS_SCOPE]] {{R[0-9]+}}, {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}} +; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX-NOT: {{.*}}LD.E{{.*}}.STRONG{{.*}} +; SMXX: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/store_apis.cu b/libcudacxx/test/atomic_codegen/sass/store_apis.cu new file mode 100644 index 000000000000..421665fe697a --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/store_apis.cu @@ -0,0 +1,41 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% TEMPLATE,SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE api cab=ca,tsb,CTA,block:cad=ca,tsd,GPU,non_block:cas=ca,tss,SYS,non_block:carb=car,tsb,CTA,block:card=car,tsd,GPU,non_block:cars=car,tss,SYS,non_block:csa=csa,tss,SYS,non_block:csar=csar,tss,SYS,non_block +// %PARAM% ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ORDER order relaxed=mor,,non_seq_cst,no_membar:release=more,ALL,non_seq_cst,membar:seq_cst=mosc,SC,seq_cst,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ void atomic_codegen_test(TEMPLATE& atom, int32_t value) +{ + atom.store(value, ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; BLOCK: {{.*}}ST.E.STRONG.{{CTA|SM}} {{.*}}, {{R[0-9]+}}{{.*}} +; NON_BLOCK: {{.*}}ST.E.STRONG.[[SASS_SCOPE]] {{.*}}, {{R[0-9]+}}{{.*}} +; SMXX-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NEXT: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/store_types.cu b/libcudacxx/test/atomic_codegen/sass/store_types.cu new file mode 100644 index 000000000000..c740962872fd --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/store_types.cu @@ -0,0 +1,42 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% TEMPLATE,SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE api cab=ca,tsb,CTA,block:card=car,tsd,GPU,non_block:csa=csa,tss,SYS,non_block +// %PARAM% TYPE,SASS_SIZE type i32=int32_t,:u32=uint32_t,:i64=int64_t,.64:u64=uint64_t,.64:f32=float,:f64=double,.64:ptr1=char*,.64:ptr4=int32_t*,.64 +// %PARAM% ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ORDER order relaxed=mor,,non_seq_cst,no_membar:release=more,ALL,non_seq_cst,membar:seq_cst=mosc,SC,seq_cst,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ void atomic_codegen_test(TEMPLATE& atom, TYPE value) +{ + atom.store(value, ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; BLOCK: {{.*}}ST.E[[SASS_SIZE]].STRONG.{{CTA|SM}} {{.*}}, {{R[0-9]+}}{{.*}} +; NON_BLOCK: {{.*}}ST.E[[SASS_SIZE]].STRONG.[[SASS_SCOPE]] {{.*}}, {{R[0-9]+}}{{.*}} +; SMXX-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NEXT: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/store_types_128_atomic.cu b/libcudacxx/test/atomic_codegen/sass/store_types_128_atomic.cu new file mode 100644 index 000000000000..35faefd0c627 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/store_types_128_atomic.cu @@ -0,0 +1,43 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope block=tsb,CTA,block:device=tsd,GPU,non_block:system=tss,SYS,non_block +// %PARAM% TYPE type i128:u128 +// %PARAM% ORDER order mor:more:mosc +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ void atomic_codegen_test(cuda::atomic& atom, TYPE value) +{ + atom.store(value, ORDER); +} + +// clang-format off +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}MEMBAR.SC.{{.*}} +; SMXX: {{.*}}BSSY [[SYNC:B[0-9]+]], {{.*}} +; BLOCK: {{.*}}ATOM.E.EXCH.STRONG.{{CTA|SM}} PT, [[LOCK_STATE:R[0-9]+]], {{.*\[}}[[BASE_ADDR:R[0-9]+]]{{(\.64)?\+0x10\].*}}, {{R[0-9]+}}{{.*}} +; NON_BLOCK: {{.*}}ATOM.E.EXCH.STRONG.[[SASS_SCOPE]] PT, [[LOCK_STATE:R[0-9]+]], {{.*\[}}[[BASE_ADDR:R[0-9]+]]{{(\.64)?\+0x10\].*}}, {{R[0-9]+}}{{.*}} +; SMXX-DAG: {{.*}}ISETP.NE{{(\.U32)?}}.AND [[LOCK_FREE:P[0-9]+]], PT, [[LOCK_STATE]], 0x1, PT{{.*}} +; NON_BLOCK-DAG: {{.*}}CCTL.IVALL{{.*}} +; SMXX: {{.*}}{{@!?}}[[LOCK_FREE]] BRA{{(\.U)?}} {{.*}} +; SMXX: {{.*}}BSYNC [[SYNC]]{{.*}} +; SMXX: {{.*}}ST.E.128{{(\.SYS)?}} {{.*\[}}[[BASE_ADDR]]{{(\.64)?\].*}}, {{R[0-9]+}}{{.*}} +; SMXX: {{.*}}MEMBAR.ALL.[[SASS_SCOPE]]{{.*}} +; BLOCK: {{.*}}ST.E.STRONG.{{CTA|SM}} {{.*\[}}[[BASE_ADDR]]{{(\.64)?\+0x10\].*}}, RZ{{.*}} +; NON_BLOCK: {{.*}}ST.E.STRONG.[[SASS_SCOPE]] {{.*\[}}[[BASE_ADDR]]{{(\.64)?\+0x10\].*}}, RZ{{.*}} +; SMXX-NEXT: {{.*}}RET.ABS.NODEC{{.*}} + +*/ +// clang-format on diff --git a/libcudacxx/test/atomic_codegen/sass/store_types_128_atomic_ref.cu b/libcudacxx/test/atomic_codegen/sass/store_types_128_atomic_ref.cu new file mode 100644 index 000000000000..24ed8a8d0aca --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/store_types_128_atomic_ref.cu @@ -0,0 +1,43 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope block=tsb,CTA,block:device=tsd,GPU,non_block:system=tss,SYS,non_block +// %PARAM% TYPE type i128:u128 +// %PARAM% ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ORDER order relaxed=mor,,non_seq_cst,no_membar:release=more,ALL,non_seq_cst,membar:seq_cst=mosc,SC,seq_cst,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ void atomic_codegen_test(cuda::atomic_ref& atom, TYPE value) +{ + atom.store(value, ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX: {{.*}}LD.E.64{{(\.SYS)?}} [[ATOM_ADDR:R[0-9]+]], {{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; BLOCK: {{.*}}ST.E.128.STRONG.{{CTA|SM}} {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}}, {{R[0-9]+}}{{.*}} +; NON_BLOCK: {{.*}}ST.E.128.STRONG.[[SASS_SCOPE]] {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}}, {{R[0-9]+}}{{.*}} +; SMXX-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NEXT: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/store_types_8_16_atomic.cu b/libcudacxx/test/atomic_codegen/sass/store_types_8_16_atomic.cu new file mode 100644 index 000000000000..ec7d3ce92371 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/store_types_8_16_atomic.cu @@ -0,0 +1,45 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope block=tsb,CTA,block:device=tsd,GPU,non_block:system=tss,SYS,non_block +// %PARAM% TYPE type int8_t:uint8_t:int16_t:uint16_t:f16:bf16 +// %PARAM% ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ORDER order relaxed=mor,,non_seq_cst,no_membar:release=more,ALL,non_seq_cst,membar:seq_cst=mosc,SC,seq_cst,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// clang-format on + +#include +#include + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ void atomic_codegen_test(cuda::atomic& atom, TYPE value) +{ + atom.store(value, ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; BLOCK: {{.*}}ST.E.STRONG.{{CTA|SM}}{{.*}} +; NON_BLOCK: {{.*}}ST.E.STRONG.[[SASS_SCOPE]]{{.*}} +; SMXX-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NEXT: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/store_types_8_16_atomic_ref.cu b/libcudacxx/test/atomic_codegen/sass/store_types_8_16_atomic_ref.cu new file mode 100644 index 000000000000..ab19dd210d96 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/store_types_8_16_atomic_ref.cu @@ -0,0 +1,46 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope block=tsb,CTA,block:device=tsd,GPU,non_block:system=tss,SYS,non_block +// %PARAM% TYPE,SASS_SIZE type i8=int8_t,.U8:u8=uint8_t,.U8:i16=int16_t,.U16:u16=uint16_t,.U16:f16=f16,.U16:bf16=bf16,.U16 +// %PARAM% ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ORDER order relaxed=mor,,non_seq_cst,no_membar:release=more,ALL,non_seq_cst,membar:seq_cst=mosc,SC,seq_cst,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// clang-format on + +#include +#include + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ void atomic_codegen_test(cuda::atomic_ref& atom, TYPE value) +{ + atom.store(value, ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX: {{.*}}LD.E.64{{(\.SYS)?}} [[ATOM_ADDR:R[0-9]+]], {{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; BLOCK: {{.*}}ST.E[[SASS_SIZE]].STRONG.{{CTA|SM}} {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}} +; NON_BLOCK: {{.*}}ST.E[[SASS_SIZE]].STRONG.[[SASS_SCOPE]] {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}} +; SMXX-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NEXT: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/store_volatile_apis.cu b/libcudacxx/test/atomic_codegen/sass/store_volatile_apis.cu new file mode 100644 index 000000000000..0aebe4e91201 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/store_volatile_apis.cu @@ -0,0 +1,44 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% TEMPLATE,SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE api vcab=vca,tsb,CTA,block:vcad=vca,tsd,GPU,non_block:vcas=vca,tss,SYS,non_block:vcarb=vcar,tsb,CTA,block:vcard=vcar,tsd,GPU,non_block:vcars=vcar,tss,SYS,non_block:vcsa=vcsa,tss,SYS,non_block:vcsar=vcsar,tss,SYS,non_block +// %PARAM% ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ORDER order relaxed=mor,,non_seq_cst,no_membar:release=more,ALL,non_seq_cst,membar:seq_cst=mosc,SC,seq_cst,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ void atomic_codegen_test(TEMPLATE& atom, int32_t value) +{ + atom.store(value, ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX-NOT: {{.*}}ST.E{{.*}}.STRONG{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX-NOT: {{.*}}ST.E{{.*}}.STRONG{{.*}} +; BLOCK: {{.*}}ST.E.STRONG.{{CTA|SM}} {{.*}}, R6{{.*}} +; NON_BLOCK: {{.*}}ST.E.STRONG.[[SASS_SCOPE]] {{.*}}, R6{{.*}} +; SMXX-NOT: {{.*}}ST.E{{.*}}.STRONG{{.*}} +; SMXX-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NEXT: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/store_volatile_types.cu b/libcudacxx/test/atomic_codegen/sass/store_volatile_types.cu new file mode 100644 index 000000000000..6ce2b3480e0a --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/store_volatile_types.cu @@ -0,0 +1,45 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% TEMPLATE,SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE api vcad=vca,tsd,GPU,non_block:vcard=vcar,tsd,GPU,non_block +// %PARAM% TYPE,SASS_SIZE type i32=int32_t,:u32=uint32_t,:i64=int64_t,.64:u64=uint64_t,.64:f32=float,:f64=double,.64:ptr1=char*,.64:ptr4=int32_t*,.64 +// %PARAM% ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ORDER order relaxed=mor,,non_seq_cst,no_membar:release=more,ALL,non_seq_cst,membar:seq_cst=mosc,SC,seq_cst,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ void atomic_codegen_test(TEMPLATE& atom, TYPE value) +{ + atom.store(value, ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX-NOT: {{.*}}ST.E{{.*}}.STRONG{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX-NOT: {{.*}}ST.E{{.*}}.STRONG{{.*}} +; BLOCK: {{.*}}ST.E[[SASS_SIZE]].STRONG.{{CTA|SM}} {{.*}}, R6{{.*}} +; NON_BLOCK: {{.*}}ST.E[[SASS_SIZE]].STRONG.[[SASS_SCOPE]] {{.*}}, R6{{.*}} +; SMXX-NOT: {{.*}}ST.E{{.*}}.STRONG{{.*}} +; SMXX-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NEXT: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/store_volatile_types_128_atomic_ref.cu b/libcudacxx/test/atomic_codegen/sass/store_volatile_types_128_atomic_ref.cu new file mode 100644 index 000000000000..20ef0ae98ed0 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/store_volatile_types_128_atomic_ref.cu @@ -0,0 +1,46 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope block=tsb,CTA,block:device=tsd,GPU,non_block:system=tss,SYS,non_block +// %PARAM% TYPE type i128:u128 +// %PARAM% ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ORDER order relaxed=mor,,non_seq_cst,no_membar:release=more,ALL,non_seq_cst,membar:seq_cst=mosc,SC,seq_cst,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// clang-format on + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ void atomic_codegen_test(cuda::atomic_ref& atom, TYPE value) +{ + atom.store(value, ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX: {{.*}}LD.E.64{{(\.SYS)?}} [[ATOM_ADDR:R[0-9]+]], {{.*}} +; SMXX-NOT: {{.*}}ST.E{{.*\[}}[[ATOM_ADDR]]{{(\.64)?(\+0x[0-9a-f]+)?\].*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ST.E{{.*\[}}[[ATOM_ADDR]]{{(\.64)?(\+0x[0-9a-f]+)?\].*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; BLOCK: {{.*}}ST.E.128.STRONG.{{CTA|SM}} {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}}, R8{{.*}} +; NON_BLOCK: {{.*}}ST.E.128.STRONG.[[SASS_SCOPE]] {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}}, R8{{.*}} +; SMXX-NOT: {{.*}}ST.E{{.*\[}}[[ATOM_ADDR]]{{(\.64)?(\+0x[0-9a-f]+)?\].*}} +; SMXX-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NEXT: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/store_volatile_types_8_16_atomic.cu b/libcudacxx/test/atomic_codegen/sass/store_volatile_types_8_16_atomic.cu new file mode 100644 index 000000000000..ebad6cd30b23 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/store_volatile_types_8_16_atomic.cu @@ -0,0 +1,48 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope device=tsd,GPU,non_block +// %PARAM% TYPE type int8_t:uint8_t:int16_t:uint16_t:f16:bf16 +// %PARAM% ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ORDER order relaxed=mor,,non_seq_cst,no_membar:release=more,ALL,non_seq_cst,membar:seq_cst=mosc,SC,seq_cst,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// clang-format on + +#include +#include + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ void atomic_codegen_test(volatile cuda::atomic& atom, TYPE value) +{ + atom.store(value, ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX-NOT: {{.*}}ST.E{{.*}}.STRONG{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX-NOT: {{.*}}ST.E{{.*}}.STRONG{{.*}} +; BLOCK: {{.*}}ST.E.STRONG.{{CTA|SM}}{{.*}} +; NON_BLOCK: {{.*}}ST.E.STRONG.[[SASS_SCOPE]]{{.*}} +; SMXX-NOT: {{.*}}ST.E{{.*}}.STRONG{{.*}} +; SMXX-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NEXT: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/atomic_codegen/sass/store_volatile_types_8_16_atomic_ref.cu b/libcudacxx/test/atomic_codegen/sass/store_volatile_types_8_16_atomic_ref.cu new file mode 100644 index 000000000000..39d22a2f8e20 --- /dev/null +++ b/libcudacxx/test/atomic_codegen/sass/store_volatile_types_8_16_atomic_ref.cu @@ -0,0 +1,49 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++ in the CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// clang-format off +// %PARAM% SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE scope device=tsd,GPU,non_block +// %PARAM% TYPE,SASS_SIZE type i8=int8_t,.U8:u8=uint8_t,.U8:i16=int16_t,.U16:u16=uint16_t,.U16:f16=f16,.U16:bf16=bf16,.U16 +// %PARAM% ORDER,SASS_MEMBAR,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ORDER order relaxed=mor,,non_seq_cst,no_membar:release=more,ALL,non_seq_cst,membar:seq_cst=mosc,SC,seq_cst,membar +// %FILECHECK% PREFIX_COMBINE non_block,seq_cst +// clang-format on + +#include +#include + +#include "atomic_codegen_helpers.h" + +extern "C" __device__ void atomic_codegen_test(cuda::atomic_ref& atom, TYPE value) +{ + atom.store(value, ORDER); +} + +/* + +; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX: {{.*}}LD.E.64{{(\.SYS)?}} [[ATOM_ADDR:R[0-9]+]], {{.*}} +; SMXX-NOT: {{.*}}ST.E{{.*}}.STRONG{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} +; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} +; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} +; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} +; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} +; SMXX-NOT: {{.*}}ATOM.{{.*}} +; SMXX-NOT: {{.*}}ST.E{{.*}}.STRONG{{.*}} +; BLOCK: {{.*}}ST.E[[SASS_SIZE]].STRONG.{{CTA|SM}} {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}} +; NON_BLOCK: {{.*}}ST.E[[SASS_SIZE]].STRONG.[[SASS_SCOPE]] {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}} +; SMXX-NOT: {{.*}}ST.E{{.*}}.STRONG{{.*}} +; SMXX-NOT: {{.*}}CCTL.IVALL{{.*}} +; SMXX-NEXT: {{.*}}RET.ABS.NODEC{{.*}} + +*/ diff --git a/libcudacxx/test/cmake/CodegenTest.cmake b/libcudacxx/test/cmake/CodegenTest.cmake index 0fa9a72fa449..c1f35fc05d9e 100644 --- a/libcudacxx/test/cmake/CodegenTest.cmake +++ b/libcudacxx/test/cmake/CodegenTest.cmake @@ -8,6 +8,8 @@ ## ##===----------------------------------------------------------------------===## +include("${CMAKE_CURRENT_LIST_DIR}/../../../cmake/CCCLTestParams.cmake") + option( LIBCUDACXX_REQUIRE_CODEGEN_TEST_TOOLS "Fail configuration when tools required by libcu++ codegen tests are missing." @@ -15,6 +17,12 @@ option( ) function(libcudacxx_codegen_check_tools out_var) + if (LIBCUDACXX_REQUIRE_CODEGEN_TEST_TOOLS) + set(require_codegen_tools TRUE) + else() + set(require_codegen_tools FALSE) + endif() + get_property( tools_available GLOBAL @@ -51,7 +59,7 @@ function(libcudacxx_codegen_check_tools out_var) if (missing_tools) list(JOIN missing_tools ", " missing_tools) - if (LIBCUDACXX_REQUIRE_CODEGEN_TEST_TOOLS) + if (require_codegen_tools) message( FATAL_ERROR "Tools required by libcu++ codegen tests were not found: ${missing_tools}" @@ -80,6 +88,29 @@ set( "${CMAKE_CURRENT_LIST_DIR}/../codegen/dump_and_check.bash" ) +# Return the concrete CUDA architectures requested for codegen testing. +# Virtual-only architectures do not produce the SASS consumed by these tests; +# special values such as "native" and "all" do not identify a stable FileCheck +# prefix and should be replaced with explicit architectures by the caller. +function(libcudacxx_codegen_get_cuda_architectures out_var) + set(cuda_architectures) + foreach (arch IN LISTS CMAKE_CUDA_ARCHITECTURES) + if (arch MATCHES "^([0-9]+[af]?)(-real)?$") + list(APPEND cuda_architectures "${CMAKE_MATCH_1}") + endif() + endforeach() + list(REMOVE_DUPLICATES cuda_architectures) + + if (NOT cuda_architectures) + message( + FATAL_ERROR + "Libcu++ FileCheck codegen tests require at least one explicit real CUDA architecture; got '${CMAKE_CUDA_ARCHITECTURES}'" + ) + endif() + + set(${out_var} "${cuda_architectures}" PARENT_SCOPE) +endfunction() + function(libcudacxx_codegen_set_cuda_arch target_name arch) if (arch MATCHES "[af]$") set_target_properties(${target_name} PROPERTIES CUDA_ARCHITECTURES OFF) @@ -159,6 +190,94 @@ function( set(${out_has_specific_checks} "${has_specific_checks}" PARENT_SCOPE) endfunction() +function( + libcudacxx_codegen_resolve_filecheck_prefixes + out_target_prefixes + out_filecheck_prefixes + test_path + input_prefixes +) + file(READ "${test_path}" test_contents) + file( + STRINGS "${test_path}" + prefix_combine_directives + REGEX "%FILECHECK%[ ]+PREFIX_COMBINE" + ) + + if (NOT prefix_combine_directives) + set(${out_target_prefixes} "${input_prefixes}" PARENT_SCOPE) + set(${out_filecheck_prefixes} "${input_prefixes}" PARENT_SCOPE) + return() + endif() + + string(REPLACE "," ";" active_prefixes "${input_prefixes}") + set(combination_components) + set(active_combinations) + + foreach (directive IN LISTS prefix_combine_directives) + if ( + NOT + directive + MATCHES + "^[ ]*//[ ]+%FILECHECK%[ ]+PREFIX_COMBINE[ ]+([A-Za-z][A-Za-z0-9_-]*([ ]*,[ ]*[A-Za-z][A-Za-z0-9_-]*)+)[ ]*$" + ) + message( + FATAL_ERROR + "Malformed PREFIX_COMBINE directive in ${test_path}: ${directive}" + ) + endif() + + set(components "${CMAKE_MATCH_1}") + string(REPLACE " " "" components "${components}") + string(REPLACE "," ";" components "${components}") + set(normalized_components) + set(combination_is_active TRUE) + foreach (component IN LISTS components) + string(TOUPPER "${component}" component) + list(APPEND normalized_components "${component}") + list(APPEND combination_components "${component}") + if (NOT component IN_LIST active_prefixes) + set(combination_is_active FALSE) + endif() + endforeach() + + if (combination_is_active) + list(JOIN normalized_components "_" combined_prefix) + list(APPEND active_combinations "${combined_prefix}") + endif() + endforeach() + + list(REMOVE_DUPLICATES combination_components) + list(REMOVE_DUPLICATES active_combinations) + + # A prefix used only as a PREFIX_COMBINE component is an input to a combined + # prefix, not a standalone FileCheck prefix. This permits semantic markers + # such as ACQUIRE without requiring a dummy ACQUIRE directive. + set(standalone_prefixes) + foreach (prefix IN LISTS active_prefixes) + if (prefix IN_LIST combination_components) + string( + REGEX MATCH + "; ${prefix}(:|-[A-Z]+:)" + has_standalone_directive + "${test_contents}" + ) + if (NOT has_standalone_directive) + continue() + endif() + endif() + list(APPEND standalone_prefixes "${prefix}") + endforeach() + + set(filecheck_prefixes ${standalone_prefixes} ${active_combinations}) + list(REMOVE_DUPLICATES standalone_prefixes) + list(REMOVE_DUPLICATES filecheck_prefixes) + list(JOIN standalone_prefixes "," standalone_prefixes) + list(JOIN filecheck_prefixes "," filecheck_prefixes) + set(${out_target_prefixes} "${standalone_prefixes}" PARENT_SCOPE) + set(${out_filecheck_prefixes} "${filecheck_prefixes}" PARENT_SCOPE) +endfunction() + function( libcudacxx_codegen_add_check_target target_path @@ -166,9 +285,34 @@ function( test_path check_prefixes ) - string(REGEX REPLACE "[^A-Za-z0-9_]" "_" check_suffix "${check_prefixes}") + cmake_parse_arguments( + arg + "" + "DUMP_MODE;DUMP_FUNCTIONS" + "CHECK_DEFINITIONS" + ${ARGN} + ) + + libcudacxx_codegen_resolve_filecheck_prefixes( + target_check_prefixes + filecheck_prefixes + "${test_path}" + "${check_prefixes}" + ) + string( + REGEX REPLACE + "[^A-Za-z0-9_]" + "_" + check_suffix + "${target_check_prefixes}" + ) set(check_target_name "${target_path}.${check_suffix}.check") + set(filecheck_definitions) + foreach (definition IN LISTS arg_CHECK_DEFINITIONS) + list(APPEND filecheck_definitions "-D${definition}") + endforeach() + add_custom_target( ${check_target_name} DEPENDS ${target_name} @@ -176,13 +320,15 @@ function( COMMAND "${CMAKE_COMMAND}" -E env "CUOBJDUMP=${libcudacxx_codegen_cuobjdump}" + "CUOBJDUMP_FUNCTIONS=${arg_DUMP_FUNCTIONS}" "FILECHECK=${libcudacxx_codegen_filecheck}" "${libcudacxx_codegen_bash}" "${libcudacxx_codegen_dump_and_check}" $ "${test_path}" - "${check_prefixes}" - ${ARGN} + "${filecheck_prefixes}" + "${arg_DUMP_MODE}" + ${filecheck_definitions} # gersemi: on ) cccl_ensure_metatargets(${check_target_name}) @@ -197,8 +343,10 @@ function(libcudacxx_codegen_add_test) CODE_KIND ARCH TEST + VARIANT + DUMP_FUNCTIONS ) - set(multi_value_args CHECK_PREFIXES COMPILE_DEFINITIONS) + set(multi_value_args CHECK_PREFIXES CHECK_DEFINITIONS COMPILE_DEFINITIONS) cmake_parse_arguments( arg "${options}" @@ -214,6 +362,10 @@ function(libcudacxx_codegen_add_test) target_name "${arg_TARGET_PREFIX}_${arg_CODE_KIND}_sm${arg_ARCH}_${test_name}" ) + if (arg_VARIANT) + string(APPEND test_target_path ".${arg_VARIANT}") + string(APPEND target_name ".${arg_VARIANT}") + endif() add_library(${target_name} STATIC "${arg_TEST}") libcudacxx_codegen_set_cuda_arch(${target_name} "${arg_ARCH}") @@ -230,7 +382,7 @@ function(libcudacxx_codegen_add_test) ) endif() - set(dump_options) + set(dump_mode --dump-ptx) if (arg_CODE_KIND STREQUAL "ptx") # Clang stopped emitting PTX in clang20. Add flags to re-enable it. if ( @@ -243,7 +395,7 @@ function(libcudacxx_codegen_add_test) ) endif() else() - list(APPEND dump_options --dump-sass) + set(dump_mode --dump-sass) if (arg_SEPARABLE_COMPILATION) # CMake supplies the compile phase, so request relocatable device code # without adding a second compile-phase option via -dc. @@ -257,11 +409,53 @@ function(libcudacxx_codegen_add_test) ${target_name} "${arg_TEST}" "${check_prefix}" - ${dump_options} + DUMP_MODE "${dump_mode}" + DUMP_FUNCTIONS "${arg_DUMP_FUNCTIONS}" + CHECK_DEFINITIONS ${arg_CHECK_DEFINITIONS} ) endforeach() endfunction() +# Separate variant definitions used as codegen test metadata from ordinary +# definitions. FILECHECK_PREFIX_= adds the upper-case form of +# to the FileCheck invocation for that variant; all other definitions +# are available to both the compiler and FileCheck. +function( + libcudacxx_codegen_get_variant_options + out_compile_definitions + out_check_definitions + out_check_prefixes +) + set(compile_definitions) + set(check_definitions) + set(check_prefixes) + + foreach (definition IN LISTS ARGN) + if (definition MATCHES "^FILECHECK_PREFIX_[A-Za-z0-9_]+=(.*)$") + set(check_prefix "${CMAKE_MATCH_1}") + if (check_prefix STREQUAL "") + continue() + endif() + if (NOT check_prefix MATCHES "^[A-Za-z][A-Za-z0-9_-]*$") + message( + FATAL_ERROR + "Invalid FileCheck prefix '${check_prefix}' in '${definition}'" + ) + endif() + string(TOUPPER "${check_prefix}" check_prefix) + list(APPEND check_prefixes "${check_prefix}") + else() + list(APPEND compile_definitions "${definition}") + list(APPEND check_definitions "${definition}") + endif() + endforeach() + + list(REMOVE_DUPLICATES check_prefixes) + set(${out_compile_definitions} "${compile_definitions}" PARENT_SCOPE) + set(${out_check_definitions} "${check_definitions}" PARENT_SCOPE) + set(${out_check_prefixes} "${check_prefixes}" PARENT_SCOPE) +endfunction() + function(libcudacxx_codegen_add_ptx_tests) set(options) set(one_value_args AGGREGATE_TARGET TARGET_PREFIX ARCH) @@ -289,8 +483,8 @@ endfunction() function(libcudacxx_codegen_add_sass_tests) set(options) - set(one_value_args AGGREGATE_TARGET TARGET_PREFIX) - set(multi_value_args ARCHITECTURES TESTS COMPILE_DEFINITIONS) + set(one_value_args AGGREGATE_TARGET TARGET_PREFIX DUMP_FUNCTIONS) + set(multi_value_args ARCHITECTURES CHECK_PREFIXES TESTS COMPILE_DEFINITIONS) cmake_parse_arguments( arg "${options}" @@ -300,7 +494,24 @@ function(libcudacxx_codegen_add_sass_tests) ) foreach (test_path IN LISTS arg_TESTS) + set_property( + DIRECTORY + APPEND + PROPERTY CMAKE_CONFIGURE_DEPENDS "${test_path}" + ) file(READ "${test_path}" test_contents) + cccl_parse_variant_params( + "${test_path}" + num_variants + variant_labels + variant_definitions + ) + cccl_log_variant_params( + "${test_path}" + ${num_variants} + variant_labels + variant_definitions + ) string( REGEX MATCH "; SM[0-9]+[af]-PLUS(:|-[A-Z]+:)" @@ -315,13 +526,17 @@ function(libcudacxx_codegen_add_sass_tests) endif() set(test_archs) + set(has_arch_specific_checks FALSE) foreach (arch IN LISTS arg_ARCHITECTURES) libcudacxx_codegen_get_sass_check_prefixes( check_prefixes - has_arch_specific_checks + arch_has_specific_checks "${test_contents}" "${arch}" ) + if (arch_has_specific_checks) + set(has_arch_specific_checks TRUE) + endif() if (NOT "${check_prefixes}" STREQUAL "SMXX") list(APPEND test_archs "${arch}") endif() @@ -351,17 +566,72 @@ function(libcudacxx_codegen_add_sass_tests) "${test_contents}" "${arch}" ) + string(REPLACE "," ";" common_check_prefixes "${check_prefixes}") + foreach (check_prefix IN LISTS arg_CHECK_PREFIXES) + string( + REGEX MATCH + "; ${check_prefix}(:|-[A-Z]+:)" + has_check_prefix + "${test_contents}" + ) + if (has_check_prefix) + list(APPEND common_check_prefixes "${check_prefix}") + endif() + endforeach() + list(REMOVE_DUPLICATES common_check_prefixes) - libcudacxx_codegen_add_test( - ${test_options} - AGGREGATE_TARGET ${arg_AGGREGATE_TARGET} - TARGET_PREFIX ${arg_TARGET_PREFIX} - CODE_KIND sass - ARCH ${arch} - TEST "${test_path}" - CHECK_PREFIXES "${check_prefixes}" - COMPILE_DEFINITIONS ${arg_COMPILE_DEFINITIONS} - ) + if (num_variants EQUAL 0) + list(JOIN common_check_prefixes "," combined_check_prefixes) + libcudacxx_codegen_add_test( + ${test_options} + AGGREGATE_TARGET ${arg_AGGREGATE_TARGET} + TARGET_PREFIX ${arg_TARGET_PREFIX} + CODE_KIND sass + ARCH ${arch} + TEST "${test_path}" + DUMP_FUNCTIONS "${arg_DUMP_FUNCTIONS}" + CHECK_PREFIXES "${combined_check_prefixes}" + COMPILE_DEFINITIONS ${arg_COMPILE_DEFINITIONS} + ) + else() + math(EXPR last_variant "${num_variants} - 1") + foreach (variant_index RANGE ${last_variant}) + cccl_get_variant_data( + variant_labels + variant_definitions + ${variant_index} + variant_label + definitions + ) + libcudacxx_codegen_get_variant_options( + variant_compile_definitions + variant_check_definitions + variant_check_prefixes + ${definitions} + ) + + set(combined_check_prefixes ${common_check_prefixes}) + list(APPEND combined_check_prefixes ${variant_check_prefixes}) + list(REMOVE_DUPLICATES combined_check_prefixes) + list(JOIN combined_check_prefixes "," combined_check_prefixes) + + libcudacxx_codegen_add_test( + ${test_options} + AGGREGATE_TARGET ${arg_AGGREGATE_TARGET} + TARGET_PREFIX ${arg_TARGET_PREFIX} + CODE_KIND sass + ARCH ${arch} + TEST "${test_path}" + VARIANT "${variant_label}" + DUMP_FUNCTIONS "${arg_DUMP_FUNCTIONS}" + CHECK_PREFIXES "${combined_check_prefixes}" + CHECK_DEFINITIONS ${variant_check_definitions} + COMPILE_DEFINITIONS + ${arg_COMPILE_DEFINITIONS} + ${variant_compile_definitions} + ) + endforeach() + endif() endforeach() endforeach() endfunction() diff --git a/libcudacxx/test/codegen/dump_and_check.bash b/libcudacxx/test/codegen/dump_and_check.bash index e837b37e4fbe..8ac638b94f97 100755 --- a/libcudacxx/test/codegen/dump_and_check.bash +++ b/libcudacxx/test/codegen/dump_and_check.bash @@ -1,13 +1,31 @@ #!/usr/bin/env bash set -euo pipefail -## Usage: dump_and_check test.a test.cu PREFIXES [cuobjdump-mode] +## Usage: dump_and_check test.a test.cu PREFIXES cuobjdump-mode [FileCheck-options...] input_archive="${1}" input_testfile="${2}" input_prefix="${3}" -dump_mode="${4:---dump-ptx}" +shift 3 +dump_mode="${1:---dump-ptx}" +if (( $# > 0 )); then + shift +fi +dump_functions="${CUOBJDUMP_FUNCTIONS:-}" filecheck="${FILECHECK:-FileCheck}" cuobjdump="${CUOBJDUMP:-cuobjdump}" -"${cuobjdump}" "${dump_mode}" "${input_archive}" \ - | "${filecheck}" --match-full-lines --check-prefixes="${input_prefix}" "${input_testfile}" +cuobjdump_options=("${dump_mode}") +if [[ -n "${dump_functions}" ]]; then + cuobjdump_options+=(--function "${dump_functions}") +fi + +if [[ "${dump_mode}" == "--dump-sass" ]]; then + # cuobjdump prints each instruction's control word on a separate line. Remove + # those lines so FileCheck -NEXT describes adjacent SASS instructions. + "${cuobjdump}" "${cuobjdump_options[@]}" "${input_archive}" \ + | sed -E '\|^[[:space:]]*/\* 0x[[:xdigit:]]+ \*/[[:space:]]*$|d' \ + | "${filecheck}" --match-full-lines --check-prefixes="${input_prefix}" "$@" "${input_testfile}" +else + "${cuobjdump}" "${cuobjdump_options[@]}" "${input_archive}" \ + | "${filecheck}" --match-full-lines --check-prefixes="${input_prefix}" "$@" "${input_testfile}" +fi diff --git a/libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/atomic_ref_volatile.pass.cpp b/libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/atomic_ref_volatile.pass.cpp new file mode 100644 index 000000000000..0d415ccf6e5d --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/atomic_ref_volatile.pass.cpp @@ -0,0 +1,58 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: libcpp-has-no-threads, pre-sm-60 +// UNSUPPORTED: windows && pre-sm-70 + +// UNSUPPORTED: force-tile +// error: asm statement is unsupported in tile code + +#include +#include +#include +#include + +#include "test_macros.h" + +template +TEST_HOST_DEVICE_FUNC void test() +{ + static_assert(cuda::std::is_same_v); + static_assert(AtomicRef::is_always_lock_free); + + volatile int value = 0; + AtomicRef atom(value); + + atom.store(1, cuda::std::memory_order_release); + assert(atom.load(cuda::std::memory_order_acquire) == 1); + assert(atom.exchange(2, cuda::std::memory_order_acq_rel) == 1); + + int expected = 2; + assert(atom.compare_exchange_strong(expected, 3, cuda::std::memory_order_acq_rel, cuda::std::memory_order_acquire)); + assert(expected == 2); + + expected = 2; + assert(!atom.compare_exchange_strong(expected, 4, cuda::std::memory_order_seq_cst)); + assert(expected == 3); + + expected = 3; + while (!atom.compare_exchange_weak(expected, 4, cuda::std::memory_order_relaxed)) + { + } + assert(atom.load(cuda::std::memory_order_relaxed) == 4); +} + +int main(int, char**) +{ + test>(); + test>(); + test>(); + test>(); + + return 0; +} diff --git a/libcudacxx/test/simd_codegen/CMakeLists.txt b/libcudacxx/test/simd_codegen/CMakeLists.txt index 376a724cc420..f8de6e600d21 100644 --- a/libcudacxx/test/simd_codegen/CMakeLists.txt +++ b/libcudacxx/test/simd_codegen/CMakeLists.txt @@ -8,8 +8,12 @@ ## ##===----------------------------------------------------------------------===## -add_custom_target(libcudacxx.test.simd.ptx) -add_custom_target(libcudacxx.test.simd.sass) +if (libcudacxx_codegen_filecheck_simd_ptx) + add_custom_target(libcudacxx.test.simd.ptx) +endif() +if (libcudacxx_codegen_filecheck_simd_sass) + add_custom_target(libcudacxx.test.simd.sass) +endif() #----------------------------------------------------------------------------------------------------------------------- # SETUP: skip unsupported compilers, find tools, set up CUDA architectures @@ -51,44 +55,36 @@ if ("Clang" STREQUAL "${CMAKE_CUDA_COMPILER_ID}") return() endif() -set(simd_codegen_sass_cuda_archs 80 90) -if ( - CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 12.8 - AND NOT CMAKE_CUDA_COMPILER_ID STREQUAL Clang -) - list(APPEND simd_codegen_sass_cuda_archs 100 120) -endif() +libcudacxx_codegen_get_cuda_architectures(simd_codegen_cuda_archs) -if ( - CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 12.9 - AND NOT CMAKE_CUDA_COMPILER_ID STREQUAL Clang -) - list(APPEND simd_codegen_sass_cuda_archs 103 120f) +if (libcudacxx_codegen_filecheck_simd_ptx) + add_subdirectory(load_store) endif() -set(simd_codegen_sass_tests) -file( - GLOB simd_codegen_sass_tests - "floating_point/*.cu" - "integer/*.cu" - "min_max/*.cu" -) - -# The 8-bit arithmetic and min/max tests require the PTX ISA 9.2 8-bit SIMD -# path available with CUDA 13.2 and newer. -if (CMAKE_CUDA_COMPILER_VERSION VERSION_LESS 13.2) - list( - REMOVE_ITEM simd_codegen_sass_tests - "${CMAKE_CURRENT_SOURCE_DIR}/integer/arithmetic_u8x4.cu" - "${CMAKE_CURRENT_SOURCE_DIR}/min_max/min_max_i8x4.cu" +if (libcudacxx_codegen_filecheck_simd_sass) + set(simd_codegen_sass_tests) + file( + GLOB simd_codegen_sass_tests + CONFIGURE_DEPENDS + "floating_point/*.cu" + "integer/*.cu" + "min_max/*.cu" ) -endif() -add_subdirectory(load_store) + # The 8-bit arithmetic and min/max tests require the PTX ISA 9.2 8-bit SIMD + # path available with CUDA 13.2 and newer. + if (CMAKE_CUDA_COMPILER_VERSION VERSION_LESS 13.2) + list( + REMOVE_ITEM simd_codegen_sass_tests + "${CMAKE_CURRENT_SOURCE_DIR}/integer/arithmetic_u8x4.cu" + "${CMAKE_CURRENT_SOURCE_DIR}/min_max/min_max_i8x4.cu" + ) + endif() -libcudacxx_codegen_add_sass_tests( - AGGREGATE_TARGET libcudacxx.test.simd.sass - TARGET_PREFIX simd_codegen - ARCHITECTURES ${simd_codegen_sass_cuda_archs} - TESTS ${simd_codegen_sass_tests} -) + libcudacxx_codegen_add_sass_tests( + AGGREGATE_TARGET libcudacxx.test.simd.sass + TARGET_PREFIX simd_codegen + ARCHITECTURES ${simd_codegen_cuda_archs} + TESTS ${simd_codegen_sass_tests} + ) +endif() diff --git a/libcudacxx/test/simd_codegen/load_store/CMakeLists.txt b/libcudacxx/test/simd_codegen/load_store/CMakeLists.txt index 8e6748913222..a7c2a3fb5967 100644 --- a/libcudacxx/test/simd_codegen/load_store/CMakeLists.txt +++ b/libcudacxx/test/simd_codegen/load_store/CMakeLists.txt @@ -20,40 +20,22 @@ set( "${CMAKE_CURRENT_SOURCE_DIR}/load_store_f64.cu" ) -set(simd_codegen_default_prefixes SMXXX) -libcudacxx_codegen_add_ptx_tests( - AGGREGATE_TARGET libcudacxx.test.simd.ptx - TARGET_PREFIX simd_codegen - ARCH 80 - CHECK_PREFIXES ${simd_codegen_default_prefixes} - TESTS ${simd_codegen_load_store_tests} -) -libcudacxx_codegen_add_ptx_tests( - AGGREGATE_TARGET libcudacxx.test.simd.ptx - TARGET_PREFIX simd_codegen - ARCH 90 - CHECK_PREFIXES ${simd_codegen_default_prefixes} - TESTS ${simd_codegen_load_store_tests} -) - -# SM100/SM120 add 32-byte vectorized checks. -if ( - CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 12.9 - AND NOT CMAKE_CUDA_COMPILER_ID STREQUAL Clang -) - set(simd_codegen_sm100_prefixes SMXXX SM100-PLUS) - libcudacxx_codegen_add_ptx_tests( - AGGREGATE_TARGET libcudacxx.test.simd.ptx - TARGET_PREFIX simd_codegen - ARCH 100 - CHECK_PREFIXES ${simd_codegen_sm100_prefixes} - TESTS ${simd_codegen_load_store_tests} +foreach (arch IN LISTS simd_codegen_cuda_archs) + set(simd_codegen_check_prefixes SMXXX) + if ( + arch MATCHES "^([0-9]+)" + AND CMAKE_MATCH_1 GREATER_EQUAL 100 + AND CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 12.9 + AND NOT CMAKE_CUDA_COMPILER_ID STREQUAL Clang ) + list(APPEND simd_codegen_check_prefixes SM100-PLUS) + endif() + libcudacxx_codegen_add_ptx_tests( AGGREGATE_TARGET libcudacxx.test.simd.ptx TARGET_PREFIX simd_codegen - ARCH 120 - CHECK_PREFIXES ${simd_codegen_sm100_prefixes} + ARCH "${arch}" + CHECK_PREFIXES ${simd_codegen_check_prefixes} TESTS ${simd_codegen_load_store_tests} ) -endif() +endforeach()