From 67c051425b2e7d560b593ebaa0a4117b7074f335 Mon Sep 17 00:00:00 2001 From: "pre-commit-ci[bot]" <66853113+pre-commit-ci[bot]@users.noreply.github.com> Date: Mon, 6 Jul 2026 21:29:37 +0000 Subject: [PATCH 1/3] [pre-commit.ci] pre-commit autoupdate MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit updates: - [github.com/pre-commit/pre-commit-hooks: v5.0.0 → v6.0.0](https://github.com/pre-commit/pre-commit-hooks/compare/v5.0.0...v6.0.0) - [github.com/pre-commit/mirrors-clang-format: v20.1.7 → v22.1.5](https://github.com/pre-commit/mirrors-clang-format/compare/v20.1.7...v22.1.5) - [github.com/astral-sh/ruff-pre-commit: v0.12.2 → v0.15.20](https://github.com/astral-sh/ruff-pre-commit/compare/v0.12.2...v0.15.20) - [github.com/codespell-project/codespell: v2.4.1 → v2.4.2](https://github.com/codespell-project/codespell/compare/v2.4.1...v2.4.2) - [github.com/pre-commit/mirrors-mypy: v1.16.1 → v2.1.0](https://github.com/pre-commit/mirrors-mypy/compare/v1.16.1...v2.1.0) --- .pre-commit-config.yaml | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml index a83defe6090..3dd5e6b4115 100644 --- a/.pre-commit-config.yaml +++ b/.pre-commit-config.yaml @@ -11,7 +11,7 @@ ci: repos: - repo: https://github.com/pre-commit/pre-commit-hooks - rev: v5.0.0 + rev: v6.0.0 hooks: - id: end-of-file-fixer - id: mixed-line-ending @@ -40,7 +40,7 @@ repos: ) - repo: https://github.com/pre-commit/mirrors-clang-format - rev: v20.1.7 + rev: v22.1.5 hooks: - id: clang-format types_or: [file] @@ -62,7 +62,7 @@ repos: # TODO/REMINDER: add the Ruff vscode extension to the devcontainers # Ruff, the Python auto-correcting linter/formatter written in Rust - repo: https://github.com/astral-sh/ruff-pre-commit - rev: v0.12.2 + rev: v0.15.20 hooks: - id: ruff # linter - id: ruff-format # formatter @@ -74,7 +74,7 @@ repos: - id: gersemi - repo: https://github.com/codespell-project/codespell - rev: v2.4.1 + rev: v2.4.2 hooks: - id: codespell additional_dependencies: [tomli] @@ -86,7 +86,7 @@ repos: ) - repo: https://github.com/pre-commit/mirrors-mypy - rev: 'v1.16.1' + rev: 'v2.1.0' hooks: - id: mypy # numpy is intentionally not installed here: numpy>=2.3 ships PEP 695 From 855d640b2ff94c09336ad6bc8c37463521ccc840 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Tue, 14 Jul 2026 11:59:47 +0200 Subject: [PATCH 2/3] Run pre-commit run --all-files --- benchmarks/scripts/cccl/bench/search.py | 14 +- benchmarks/scripts/sol.py | 4 +- c/parallel/test/test_merge_sort.cpp | 4 +- c/parallel/test/test_reduce.cpp | 4 +- c/parallel/test/test_scan.cpp | 4 +- c/parallel/test/test_transform.cpp | 4 +- c/parallel/test/test_unique_by_key.cpp | 6 +- c/parallel/test/test_util.h | 16 +- cub/cub/agent/agent_batched_topk.cuh | 2 +- cub/cub/agent/agent_select_if.cuh | 4 +- cub/cub/agent/agent_three_way_partition.cuh | 2 +- cub/cub/block/block_radix_sort.cuh | 3 +- cub/cub/block/radix_rank_sort_operations.cuh | 9 +- cub/cub/detail/fast_modulo_division.cuh | 2 +- cub/cub/device/device_scan.cuh | 8 +- .../dispatch/dispatch_adjacent_difference.cuh | 34 ++-- .../device/dispatch/dispatch_batch_memcpy.cuh | 6 +- .../device/dispatch/dispatch_histogram.cuh | 12 +- .../dispatch/dispatch_reduce_by_key.cuh | 24 +-- cub/cub/device/dispatch/dispatch_rle.cuh | 14 +- cub/cub/device/dispatch/dispatch_scan.cuh | 10 +- .../device/dispatch/dispatch_scan_by_key.cuh | 54 +++---- .../dispatch_segmented_radix_sort.cuh | 2 +- .../dispatch/dispatch_segmented_reduce.cuh | 32 ++-- .../dispatch/dispatch_segmented_scan.cuh | 22 +-- .../device/dispatch/dispatch_select_if.cuh | 16 +- .../dispatch/dispatch_three_way_partition.cuh | 24 +-- .../dispatch/kernels/kernel_histogram.cuh | 8 +- .../dispatch/kernels/kernel_radix_sort.cuh | 38 ++--- .../kernels/kernel_segmented_sort.cuh | 72 ++++----- .../device/dispatch/tuning/tuning_merge.cuh | 8 +- .../dispatch/tuning/tuning_segmented_sort.cuh | 28 ++-- cub/cub/util_device.cuh | 4 +- cub/cub/util_ptx.cuh | 4 +- cub/cub/warp/specializations/warp_redux.cuh | 4 +- .../example_device_radix_sort_custom.cu | 3 +- cub/test/catch2_test_block_reduce.cu | 2 +- cub/test/catch2_test_device_partition_env.cu | 4 +- ...est_device_radix_sort_custom_policy_hub.cu | 16 +- cub/test/catch2_test_device_radix_sort_env.cu | 18 +-- ..._device_reduce_by_key_custom_policy_hub.cu | 20 +-- ...atch2_test_device_reduce_by_key_env_api.cu | 2 +- cub/test/catch2_test_device_reduce_env.cu | 102 ++++++------ ...atch2_test_device_run_length_encode_env.cu | 4 +- .../catch2_test_device_scan_by_key_env.cu | 2 +- .../catch2_test_device_scan_by_key_env_api.cu | 2 +- cub/test/catch2_test_device_scan_env.cu | 2 +- ..._segmented_radix_sort_custom_policy_hub.cu | 30 ++-- ...vice_segmented_reduce_custom_policy_hub.cu | 18 +-- .../catch2_test_device_segmented_scan_api.cu | 8 +- cub/test/catch2_test_device_select_env.cu | 4 +- cub/test/catch2_test_device_select_env_api.cu | 2 +- ...e_three_way_partition_custom_policy_hub.cu | 18 +-- cub/test/test_block_radix_rank.cu | 2 +- .../experimental/__copy/copy_optimized.cuh | 18 +-- .../cuda/experimental/__cuco/hyperloglog.cuh | 12 +- .../__execution/get_completion_signatures.cuh | 2 +- .../__graph/graph_builder_ref.cuh | 5 +- cudax/test/execution/test_bulk.cu | 48 +++--- cudax/test/execution/test_continues_on.cu | 6 +- cudax/test/execution/test_let_value.cu | 22 +-- cudax/test/execution/test_on.cu | 50 +++--- cudax/test/execution/test_starts_on.cu | 2 +- cudax/test/execution/test_task_scheduler.cu | 2 +- cudax/test/execution/test_when_all.cu | 2 +- cudax/test/stf/examples/07-cholesky-redux.cu | 10 +- cudax/test/stf/graph/epoch.cu | 4 +- .../cuda/__barrier/barrier_block_scope.h | 8 +- libcudacxx/include/cuda/__device/arch_id.h | 4 +- .../generated/clusterlaunchcontrol.h | 19 +-- .../cp_async_bulk_tensor_multicast.h | 16 +- .../__ptx/instructions/generated/tcgen05_ld.h | 148 +++++++++--------- .../instructions/generated/tcgen05_mma.h | 8 +- .../__ptx/instructions/generated/tcgen05_st.h | 148 +++++++++--------- .../include/cuda/__utility/__basic_any/rtti.h | 2 +- .../std/__algorithm/iterator_operations.h | 6 +- .../include/cuda/std/__algorithm/sort.h | 6 +- libcudacxx/include/cuda/std/__cccl/compiler.h | 12 +- .../include/cuda/std/__cccl/host_std_lib.h | 2 +- .../include/cuda/std/__cccl/preprocessor.h | 16 +- .../include/cuda/std/__cccl/visibility.h | 2 +- .../include/cuda/std/__chrono/duration.h | 3 +- .../include/cuda/std/__chrono/time_point.h | 8 +- .../include/cuda/std/__complex/complex.h | 20 +-- .../cuda/std/__concepts/concept_macros.h | 26 +-- libcudacxx/include/cuda/std/__format/buffer.h | 6 +- .../include/cuda/std/__iterator/concepts.h | 14 +- .../include/cuda/std/__memory/construct_at.h | 2 +- .../std/__memory/uninitialized_algorithms.h | 4 +- .../include/cuda/std/__memory/unique_ptr.h | 9 +- .../include/cuda/std/__pstl/cuda/merge.h | 2 +- .../cuda/std/__random/cauchy_distribution.h | 4 +- .../cuda/std/__random/fisher_f_distribution.h | 4 +- .../cuda/std/__random/gamma_distribution.h | 4 +- .../include/cuda/std/__tuple_dir/tuple.h | 8 +- .../include/cuda/std/__type_traits/promote.h | 2 +- .../cuda/std/__utility/constant_wrapper.h | 67 ++++---- libcudacxx/include/cuda/std/array | 4 +- libcudacxx/include/cuda/std/bitset | 4 +- .../tuple/forward_as_tuple_interop.pass.cpp | 2 +- libcudacxx/test/libcudacxx/force_include.h | 6 +- .../equality_comparable.compile.pass.cpp | 8 +- .../totally_ordered.pass.cpp | 4 +- .../concepts.object/copyable.compile.pass.cpp | 8 +- .../concepts.object/movable.compile.pass.cpp | 8 +- .../concepts.object/regular.compile.pass.cpp | 16 +- .../semiregular.compile.pass.cpp | 8 +- .../sequences/array/iterators.pass.cpp | 3 +- .../mdspan/layout_stride/comparison.pass.cpp | 2 +- .../indirectly_readable.compile.pass.cpp | 2 +- .../utilities/const.wrap.class/call.pass.cpp | 2 +- .../func.invoke/invoke.pass.cpp | 4 +- .../optional.object/triviality.pass.cpp | 8 +- .../optional/optional.specalg/swap.pass.cpp | 3 +- .../bitset.members/to_ullong.pass.cpp | 2 +- .../bitset.members/to_ulong.pass.cpp | 2 +- .../nvtarget/arch_specific/arch_specific.cu | 3 +- .../test/nvtarget/dialect/nvtargettest.cu | 3 +- .../family_specific/family_specific.cu | 3 +- .../test/support/test_execution_policies.h | 23 +-- .../benchmarks/compute/histogram/even.py | 2 +- .../compute/host/host_benchmark_cases.py | 2 +- .../benchmarks/compute/merge_sort/keys.py | 2 +- .../benchmarks/compute/merge_sort/pairs.py | 2 +- .../benchmarks/compute/partition/three_way.py | 2 +- .../benchmarks/compute/radix_sort/keys.py | 2 +- .../benchmarks/compute/radix_sort/pairs.py | 2 +- .../benchmarks/compute/reduce/custom.py | 2 +- .../benchmarks/compute/reduce/min.py | 2 +- .../compute/reduce/nondeterministic.py | 2 +- .../benchmarks/compute/reduce/sum.py | 2 +- .../compute/scan/exclusive/custom.py | 2 +- .../benchmarks/compute/scan/exclusive/sum.py | 2 +- .../compute/segmented_reduce/variable_sum.py | 2 +- .../benchmarks/compute/segmented_sort/keys.py | 2 +- .../cuda_cccl/benchmarks/compute/select/if.py | 2 +- .../compute/select/unique_by_key.py | 2 +- .../compute/transform/babelstream.py | 2 +- .../compute/transform/complex_cmp.py | 2 +- .../benchmarks/compute/transform/fib.py | 2 +- .../benchmarks/compute/transform/fill.py | 2 +- .../benchmarks/compute/transform/grayscale.py | 2 +- .../benchmarks/compute/transform/heavy.py | 2 +- .../compute/transform_reduce/sum.py | 2 +- python/cuda_cccl/benchmarks/compute/utils.py | 3 +- .../coop/_experimental/bench_warp_reduce.py | 3 +- .../_experimental/device_side_benchmark.py | 4 +- python/cuda_cccl/cuda/compute/_bindings.py | 3 +- python/cuda_cccl/cuda/compute/_cpp_compile.py | 3 +- python/cuda_cccl/tests/_utils/device_array.py | 3 +- python/cuda_cccl/tests/compute/conftest.py | 1 - .../compute/examples/raw_op/cpp_stateful.py | 2 +- .../compute/examples/raw_op/cpp_stateless.py | 2 +- .../compute/test_multi_cc_serialization.py | 2 +- python/cuda_cccl/tests/compute/test_raw_op.py | 2 +- thrust/testing/cuda/transform_iterator.cu | 4 +- thrust/testing/dereference.cu | 3 +- thrust/testing/unittest/special_types.h | 3 +- thrust/testing/unittest/testframework.h | 10 +- thrust/testing/zip_iterator.cu | 4 +- thrust/thrust/iterator/constant_iterator.h | 4 +- thrust/thrust/iterator/iterator_facade.h | 72 ++++----- thrust/thrust/iterator/permutation_iterator.h | 4 +- .../iterator/tabulate_output_iterator.h | 4 +- thrust/thrust/system/cuda/detail/util.h | 4 +- 165 files changed, 910 insertions(+), 930 deletions(-) diff --git a/benchmarks/scripts/cccl/bench/search.py b/benchmarks/scripts/cccl/bench/search.py index 259cb586630..59bcbd095ce 100644 --- a/benchmarks/scripts/cccl/bench/search.py +++ b/benchmarks/scripts/cccl/bench/search.py @@ -80,12 +80,14 @@ def filter_benchmark_space_for_p0(algname, ct_space, rt_values): ]: ct_space = list( filter( - lambda variant: not ( - ("OffsetT{ct}=I64" in variant) - or ("KeyT{ct}=I16" in variant) - or ("ValueT{ct}=I16" in variant) - or ("KeyT{ct}=I128" in variant) - or ("ValueT{ct}=I128" in variant) + lambda variant: ( + not ( + ("OffsetT{ct}=I64" in variant) + or ("KeyT{ct}=I16" in variant) + or ("ValueT{ct}=I16" in variant) + or ("KeyT{ct}=I128" in variant) + or ("ValueT{ct}=I128" in variant) + ) ), ct_space, ) diff --git a/benchmarks/scripts/sol.py b/benchmarks/scripts/sol.py index 9d94a8bcef7..19eeb0b0b24 100755 --- a/benchmarks/scripts/sol.py +++ b/benchmarks/scripts/sol.py @@ -103,9 +103,7 @@ def alg_bws(dfs, verbose): medians["hue"] = medians["hue"] + gpuname + " " if medians["variant"].unique().size > 1: variant = ( - medians["variant"] - .astype(str) - .map(lambda x: (" " + x if x != "base" else "")) + medians["variant"].astype(str).map(lambda x: " " + x if x != "base" else "") ) medians["hue"] = medians["hue"] + variant + " " if verbose or medians["ctk"].unique().size > 1: diff --git a/c/parallel/test/test_merge_sort.cpp b/c/parallel/test/test_merge_sort.cpp index dd25747dcf4..61552f77ce6 100644 --- a/c/parallel/test/test_merge_sort.cpp +++ b/c/parallel/test/test_merge_sort.cpp @@ -252,7 +252,7 @@ C2H_TEST("DeviceMergeSort:SortPairsCopy works with custom types", "[merge_sort]" { const size_t num_items = GENERATE_COPY(take(2, random(1, 100000)), values({5, 10000, 100000})); operation_t op = make_operation("op", - R"(struct key_pair { short a; size_t b; }; + R"(struct key_pair { short a; size_t b; }; extern "C" __device__ void op(void* lhs_ptr, void* rhs_ptr, bool* out_ptr) { key_pair* lhs = static_cast(lhs_ptr); key_pair* rhs = static_cast(rhs_ptr); @@ -308,7 +308,7 @@ C2H_TEST("DeviceMergeSort:SortPairsCopy works with custom types with well-known { const size_t num_items = GENERATE_COPY(take(2, random(1, 100000)), values({5, 10000, 100000})); operation_t op_state = make_operation("op", - R"(struct key_pair { short a; size_t b; }; + R"(struct key_pair { short a; size_t b; }; extern "C" __device__ void op(void* lhs_ptr, void* rhs_ptr, bool* out_ptr) { key_pair* lhs = static_cast(lhs_ptr); key_pair* rhs = static_cast(rhs_ptr); diff --git a/c/parallel/test/test_reduce.cpp b/c/parallel/test/test_reduce.cpp index a9c6b64447e..b39e8b01fc8 100644 --- a/c/parallel/test/test_reduce.cpp +++ b/c/parallel/test/test_reduce.cpp @@ -229,7 +229,7 @@ C2H_TEST("Reduce works with custom types", "[reduce]") const std::size_t num_items = GENERATE(0, 42, take(4, random(1 << 12, 1 << 24))); operation_t op = make_operation("op", - R"(struct pair { short a; size_t b; }; + R"(struct pair { short a; size_t b; }; extern "C" __device__ void op(void* lhs_ptr, void* rhs_ptr, void* out_ptr) { pair* lhs = static_cast(lhs_ptr); pair* rhs = static_cast(rhs_ptr); @@ -266,7 +266,7 @@ C2H_TEST("Reduce works with custom types with well-known operations", "[reduce][ const std::size_t num_items = GENERATE(0, 42, take(4, random(1 << 12, 1 << 24))); operation_t op_state = make_operation("op", - R"(struct pair { short a; size_t b; }; + R"(struct pair { short a; size_t b; }; extern "C" __device__ void op(void* lhs_ptr, void* rhs_ptr, void* out_ptr) { pair* lhs = static_cast(lhs_ptr); pair* rhs = static_cast(rhs_ptr); diff --git a/c/parallel/test/test_scan.cpp b/c/parallel/test/test_scan.cpp index 95e76e01305..6e2034c9b40 100644 --- a/c/parallel/test/test_scan.cpp +++ b/c/parallel/test/test_scan.cpp @@ -399,7 +399,7 @@ C2H_TEST("Scan works with custom types", "[scan]") const std::size_t num_items = GENERATE(0, 42, take(4, random(1 << 12, 1 << 24))); operation_t op = make_operation("op", - R"(struct pair { short a; size_t b; }; + R"(struct pair { short a; size_t b; }; extern "C" __device__ void op(void* lhs_ptr, void* rhs_ptr, void* out_ptr) { pair* lhs = static_cast(lhs_ptr); pair* rhs = static_cast(rhs_ptr); @@ -439,7 +439,7 @@ C2H_TEST("Scan works with custom types with well-known operations", "[scan][well const std::size_t num_items = GENERATE(0, 42, take(4, random(1 << 12, 1 << 24))); operation_t op_state = make_operation("op", - R"(struct pair { short a; size_t b; }; + R"(struct pair { short a; size_t b; }; extern "C" __device__ void op(void* lhs_ptr, void* rhs_ptr, void* out_ptr) { pair* lhs = static_cast(lhs_ptr); pair* rhs = static_cast(rhs_ptr); diff --git a/c/parallel/test/test_transform.cpp b/c/parallel/test/test_transform.cpp index 6faedbc0425..b9016ddda76 100644 --- a/c/parallel/test/test_transform.cpp +++ b/c/parallel/test/test_transform.cpp @@ -355,7 +355,7 @@ C2H_TEST("Transform works with output of different type", "[transform]") const std::size_t num_items = GENERATE(0, 42, take(4, random(1 << 12, 1 << 24))); operation_t op = make_operation("op", - R"(struct pair { short a; size_t b; }; + R"(struct pair { short a; size_t b; }; extern "C" __device__ void op(void* x_ptr, void* out_ptr) { int* x = static_cast(x_ptr); pair* out = static_cast(out_ptr); @@ -443,7 +443,7 @@ C2H_TEST("Transform works with custom types", "[transform]") const std::size_t num_items = GENERATE(0, 42, take(4, random(1 << 12, 1 << 24))); operation_t op = make_operation("op", - R"(struct pair { short a; size_t b; }; + R"(struct pair { short a; size_t b; }; extern "C" __device__ void op(void* x_ptr, void* out_ptr) { pair* x = static_cast(x_ptr); pair* out = static_cast(out_ptr); diff --git a/c/parallel/test/test_unique_by_key.cpp b/c/parallel/test/test_unique_by_key.cpp index c22bafb5f73..b8aebbb3d4e 100644 --- a/c/parallel/test/test_unique_by_key.cpp +++ b/c/parallel/test/test_unique_by_key.cpp @@ -468,7 +468,7 @@ C2H_TEST("DeviceSelect::UniqueByKey works with custom types", "[device][select_u const int num_items = GENERATE_COPY(take(2, random(1, 1000000))); operation_t op = make_operation("op", - R"(struct key_pair { short a; size_t b; }; + R"(struct key_pair { short a; size_t b; }; extern "C" __device__ void op(void* lhs_ptr, void* rhs_ptr, bool* out_ptr) { key_pair* lhs = static_cast(lhs_ptr); key_pair* rhs = static_cast(rhs_ptr); @@ -541,7 +541,7 @@ C2H_TEST("DeviceSelect::UniqueByKey works with custom types with well-known oper const int num_items = GENERATE_COPY(take(2, random(1, 1000000))); operation_t op_state = make_operation("op", - R"(struct key_pair { short a; size_t b; }; + R"(struct key_pair { short a; size_t b; }; extern "C" __device__ void op(void* lhs_ptr, void* rhs_ptr, bool* out_ptr) { key_pair* lhs = static_cast(lhs_ptr); key_pair* rhs = static_cast(rhs_ptr); @@ -706,7 +706,7 @@ C2H_TEST("DeviceSelect::UniqueByKey fails to build for large types due to no vsm const int num_items = 1; operation_t op = make_operation("op", - R"(struct large_key_pair { int a; char c[500]; }; + R"(struct large_key_pair { int a; char c[500]; }; extern "C" __device__ bool op(large_key_pair lhs, large_key_pair rhs) { return lhs.a == rhs.a; })"); diff --git a/c/parallel/test/test_util.h b/c/parallel/test/test_util.h index 00770f44151..fbe7cefe563 100644 --- a/c/parallel/test/test_util.h +++ b/c/parallel/test/test_util.h @@ -1099,10 +1099,10 @@ inline std::tuple make_counting_iterator_ std::string iterator_state_def_src = std::format("struct {0} {{ {1} value; }};\n", iterator_state_name, value_type); std::string advance_fn_def_src = std::format( "extern \"C\" __device__ void {0}(void* state, const void* offset) {{\n" - " auto* typed_state = static_cast<{1}*>(state);\n" - " auto offset_val = *static_cast(offset);\n" - " typed_state->value += offset_val;\n" - "}}", + " auto* typed_state = static_cast<{1}*>(state);\n" + " auto offset_val = *static_cast(offset);\n" + " typed_state->value += offset_val;\n" + "}}", advance_fn_name, iterator_state_name); @@ -1186,10 +1186,10 @@ inline std::tuple make_reverse_iterator_s std::string iterator_state_src = std::format("struct {0} {{ {1}* data; }};\n", iterator_state_name, value_type); std::string advance_fn_src = std::format( "extern \"C\" __device__ void {0}(void* state, const void* offset) {{\n" - " auto* typed_state = static_cast<{1}*>(state);\n" - " auto offset_val = *static_cast(offset);\n" - " typed_state->data -= offset_val;\n" - "}}", + " auto* typed_state = static_cast<{1}*>(state);\n" + " auto offset_val = *static_cast(offset);\n" + " typed_state->data -= offset_val;\n" + "}}", advance_fn_name, iterator_state_name); std::string dereference_fn_src; diff --git a/cub/cub/agent/agent_batched_topk.cuh b/cub/cub/agent/agent_batched_topk.cuh index d2a99cfc809..4f09f0abaa5 100644 --- a/cub/cub/agent/agent_batched_topk.cuh +++ b/cub/cub/agent/agent_batched_topk.cuh @@ -216,7 +216,7 @@ struct agent_batched_topk_worker_per_segment else { // Process small segment - const auto k = (::cuda::std::min) (params::get_param(k_param, segment_id), + const auto k = (::cuda::std::min) (params::get_param(k_param, segment_id), static_cast(segment_size)); const auto direction = select_directions.get_param(segment_id); diff --git a/cub/cub/agent/agent_select_if.cuh b/cub/cub/agent/agent_select_if.cuh index 44558013953..c81ccf4fe50 100644 --- a/cub/cub/agent/agent_select_if.cuh +++ b/cub/cub/agent/agent_select_if.cuh @@ -893,7 +893,7 @@ struct AgentSelectIf 0, 0, num_tile_selections, - bool_constant_v < SelectionOpt == SelectImpl::Partition >); + bool_constant_v); return num_tile_selections; } @@ -976,7 +976,7 @@ struct AgentSelectIf num_selections_prefix, num_rejected_prefix, num_selections, - bool_constant_v < SelectionOpt == SelectImpl::Partition >); + bool_constant_v); return num_selections; } diff --git a/cub/cub/agent/agent_three_way_partition.cuh b/cub/cub/agent/agent_three_way_partition.cuh index 505edc55740..0e11783b308 100644 --- a/cub/cub/agent/agent_three_way_partition.cuh +++ b/cub/cub/agent/agent_three_way_partition.cuh @@ -335,7 +335,7 @@ struct AgentThreeWayPartition // Medium item int local_selection_idx = (first_items_selection_indices - num_first_selections_prefix) + (second_items_selection_indices - num_second_selections_prefix); - local_scatter_offset = second_item_end + item_idx - local_selection_idx; + local_scatter_offset = second_item_end + item_idx - local_selection_idx; } temp_storage.raw_exchange.Alias()[local_scatter_offset] = items[ITEM]; diff --git a/cub/cub/block/block_radix_sort.cuh b/cub/cub/block/block_radix_sort.cuh index 86f17c5e655..4c7ae09c27f 100644 --- a/cub/cub/block/block_radix_sort.cuh +++ b/cub/cub/block/block_radix_sort.cuh @@ -741,8 +741,7 @@ public: template _CCCL_DEVICE _CCCL_FORCEINLINE // ::cuda::std::enable_if_t< // - !::cuda::std::is_convertible_v> - Sort(KeyT (&keys)[ItemsPerThread], DecomposerT decomposer) + !::cuda::std::is_convertible_v> Sort(KeyT (&keys)[ItemsPerThread], DecomposerT decomposer) { Sort(keys, decomposer, 0, detail::radix::traits_t::default_end_bit(decomposer)); } diff --git a/cub/cub/block/radix_rank_sort_operations.cuh b/cub/cub/block/radix_rank_sort_operations.cuh index 6a47d2091ad..13b0cf5c800 100644 --- a/cub/cub/block/radix_rank_sort_operations.cuh +++ b/cub/cub/block/radix_rank_sort_operations.cuh @@ -544,8 +544,7 @@ private: public: template static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE // - bit_ordered_type - In(bit_ordered_type key, DecomposerT decomposer = {}) + bit_ordered_type In(bit_ordered_type key, DecomposerT decomposer = {}) { key = bit_ordered_conversion_policy::to_bit_ordered(decomposer, key); if constexpr (IS_DESCENDING) @@ -557,8 +556,7 @@ public: template static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE // - bit_ordered_type - Out(bit_ordered_type key, DecomposerT decomposer = {}) + bit_ordered_type Out(bit_ordered_type key, DecomposerT decomposer = {}) { if constexpr (IS_DESCENDING) { @@ -570,8 +568,7 @@ public: template static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE // - bit_ordered_type - DefaultKey(DecomposerT decomposer = {}) + bit_ordered_type DefaultKey(DecomposerT decomposer = {}) { return IS_DESCENDING ? traits::min_raw_binary_key(decomposer) : traits::max_raw_binary_key(decomposer); } diff --git a/cub/cub/detail/fast_modulo_division.cuh b/cub/cub/detail/fast_modulo_division.cuh index 8b26b0bb9da..6ae78f916b0 100644 --- a/cub/cub/detail/fast_modulo_division.cuh +++ b/cub/cub/detail/fast_modulo_division.cuh @@ -160,7 +160,7 @@ public: _CCCL_ASSERT(static_cast(num_bits + BitSize - BitOffset) < sizeof(larger_t) * CHAR_BIT, "overflow error"); // without explicit power-of-two check, num_bits needs to replace +1 with !::cuda::is_power_of_two(udivisor) _multiplier = static_cast(::cuda::ceil_div(larger_t{1} << (num_bits + BitSize - BitOffset), // - static_cast(divisor))); + static_cast(divisor))); _shift_right = num_bits - BitOffset; _CCCL_ASSERT(_multiplier != 0, "overflow error"); } diff --git a/cub/cub/device/device_scan.cuh b/cub/cub/device/device_scan.cuh index 106db7e6d02..d29ddf36cb0 100644 --- a/cub/cub/device/device_scan.cuh +++ b/cub/cub/device/device_scan.cuh @@ -249,10 +249,10 @@ struct DeviceScan { using offset_t = detail::choose_offset_t; using accum_t = ::cuda::std::__accumulator_t< - ScanOpT, - cub::detail::it_value_t, - ::cuda::std:: - _If<::cuda::std::is_same_v, cub::detail::it_value_t, InitValueT>>; + ScanOpT, + cub::detail::it_value_t, + ::cuda::std:: + _If<::cuda::std::is_same_v, cub::detail::it_value_t, InitValueT>>; using default_policy_selector_t = detail::scan_by_key::policy_selector_from_types, diff --git a/cub/cub/device/dispatch/dispatch_adjacent_difference.cuh b/cub/cub/device/dispatch/dispatch_adjacent_difference.cuh index 6e8b1226327..350c560dcf0 100644 --- a/cub/cub/device/dispatch/dispatch_adjacent_difference.cuh +++ b/cub/cub/device/dispatch/dispatch_adjacent_difference.cuh @@ -249,14 +249,15 @@ struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceAdjacentDifference" error = CubDebug( THRUST_NS_QUALIFIER::cuda_cub::detail::triple_chevron( num_tiles, AdjacentDifferencePolicyT::BLOCK_THREADS, 0, stream) - .doit(detail::adjacent_difference::DeviceAdjacentDifferenceDifferenceKernel < KernelPolicySelector, - InputIteratorT, - OutputIteratorT, - DifferenceOpT, - OffsetT, - InputT, - AliasOpt == MayAlias::Yes, - ReadOpt == ReadOption::Left >, + .doit(detail::adjacent_difference::DeviceAdjacentDifferenceDifferenceKernel< + KernelPolicySelector, + InputIteratorT, + OutputIteratorT, + DifferenceOpT, + OffsetT, + InputT, + AliasOpt == MayAlias::Yes, + ReadOpt == ReadOption::Left>, d_input, first_tile_previous, d_output, @@ -441,14 +442,15 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE auto dispatch( if (const auto error = CubDebug( THRUST_NS_QUALIFIER::cuda_cub::detail::triple_chevron(num_tiles, active_policy.threads_per_block, 0, stream) - .doit(DeviceAdjacentDifferenceDifferenceKernel < policy_selector_t, - InputIteratorT, - OutputIteratorT, - DifferenceOpT, - offset_t, - input_t, - AliasOpt == MayAlias::Yes, - ReadOpt == ReadOption::Left >, + .doit(DeviceAdjacentDifferenceDifferenceKernel< + policy_selector_t, + InputIteratorT, + OutputIteratorT, + DifferenceOpT, + offset_t, + input_t, + AliasOpt == MayAlias::Yes, + ReadOpt == ReadOption::Left>, d_input, first_tile_previous, d_output, diff --git a/cub/cub/device/dispatch/dispatch_batch_memcpy.cuh b/cub/cub/device/dispatch/dispatch_batch_memcpy.cuh index f44333a958f..ff9d8b77ac9 100644 --- a/cub/cub/device/dispatch/dispatch_batch_memcpy.cuh +++ b/cub/cub/device/dispatch/dispatch_batch_memcpy.cuh @@ -152,8 +152,8 @@ __launch_bounds__(int(current_policy().large_buffer.threads_per_ { if (thread_offset < buffer_sizes[buffer_id]) { - const auto value = read_item < MemcpyOpt == CopyAlg::Memcpy, AliasT, - InputBufferT > (input_buffer_it[buffer_id], thread_offset); + const auto value = + read_item(input_buffer_it[buffer_id], thread_offset); write_item( output_buffer_it[buffer_id], thread_offset, value); } @@ -346,7 +346,7 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t dispatch( constexpr BlockOffsetT init_kernel_threads = 128U; const auto tile_size = static_cast(active_policy.small_buffer.threads_per_block) - * static_cast(active_policy.small_buffer.buffers_per_thread); + * static_cast(active_policy.small_buffer.buffers_per_thread); constexpr auto max_num_buffers_per_invocation = ::cuda::std::int64_t{512 * 1024 * 1024}; static_assert(max_num_buffers_per_invocation <= ::cuda::std::numeric_limits::max()); diff --git a/cub/cub/device/dispatch/dispatch_histogram.cuh b/cub/cub/device/dispatch/dispatch_histogram.cuh index 704c57add4f..673be75f6fd 100644 --- a/cub/cub/device/dispatch/dispatch_histogram.cuh +++ b/cub/cub/device/dispatch/dispatch_histogram.cuh @@ -216,12 +216,12 @@ CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE auto dispatch( if constexpr (IsDeviceInit) { return kernel_source.template HistogramSweepKernelDeviceInit< - PolicySelector, - PRIVATIZED_SMEM_BINS, - FirstLevelArrayT, - SecondLevelArrayT, - IsEven, - IsByteSample>(); + PolicySelector, + PRIVATIZED_SMEM_BINS, + FirstLevelArrayT, + SecondLevelArrayT, + IsEven, + IsByteSample>(); } else { diff --git a/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh b/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh index a67850b9fd7..82eb15754bb 100644 --- a/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh +++ b/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh @@ -205,12 +205,12 @@ __launch_bounds__(int(current_policy().threads_per_block)) { static constexpr ReduceByKeyPolicy policy = current_policy(); using AgentReduceByKeyPolicyT = agent_reduce_by_key_policy< - policy.threads_per_block, - policy.items_per_thread, - policy.load_algorithm, - policy.load_modifier, - policy.scan_algorithm, - delay_constructor_t>; + policy.threads_per_block, + policy.items_per_thread, + policy.load_algorithm, + policy.load_modifier, + policy.scan_algorithm, + delay_constructor_t>; using vsmem_helper_t = vsmem_helper_default_fallback_policy_t< AgentReduceByKeyPolicyT, @@ -651,12 +651,12 @@ _CCCL_HOST_DEVICE_API auto determine_threads_items_vsmem(PolicyGetter policy_get // TODO(bgruber): refactor this in the future constexpr ReduceByKeyPolicy policy = policy_getter(); using Policy = agent_reduce_by_key_policy< - policy.threads_per_block, - policy.items_per_thread, - policy.load_algorithm, - policy.load_modifier, - policy.scan_algorithm, - delay_constructor_t>; + policy.threads_per_block, + policy.items_per_thread, + policy.load_algorithm, + policy.load_modifier, + policy.scan_algorithm, + delay_constructor_t>; using vsmem_helper_t = vsmem_helper_default_fallback_policy_t; return ::cuda::std::tuple{vsmem_helper_t::agent_policy_t::BLOCK_THREADS, vsmem_helper_t::agent_policy_t::ITEMS_PER_THREAD, diff --git a/cub/cub/device/dispatch/dispatch_rle.cuh b/cub/cub/device/dispatch/dispatch_rle.cuh index de21b507437..73bb94efc39 100644 --- a/cub/cub/device/dispatch/dispatch_rle.cuh +++ b/cub/cub/device/dispatch/dispatch_rle.cuh @@ -187,13 +187,13 @@ __launch_bounds__(int(current_policy().threads_per_block)) { static constexpr RleNonTrivialRunsPolicy policy = current_policy(); using AgentRlePolicyT = agent_rle_policy< - policy.threads_per_block, - policy.items_per_thread, - policy.load_algorithm, - policy.load_modifier, - policy.store_with_time_slicing, - policy.scan_algorithm, - delay_constructor_t>; + policy.threads_per_block, + policy.items_per_thread, + policy.load_algorithm, + policy.load_modifier, + policy.store_with_time_slicing, + policy.scan_algorithm, + delay_constructor_t>; using AgentRleT = AgentRle, - ::cuda::std::_If<::cuda::std::is_same_v, - cub::detail::it_value_t, - typename InitValueT::value_type>>, + typename AccumT = ::cuda::std::__accumulator_t, + ::cuda::std::_If<::cuda::std::is_same_v, + cub::detail::it_value_t, + typename InitValueT::value_type>>, ForceInclusive EnforceInclusive = ForceInclusive::No, typename PolicyHub = detail::scan:: policy_hub, detail::it_value_t, AccumT, OffsetT, ScanOpT>, diff --git a/cub/cub/device/dispatch/dispatch_scan_by_key.cuh b/cub/cub/device/dispatch/dispatch_scan_by_key.cuh index 177ba984f07..7187ad5ffba 100644 --- a/cub/cub/device/dispatch/dispatch_scan_by_key.cuh +++ b/cub/cub/device/dispatch/dispatch_scan_by_key.cuh @@ -236,15 +236,15 @@ template < typename PolicyHub = policy_hub, ScanOpT>, typename PolicySelector = policy_selector_from_hub, typename KernelSource = DeviceScanByKeyKernelSource< - PolicySelector, - KeysInputIteratorT, - ValuesInputIteratorT, - ValuesOutputIteratorT, - EqualityOp, - ScanOpT, - InitValueT, - OffsetT, - AccumT>, + PolicySelector, + KeysInputIteratorT, + ValuesInputIteratorT, + ValuesOutputIteratorT, + EqualityOp, + ScanOpT, + InitValueT, + OffsetT, + AccumT>, typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> struct dispatch_scan_by_key { @@ -670,15 +670,15 @@ template < cub::detail::it_value_t, ScanOpT>, typename KernelSource = DeviceScanByKeyKernelSource< - PolicySelector, - KeysInputIteratorT, - ValuesInputIteratorT, - ValuesOutputIteratorT, - EqualityOp, - ScanOpT, - InitValueT, - OffsetT, - AccumT>, + PolicySelector, + KeysInputIteratorT, + ValuesInputIteratorT, + ValuesOutputIteratorT, + EqualityOp, + ScanOpT, + InitValueT, + OffsetT, + AccumT>, typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> #if _CCCL_HAS_CONCEPTS() requires scan_by_key_policy_selector @@ -873,15 +873,15 @@ template < detail::scan_by_key::policy_hub, ScanOpT>, typename PolicySelector = detail::scan_by_key::policy_selector_from_hub, typename KernelSource = detail::scan_by_key::DeviceScanByKeyKernelSource< - PolicySelector, - KeysInputIteratorT, - ValuesInputIteratorT, - ValuesOutputIteratorT, - EqualityOp, - ScanOpT, - InitValueT, - OffsetT, - AccumT>, + PolicySelector, + KeysInputIteratorT, + ValuesInputIteratorT, + ValuesOutputIteratorT, + EqualityOp, + ScanOpT, + InitValueT, + OffsetT, + AccumT>, typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> using DispatchScanByKey CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceScan") = detail::scan_by_key::dispatch_scan_by_key< diff --git a/cub/cub/device/dispatch/dispatch_segmented_radix_sort.cuh b/cub/cub/device/dispatch/dispatch_segmented_radix_sort.cuh index 2c9d6fbe652..ce33983c7c0 100644 --- a/cub/cub/device/dispatch/dispatch_segmented_radix_sort.cuh +++ b/cub/cub/device/dispatch/dispatch_segmented_radix_sort.cuh @@ -911,7 +911,7 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t dispatch( { using default_policy_selector_t = policy_selector_from_types; using policy_selector_t = ::cuda::std::decay_t< - ::cuda::std::execution::__query_result_or_t>; + ::cuda::std::execution::__query_result_or_t>; #if _CCCL_HAS_CONCEPTS() static_assert(segmented_radix_sort_policy_selector); #endif // _CCCL_HAS_CONCEPTS() diff --git a/cub/cub/device/dispatch/dispatch_segmented_reduce.cuh b/cub/cub/device/dispatch/dispatch_segmented_reduce.cuh index 35c432b40f0..a01cfe455e0 100644 --- a/cub/cub/device/dispatch/dispatch_segmented_reduce.cuh +++ b/cub/cub/device/dispatch/dispatch_segmented_reduce.cuh @@ -503,15 +503,15 @@ template < decltype(select_segmented_accum_t(static_cast(nullptr))), typename PolicySelector = policy_selector_from_types, typename KernelSource = DeviceSegmentedReduceKernelSource< - PolicySelector, - InputIteratorT, - OutputIteratorT, - BeginOffsetIteratorT, - EndOffsetIteratorT, - OffsetT, - ReductionOpT, - InitValueT, - AccumT>, + PolicySelector, + InputIteratorT, + OutputIteratorT, + BeginOffsetIteratorT, + EndOffsetIteratorT, + OffsetT, + ReductionOpT, + InitValueT, + AccumT>, typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> #if _CCCL_HAS_CONCEPTS() requires segmented_reduce_policy_selector @@ -711,13 +711,13 @@ template (nullptr))), typename PolicySelector = policy_selector_from_types, typename KernelSource = DeviceFixedSizeSegmentedReduceKernelSource< - PolicySelector, - InputIteratorT, - OutputIteratorT, - OffsetT, - ReductionOpT, - InitValueT, - AccumT>, + PolicySelector, + InputIteratorT, + OutputIteratorT, + OffsetT, + ReductionOpT, + InitValueT, + AccumT>, typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY, ::cuda::std::enable_if_t<::cuda::std::is_arithmetic_v, int> = 0> #if _CCCL_HAS_CONCEPTS() diff --git a/cub/cub/device/dispatch/dispatch_segmented_scan.cuh b/cub/cub/device/dispatch/dispatch_segmented_scan.cuh index c8dd92d2ed0..673c9b7fb02 100644 --- a/cub/cub/device/dispatch/dispatch_segmented_scan.cuh +++ b/cub/cub/device/dispatch/dispatch_segmented_scan.cuh @@ -96,17 +96,17 @@ template < common_iterator_value_t, typename PolicySelector = policy_selector_from_types, typename KernelSource = device_segmented_scan_kernel_source< - PolicySelector, - InputIteratorT, - OutputIteratorT, - BeginOffsetIteratorInputT, - EndOffsetIteratorInputT, - BeginOffsetIteratorOutputT, - OffsetT, - ScanOpT, - InitValueT, - AccumT, - EnforceInclusive>, + PolicySelector, + InputIteratorT, + OutputIteratorT, + BeginOffsetIteratorInputT, + EndOffsetIteratorInputT, + BeginOffsetIteratorOutputT, + OffsetT, + ScanOpT, + InitValueT, + AccumT, + EnforceInclusive>, typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> #if _CCCL_HAS_CONCEPTS() requires segmented_scan_policy_selector diff --git a/cub/cub/device/dispatch/dispatch_select_if.cuh b/cub/cub/device/dispatch/dispatch_select_if.cuh index 82c0af411ef..69c8cf33032 100644 --- a/cub/cub/device/dispatch/dispatch_select_if.cuh +++ b/cub/cub/device/dispatch/dispatch_select_if.cuh @@ -216,14 +216,14 @@ struct make_vsmem_helper { static constexpr SelectPolicy active_policy = DefaultPolicyGetter{}(); using agent_policy_t = detail::agent_select_if_policy< - active_policy.threads_per_block, - active_policy.items_per_thread, - active_policy.load_algorithm, - active_policy.load_modifier, - active_policy.scan_algorithm, - delay_constructor_t>; + active_policy.threads_per_block, + active_policy.items_per_thread, + active_policy.load_algorithm, + active_policy.load_modifier, + active_policy.scan_algorithm, + delay_constructor_t>; using type = vsmem_helper_default_fallback_policy_t< agent_policy_t, bind_selection_opt::template agent_t, diff --git a/cub/cub/device/dispatch/dispatch_three_way_partition.cuh b/cub/cub/device/dispatch/dispatch_three_way_partition.cuh index b921a4bade6..7439bf9000d 100644 --- a/cub/cub/device/dispatch/dispatch_three_way_partition.cuh +++ b/cub/cub/device/dispatch/dispatch_three_way_partition.cuh @@ -443,18 +443,18 @@ template , per_partition_offset_t>, typename KernelSource = DeviceThreeWayPartitionKernelSource< - PolicySelector, - InputIteratorT, - FirstOutputIteratorT, - SecondOutputIteratorT, - UnselectedOutputIteratorT, - NumSelectedIteratorT, - ScanTileStateT, - SelectFirstPartOp, - SelectSecondPartOp, - per_partition_offset_t, - streaming_context_t, - OffsetT>, + PolicySelector, + InputIteratorT, + FirstOutputIteratorT, + SecondOutputIteratorT, + UnselectedOutputIteratorT, + NumSelectedIteratorT, + ScanTileStateT, + SelectFirstPartOp, + SelectSecondPartOp, + per_partition_offset_t, + streaming_context_t, + OffsetT>, typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> #if _CCCL_HAS_CONCEPTS() requires three_way_partition_policy_selector diff --git a/cub/cub/device/dispatch/kernels/kernel_histogram.cuh b/cub/cub/device/dispatch/kernels/kernel_histogram.cuh index 8d8b468e1bc..e05fd9b3f3b 100644 --- a/cub/cub/device/dispatch/kernels/kernel_histogram.cuh +++ b/cub/cub/device/dispatch/kernels/kernel_histogram.cuh @@ -173,8 +173,8 @@ struct Transforms #endif // _CCCL_HAS_NVFP16() #if _CCCL_HAS_NVBF16() - _CCCL_HOST_DEVICE _CCCL_FORCEINLINE ScaleT - ComputeScale(int num_levels, __nv_bfloat16 max_level, __nv_bfloat16 min_level) + _CCCL_HOST_DEVICE + _CCCL_FORCEINLINE ScaleT ComputeScale(int num_levels, __nv_bfloat16 max_level, __nv_bfloat16 min_level) { ScaleT result; NV_IF_ELSE_TARGET( @@ -204,8 +204,8 @@ struct Transforms #endif // _CCCL_HAS_NVFP16() #if _CCCL_HAS_NVBF16() - _CCCL_HOST_DEVICE _CCCL_FORCEINLINE int - SampleIsValid(__nv_bfloat16 sample, __nv_bfloat16 max_level, __nv_bfloat16 min_level) + _CCCL_HOST_DEVICE + _CCCL_FORCEINLINE int SampleIsValid(__nv_bfloat16 sample, __nv_bfloat16 max_level, __nv_bfloat16 min_level) { NV_IF_ELSE_TARGET(NV_PROVIDES_SM_80, (return __hge(sample, min_level) && __hlt(sample, max_level);), diff --git a/cub/cub/device/dispatch/kernels/kernel_radix_sort.cuh b/cub/cub/device/dispatch/kernels/kernel_radix_sort.cuh index ce7bdfa9652..774a45a2c14 100644 --- a/cub/cub/device/dispatch/kernels/kernel_radix_sort.cuh +++ b/cub/cub/device/dispatch/kernels/kernel_radix_sort.cuh @@ -151,15 +151,15 @@ __launch_bounds__(current_policy().scan.lookback.threads_per_blo static_assert(active_policy.algorithm == ScanAlgorithm::lookback); static constexpr ScanLookbackPolicy policy = active_policy.lookback; using ScanPolicy = agent_scan_policy< - 0, - 0, - void, - policy.load_algorithm, - policy.load_modifier, - policy.store_algorithm, - policy.scan_algorithm, - NoScaling, - delay_constructor_t>; + 0, + 0, + void, + policy.load_algorithm, + policy.load_modifier, + policy.store_algorithm, + policy.scan_algorithm, + NoScaling, + delay_constructor_t>; // Parameterize the AgentScan type for the current configuration using AgentScanT = scan::AgentScan, OffsetT, OffsetT, OffsetT>; @@ -414,7 +414,7 @@ __launch_bounds__(current_policy().single_tile.threads_per_block values, current_bit, end_bit, - bool_constant_v < Order == SortOrder::Descending >, + bool_constant_v, bool_constant_v, decomposer); @@ -524,15 +524,15 @@ _CCCL_KERNEL_ATTRIBUTES void __launch_bounds__(current_policy(). { static constexpr RadixSortOnesweepPolicy policy = current_policy().onesweep; using OnesweepPolicyT = detail::agent_radix_sort_onesweep_policy< - 0, - 0, - void, - policy.rank_private_partitions, - policy.rank_algorithm, - policy.scan_algorithm, - policy.store_algorithm, - policy.radix_bits, - NoScaling>; + 0, + 0, + void, + policy.rank_private_partitions, + policy.rank_algorithm, + policy.scan_algorithm, + policy.store_algorithm, + policy.radix_bits, + NoScaling>; using AgentT = AgentRadixSortOnesweep().large_segment.threads_per_blo static constexpr SegmentedSortPolicy active_policy = current_policy(); static constexpr auto large_policy = active_policy.large_segment; using LargeSegmentPolicyT = detail::agent_radix_sort_downsweep_policy< - 0, - 0, - void, - large_policy.load_algorithm, - large_policy.load_modifier, - large_policy.rank_algorithm, - large_policy.scan_algorithm, - large_policy.radix_bits, - NoScaling>; + 0, + 0, + void, + large_policy.load_algorithm, + large_policy.load_modifier, + large_policy.rank_algorithm, + large_policy.scan_algorithm, + large_policy.radix_bits, + NoScaling>; static constexpr auto medium_policy = active_policy.medium_segment; using MediumPolicyT = agent_sub_warp_merge_sort_policy< - medium_policy.threads_per_block, - medium_policy.threads_per_warp, - medium_policy.items_per_thread, - medium_policy.load_algorithm, - medium_policy.load_modifier, - medium_policy.store_algorithm>; + medium_policy.threads_per_block, + medium_policy.threads_per_warp, + medium_policy.items_per_thread, + medium_policy.load_algorithm, + medium_policy.load_modifier, + medium_policy.store_algorithm>; const auto segment_id = static_cast(blockIdx.x); OffsetT segment_begin = d_begin_offsets[segment_id]; @@ -351,20 +351,20 @@ __launch_bounds__(current_policy().small_segment.threads_per_blo static constexpr SegmentedSortPolicy active_policy = current_policy(); static constexpr auto small_policy = active_policy.small_segment; using SmallPolicyT = agent_sub_warp_merge_sort_policy< - small_policy.threads_per_block, - small_policy.threads_per_warp, - small_policy.items_per_thread, - small_policy.load_algorithm, - small_policy.load_modifier, - small_policy.store_algorithm>; + small_policy.threads_per_block, + small_policy.threads_per_warp, + small_policy.items_per_thread, + small_policy.load_algorithm, + small_policy.load_modifier, + small_policy.store_algorithm>; static constexpr auto medium_policy = active_policy.medium_segment; using MediumPolicyT = agent_sub_warp_merge_sort_policy< - medium_policy.threads_per_block, - medium_policy.threads_per_warp, - medium_policy.items_per_thread, - medium_policy.load_algorithm, - medium_policy.load_modifier, - medium_policy.store_algorithm>; + medium_policy.threads_per_block, + medium_policy.threads_per_warp, + medium_policy.items_per_thread, + medium_policy.load_algorithm, + medium_policy.load_modifier, + medium_policy.store_algorithm>; constexpr auto threads_per_medium_segment = static_cast(MediumPolicyT::WARP_THREADS); constexpr auto threads_per_small_segment = static_cast(SmallPolicyT::WARP_THREADS); @@ -481,15 +481,15 @@ __launch_bounds__(current_policy().large_segment.threads_per_blo { static constexpr SegmentedSortRadixSortPolicy large_policy = current_policy().large_segment; using LargeSegmentPolicyT = detail::agent_radix_sort_downsweep_policy< - 0, - 0, - void, - large_policy.load_algorithm, - large_policy.load_modifier, - large_policy.rank_algorithm, - large_policy.scan_algorithm, - large_policy.radix_bits, - NoScaling>; + 0, + 0, + void, + large_policy.load_algorithm, + large_policy.load_modifier, + large_policy.rank_algorithm, + large_policy.scan_algorithm, + large_policy.radix_bits, + NoScaling>; using local_segment_index_t = local_segment_index_t; diff --git a/cub/cub/device/dispatch/tuning/tuning_merge.cuh b/cub/cub/device/dispatch/tuning/tuning_merge.cuh index e44a9bdbe40..67d414c6815 100644 --- a/cub/cub/device/dispatch/tuning/tuning_merge.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_merge.cuh @@ -90,10 +90,10 @@ struct policy_selector [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr auto operator()(::cuda::compute_capability cc) const -> MergePolicy { - const int tune_type_size = key_size + value_size; - const int ipt_800_plus = nominal_4B_items_to_items(15, tune_type_size); - const bool can_bulk_keys = (key_size == key_align) && key_is_trivially_relocatable && key_iterators_are_contiguous - && key_iterator_value_types_are_the_same; + const int tune_type_size = key_size + value_size; + const int ipt_800_plus = nominal_4B_items_to_items(15, tune_type_size); + const bool can_bulk_keys = (key_size == key_align) && key_is_trivially_relocatable && key_iterators_are_contiguous + && key_iterator_value_types_are_the_same; const bool can_bulk_values = (value_size == value_align) && value_is_trivially_relocatable && value_iterators_are_contiguous && value_iterator_value_types_are_the_same; diff --git a/cub/cub/device/dispatch/tuning/tuning_segmented_sort.cuh b/cub/cub/device/dispatch/tuning/tuning_segmented_sort.cuh index 6f730b7b2c5..5b9773d4f4a 100644 --- a/cub/cub/device/dispatch/tuning/tuning_segmented_sort.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_segmented_sort.cuh @@ -575,13 +575,13 @@ struct policy_hub static constexpr int BLOCK_THREADS = 256; static constexpr int PARTITIONING_THRESHOLD = 500; using LargeSegmentPolicy = detail::agent_radix_sort_downsweep_policy< - BLOCK_THREADS, - 23, - DominantT, - BLOCK_LOAD_TRANSPOSE, - LOAD_DEFAULT, - RADIX_RANK_MEMOIZE, - BLOCK_SCAN_WARP_SCANS, + BLOCK_THREADS, + 23, + DominantT, + BLOCK_LOAD_TRANSPOSE, + LOAD_DEFAULT, + RADIX_RANK_MEMOIZE, + BLOCK_SCAN_WARP_SCANS, (sizeof(KeyT) > 1) ? 6 : 4>; static constexpr int ITEMS_PER_SMALL_THREAD = Nominal4BItemsToItems(9); @@ -606,13 +606,13 @@ struct policy_hub static constexpr int BLOCK_THREADS = 256; static constexpr int PARTITIONING_THRESHOLD = 500; using LargeSegmentPolicy = detail::agent_radix_sort_downsweep_policy< - BLOCK_THREADS, - 23, - DominantT, - BLOCK_LOAD_TRANSPOSE, - LOAD_DEFAULT, - RADIX_RANK_MEMOIZE, - BLOCK_SCAN_WARP_SCANS, + BLOCK_THREADS, + 23, + DominantT, + BLOCK_LOAD_TRANSPOSE, + LOAD_DEFAULT, + RADIX_RANK_MEMOIZE, + BLOCK_SCAN_WARP_SCANS, (sizeof(KeyT) > 1) ? 6 : 4>; static constexpr bool LARGE_ITEMS = sizeof(DominantT) > 4; diff --git a/cub/cub/util_device.cuh b/cub/cub/util_device.cuh index d25b0ee09dc..18e2da748e5 100644 --- a/cub/cub/util_device.cuh +++ b/cub/cub/util_device.cuh @@ -819,8 +819,8 @@ private: #if !_CCCL_COMPILER(NVRTC) template - CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static constexpr cudaError_t - runtime_cc_to_compiletime(int device_ptx_version, FunctorT& op) + CUB_RUNTIME_FUNCTION + _CCCL_FORCEINLINE static constexpr cudaError_t runtime_cc_to_compiletime(int device_ptx_version, FunctorT& op) { // We instantiate find_and_invoke_policy for each CudaCcs (the arches we are compiling for), but only call the // one matching device_ptx_version. diff --git a/cub/cub/util_ptx.cuh b/cub/cub/util_ptx.cuh index dfee1213175..dda8d76c848 100644 --- a/cub/cub/util_ptx.cuh +++ b/cub/cub/util_ptx.cuh @@ -79,8 +79,8 @@ BFE(UnsignedBits source, unsigned int bit_start, unsigned int num_bits, detail:: */ template //! deprecated [Since 3.0] -CCCL_DEPRECATED_BECAUSE("Use cuda::bitfield_extract()") _CCCL_DEVICE _CCCL_FORCEINLINE unsigned int -BFE(UnsignedBits source, unsigned int bit_start, unsigned int num_bits) +CCCL_DEPRECATED_BECAUSE("Use cuda::bitfield_extract()") _CCCL_DEVICE +_CCCL_FORCEINLINE unsigned int BFE(UnsignedBits source, unsigned int bit_start, unsigned int num_bits) { return BFE(source, bit_start, num_bits, detail::constant_v); } diff --git a/cub/cub/warp/specializations/warp_redux.cuh b/cub/cub/warp/specializations/warp_redux.cuh index d527cc250ca..1db8fb69ff5 100644 --- a/cub/cub/warp/specializations/warp_redux.cuh +++ b/cub/cub/warp/specializations/warp_redux.cuh @@ -121,8 +121,8 @@ _CUB_REDUX_FLOAT_OP(max) # undef _CUB_REDUX_FLOAT_OP template -[[nodiscard]] _CCCL_DEVICE_API _CCCL_FORCEINLINE T -warp_redux_sm100af(const T input, const ::cuda::std::uint32_t mask, ReductionOp) +[[nodiscard]] _CCCL_DEVICE_API +_CCCL_FORCEINLINE T warp_redux_sm100af(const T input, const ::cuda::std::uint32_t mask, ReductionOp) { static_assert(is_warp_redux_op_supported_sm100af, "Reduction operator not supported"); _CCCL_ASSERT(mask != 0, "Mask must not be 0"); diff --git a/cub/examples/device/example_device_radix_sort_custom.cu b/cub/examples/device/example_device_radix_sort_custom.cu index a59faa1644f..d149df1b0d9 100644 --- a/cub/examples/device/example_device_radix_sort_custom.cu +++ b/cub/examples/device/example_device_radix_sort_custom.cu @@ -27,8 +27,7 @@ struct custom_t struct decomposer_t { __host__ __device__ // - ::cuda::std::tuple - operator()(custom_t& key) const + ::cuda::std::tuple operator()(custom_t& key) const { return {key.i, key.f}; } diff --git a/cub/test/catch2_test_block_reduce.cu b/cub/test/catch2_test_block_reduce.cu index 23e6035f8b7..c9c0654eca8 100644 --- a/cub/test/catch2_test_block_reduce.cu +++ b/cub/test/catch2_test_block_reduce.cu @@ -15,7 +15,7 @@ template -__launch_bounds__(BlockDimX* BlockDimY* BlockDimZ) __global__ +__launch_bounds__(BlockDimX * BlockDimY * BlockDimZ) __global__ void block_reduce_kernel(T* in, T* out, int valid_items, ActionT action) { using block_reduce_t = cub::BlockReduce; diff --git a/cub/test/catch2_test_device_partition_env.cu b/cub/test/catch2_test_device_partition_env.cu index edd91cdea23..7f808eacfcb 100644 --- a/cub/test/catch2_test_device_partition_env.cu +++ b/cub/test/catch2_test_device_partition_env.cu @@ -413,7 +413,7 @@ C2H_TEST("Test ThreeWayPartitionPolicy properties", "[partition][device]") .load_modifier = cub::CacheLoadModifier::LOAD_DEFAULT, .scan_algorithm = cub::BlockScanAlgorithm::BLOCK_SCAN_WARP_SCANS, .lookback_delay = cub::LookbackDelayPolicy{ - .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 350, .l2_write_latency = 450}}; + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 350, .l2_write_latency = 450}}; # else // _CCCL_STD_VER >= 2020 constexpr auto p2 = p1; # endif // _CCCL_STD_VER >= 2020 @@ -458,7 +458,7 @@ C2H_TEST("Test PartitionPolicy properties", "[partition][device]") .load_modifier = cub::CacheLoadModifier::LOAD_DEFAULT, .scan_algorithm = cub::BlockScanAlgorithm::BLOCK_SCAN_WARP_SCANS, .lookback_delay = cub::LookbackDelayPolicy{ - .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 350, .l2_write_latency = 450}}; + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 350, .l2_write_latency = 450}}; # else // _CCCL_STD_VER >= 2020 constexpr auto p2 = p1; # endif // _CCCL_STD_VER >= 2020 diff --git a/cub/test/catch2_test_device_radix_sort_custom_policy_hub.cu b/cub/test/catch2_test_device_radix_sort_custom_policy_hub.cu index 1346a0e1a59..db5f06db1f2 100644 --- a/cub/test/catch2_test_device_radix_sort_custom_policy_hub.cu +++ b/cub/test/catch2_test_device_radix_sort_custom_policy_hub.cu @@ -31,14 +31,14 @@ struct my_policy_hub using HistogramPolicy = AgentRadixSortHistogramPolicy<256, 8, 1, KeyT, ONESWEEP_RADIX_BITS>; using ExclusiveSumPolicy = AgentRadixSortExclusiveSumPolicy<256, ONESWEEP_RADIX_BITS>; using OnesweepPolicy = AgentRadixSortOnesweepPolicy< - 256, - 21, - DominantT, - 1, - RADIX_RANK_MATCH_EARLY_COUNTS_ANY, - BLOCK_SCAN_WARP_SCANS, - RADIX_SORT_STORE_DIRECT, - ONESWEEP_RADIX_BITS>; + 256, + 21, + DominantT, + 1, + RADIX_RANK_MATCH_EARLY_COUNTS_ANY, + BLOCK_SCAN_WARP_SCANS, + RADIX_SORT_STORE_DIRECT, + ONESWEEP_RADIX_BITS>; using ScanPolicy = AgentScanPolicy<512, 23, diff --git a/cub/test/catch2_test_device_radix_sort_env.cu b/cub/test/catch2_test_device_radix_sort_env.cu index 615aa198737..1a059cc630d 100644 --- a/cub/test/catch2_test_device_radix_sort_env.cu +++ b/cub/test/catch2_test_device_radix_sort_env.cu @@ -1143,8 +1143,8 @@ std::size_t measure_allocated_bytes(CallableT&& run, PolicySelector policy_selec size_t bytes_allocated = 0; size_t bytes_deallocated = 0; auto env = stdexec::env{device_memory_resource{stream.get(), &bytes_allocated, &bytes_deallocated}, - stream, - cuda::execution::tune(policy_selector)}; + stream, + cuda::execution::tune(policy_selector)}; REQUIRE(cudaSuccess == run(env)); stream.sync(); CHECK(bytes_allocated > 0); @@ -1603,13 +1603,13 @@ C2H_TEST("Test RadixSortPolicy properties", "[radix_sort][device]") .threads_per_block = 256, .items_per_thread = 8, .private_partitions = 1, .radix_bits = 8}; constexpr auto p2_exclusive_sum = cub::RadixSortExclusiveSumPolicy{.threads_per_block = 256, .radix_bits = 8}; constexpr auto p2_onesweep = cub::RadixSortOnesweepPolicy{ - .threads_per_block = 256, - .items_per_thread = 21, - .store_algorithm = cub::RADIX_SORT_STORE_DIRECT, - .rank_algorithm = cub::RADIX_RANK_MATCH_EARLY_COUNTS_ANY, - .scan_algorithm = cub::BLOCK_SCAN_WARP_SCANS, - .rank_private_partitions = 1, - .radix_bits = 8}; + .threads_per_block = 256, + .items_per_thread = 21, + .store_algorithm = cub::RADIX_SORT_STORE_DIRECT, + .rank_algorithm = cub::RADIX_RANK_MATCH_EARLY_COUNTS_ANY, + .scan_algorithm = cub::BLOCK_SCAN_WARP_SCANS, + .rank_private_partitions = 1, + .radix_bits = 8}; constexpr auto p2_scan = cub::ScanPolicy{ .algorithm = cub::ScanAlgorithm::lookback, .lookback = diff --git a/cub/test/catch2_test_device_reduce_by_key_custom_policy_hub.cu b/cub/test/catch2_test_device_reduce_by_key_custom_policy_hub.cu index 69bd7199c30..05699f97410 100644 --- a/cub/test/catch2_test_device_reduce_by_key_custom_policy_hub.cu +++ b/cub/test/catch2_test_device_reduce_by_key_custom_policy_hub.cu @@ -54,16 +54,16 @@ C2H_TEST("DispatchReduceByKey::Dispatch: custom policy hub", "[reduce_by_key][de using policy_hub_t = my_policy_hub; using dispatch_t = cub::DispatchReduceByKey< - key_t*, - key_t*, - value_t*, - value_t*, - offset_t*, - equality_op_t, - reduction_op_t, - offset_t, - accum_t, - policy_hub_t>; + key_t*, + key_t*, + value_t*, + value_t*, + offset_t*, + equality_op_t, + reduction_op_t, + offset_t, + accum_t, + policy_hub_t>; size_t temp_size = 0; dispatch_t::Dispatch( diff --git a/cub/test/catch2_test_device_reduce_by_key_env_api.cu b/cub/test/catch2_test_device_reduce_by_key_env_api.cu index b3cfffc5c80..f715a23ae38 100644 --- a/cub/test/catch2_test_device_reduce_by_key_env_api.cu +++ b/cub/test/catch2_test_device_reduce_by_key_env_api.cu @@ -26,7 +26,7 @@ struct ReduceByKeyPolicySelector .load_modifier = cub::LOAD_DEFAULT, .scan_algorithm = cub::BLOCK_SCAN_WARP_SCANS, .lookback_delay = cub::LookbackDelayPolicy{ - .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 832, .l2_write_latency = 1165}}; + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 832, .l2_write_latency = 1165}}; } }; // example-end reduce-by-key-policy-selector diff --git a/cub/test/catch2_test_device_reduce_env.cu b/cub/test/catch2_test_device_reduce_env.cu index c54701185ea..70ed5342572 100644 --- a/cub/test/catch2_test_device_reduce_env.cu +++ b/cub/test/catch2_test_device_reduce_env.cu @@ -641,35 +641,35 @@ C2H_TEST("Device reduce not_guaranteed falls back when output type differs from auto kernels = cuda::std::array{ reinterpret_cast( cub::detail::reduce::DeviceReduceSingleTileKernel< - policy_t, - decltype(d_in.begin()), - decltype(d_out.begin()), - offset_t, - op_t, - init_value_t, - accumulator_t, - transform_t>), + policy_t, + decltype(d_in.begin()), + decltype(d_out.begin()), + offset_t, + op_t, + init_value_t, + accumulator_t, + transform_t>), reinterpret_cast( cub::detail::reduce::DeviceReduceKernel< - policy_t, - /* StableReductionOrder */ true, - decltype(d_in.begin()), - accumulator_t*, - offset_t, - offset_t, - op_t, - accumulator_t, - init_value_t, - transform_t>), + policy_t, + /* StableReductionOrder */ true, + decltype(d_in.begin()), + accumulator_t*, + offset_t, + offset_t, + op_t, + accumulator_t, + init_value_t, + transform_t>), reinterpret_cast( cub::detail::reduce::DeviceReduceSingleTileKernel< - policy_t, - accumulator_t*, - decltype(d_out.begin()), - int, // always used with int offset - op_t, - init_value_t, - accumulator_t>)}; + policy_t, + accumulator_t*, + decltype(d_out.begin()), + int, // always used with int offset + op_t, + init_value_t, + accumulator_t>)}; auto env = stdexec::env{cuda::execution::require(cuda::execution::determinism::not_guaranteed), allowed_kernels(kernels), @@ -703,35 +703,35 @@ C2H_TEST("Device sum not_guaranteed falls back when output type differs from acc auto kernels = cuda::std::array{ reinterpret_cast( cub::detail::reduce::DeviceReduceSingleTileKernel< - policy_t, - decltype(d_in.begin()), - decltype(d_out.begin()), - offset_t, - op_t, - init_value_t, - accumulator_t, - transform_t>), + policy_t, + decltype(d_in.begin()), + decltype(d_out.begin()), + offset_t, + op_t, + init_value_t, + accumulator_t, + transform_t>), reinterpret_cast( cub::detail::reduce::DeviceReduceKernel< - policy_t, - /* StableReductionOrder */ true, - decltype(d_in.begin()), - accumulator_t*, - offset_t, - offset_t, - op_t, - accumulator_t, - init_value_t, - transform_t>), + policy_t, + /* StableReductionOrder */ true, + decltype(d_in.begin()), + accumulator_t*, + offset_t, + offset_t, + op_t, + accumulator_t, + init_value_t, + transform_t>), reinterpret_cast( cub::detail::reduce::DeviceReduceSingleTileKernel< - policy_t, - accumulator_t*, - decltype(d_out.begin()), - int, // always used with int offset - op_t, - init_value_t, - accumulator_t>)}; + policy_t, + accumulator_t*, + decltype(d_out.begin()), + int, // always used with int offset + op_t, + init_value_t, + accumulator_t>)}; auto env = stdexec::env{cuda::execution::require(cuda::execution::determinism::not_guaranteed), allowed_kernels(kernels), @@ -1182,7 +1182,7 @@ C2H_TEST("Test ReduceByKeyPolicy properties", "[reduce][device]") .load_modifier = cub::LOAD_DEFAULT, .scan_algorithm = cub::BLOCK_SCAN_WARP_SCANS, .lookback_delay = cub::LookbackDelayPolicy{ - .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 832, .l2_write_latency = 1165}}; + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 832, .l2_write_latency = 1165}}; # else // _CCCL_STD_VER >= 2020 constexpr auto p2 = p1; # endif // _CCCL_STD_VER >= 2020 diff --git a/cub/test/catch2_test_device_run_length_encode_env.cu b/cub/test/catch2_test_device_run_length_encode_env.cu index f1fc2822c7f..667e6cae9d3 100644 --- a/cub/test/catch2_test_device_run_length_encode_env.cu +++ b/cub/test/catch2_test_device_run_length_encode_env.cu @@ -327,7 +327,7 @@ C2H_TEST("Test RleEncodePolicy properties", "[run_length_encode][device]") .load_modifier = cub::LOAD_DEFAULT, .scan_algorithm = cub::BLOCK_SCAN_WARP_SCANS, .lookback_delay = cub::LookbackDelayPolicy{ - .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 832, .l2_write_latency = 1165}}; + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 832, .l2_write_latency = 1165}}; # else // _CCCL_STD_VER >= 2020 constexpr auto p2 = p1; # endif // _CCCL_STD_VER >= 2020 @@ -376,7 +376,7 @@ C2H_TEST("Test RleNonTrivialRunsPolicy properties", "[run_length_encode][device] .store_with_time_slicing = false, .scan_algorithm = cub::BlockScanAlgorithm::BLOCK_SCAN_WARP_SCANS, .lookback_delay = cub::LookbackDelayPolicy{ - .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 350, .l2_write_latency = 450}}; + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 350, .l2_write_latency = 450}}; # else // _CCCL_STD_VER >= 2020 constexpr auto p2 = p1; # endif // _CCCL_STD_VER >= 2020 diff --git a/cub/test/catch2_test_device_scan_by_key_env.cu b/cub/test/catch2_test_device_scan_by_key_env.cu index eca2946cb27..cf224711ee5 100644 --- a/cub/test/catch2_test_device_scan_by_key_env.cu +++ b/cub/test/catch2_test_device_scan_by_key_env.cu @@ -376,7 +376,7 @@ C2H_TEST("Test ScanByKeyPolicy properties", "[scan][by_key][device]") .store_algorithm = cub::BlockStoreAlgorithm::BLOCK_STORE_DIRECT, .scan_algorithm = cub::BlockScanAlgorithm::BLOCK_SCAN_RAKING, .lookback_delay = cub::LookbackDelayPolicy{ - .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 832, .l2_write_latency = 1165}}; + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 832, .l2_write_latency = 1165}}; # else // _CCCL_STD_VER >= 2020 constexpr auto p2 = p1; # endif // _CCCL_STD_VER >= 2020 diff --git a/cub/test/catch2_test_device_scan_by_key_env_api.cu b/cub/test/catch2_test_device_scan_by_key_env_api.cu index a587bfef3b8..101272192d4 100644 --- a/cub/test/catch2_test_device_scan_by_key_env_api.cu +++ b/cub/test/catch2_test_device_scan_by_key_env_api.cu @@ -144,7 +144,7 @@ struct ScanByKeyPolicySelector .store_algorithm = cub::BLOCK_STORE_WARP_TRANSPOSE, .scan_algorithm = cub::BLOCK_SCAN_WARP_SCANS, .lookback_delay = cub::LookbackDelayPolicy{ - .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 832, .l2_write_latency = 1165}}; + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 832, .l2_write_latency = 1165}}; } }; // example-end exclusive-sum-by-key-policy-selector diff --git a/cub/test/catch2_test_device_scan_env.cu b/cub/test/catch2_test_device_scan_env.cu index 9fec5ba3aa5..aa24a8fd57c 100644 --- a/cub/test/catch2_test_device_scan_env.cu +++ b/cub/test/catch2_test_device_scan_env.cu @@ -629,7 +629,7 @@ C2H_TEST("Test ScanPolicy properties", "[scan][device]") .store_algorithm = cub::BlockStoreAlgorithm::BLOCK_STORE_DIRECT, .scan_algorithm = cub::BlockScanAlgorithm::BLOCK_SCAN_RAKING, .lookback_delay = cub::LookbackDelayPolicy{ - .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 832, .l2_write_latency = 1165}}; + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 832, .l2_write_latency = 1165}}; constexpr auto p2_la = cub::ScanLookaheadPolicy{ .reduce_and_scan_warps = 3, .items_per_thread = 8, diff --git a/cub/test/catch2_test_device_segmented_radix_sort_custom_policy_hub.cu b/cub/test/catch2_test_device_segmented_radix_sort_custom_policy_hub.cu index 92e84c3787f..41e01a39259 100644 --- a/cub/test/catch2_test_device_segmented_radix_sort_custom_policy_hub.cu +++ b/cub/test/catch2_test_device_segmented_radix_sort_custom_policy_hub.cu @@ -32,14 +32,14 @@ struct my_policy_hub using HistogramPolicy = AgentRadixSortHistogramPolicy<256, 8, 1, KeyT, ONESWEEP_RADIX_BITS>; using ExclusiveSumPolicy = AgentRadixSortExclusiveSumPolicy<256, ONESWEEP_RADIX_BITS>; using OnesweepPolicy = AgentRadixSortOnesweepPolicy< - 256, - 21, - KeyT, - 1, - RADIX_RANK_MATCH_EARLY_COUNTS_ANY, - BLOCK_SCAN_WARP_SCANS, - RADIX_SORT_STORE_DIRECT, - ONESWEEP_RADIX_BITS>; + 256, + 21, + KeyT, + 1, + RADIX_RANK_MATCH_EARLY_COUNTS_ANY, + BLOCK_SCAN_WARP_SCANS, + RADIX_SORT_STORE_DIRECT, + ONESWEEP_RADIX_BITS>; using ScanPolicy = AgentScanPolicy<512, 23, @@ -118,13 +118,13 @@ C2H_TEST("DispatchSegmentedRadixSort::Dispatch: custom policy hub", "[keys][segm using policy_hub_t = my_policy_hub; using dispatch_t = DispatchSegmentedRadixSort< - SortOrder::Ascending, - key_t, - value_t, - const offset_t*, - const offset_t*, - segment_size_t, - policy_hub_t>; + SortOrder::Ascending, + key_t, + value_t, + const offset_t*, + const offset_t*, + segment_size_t, + policy_hub_t>; size_t temp_size = 0; dispatch_t::Dispatch( diff --git a/cub/test/catch2_test_device_segmented_reduce_custom_policy_hub.cu b/cub/test/catch2_test_device_segmented_reduce_custom_policy_hub.cu index 0da58bc31e9..ce58f4eb932 100644 --- a/cub/test/catch2_test_device_segmented_reduce_custom_policy_hub.cu +++ b/cub/test/catch2_test_device_segmented_reduce_custom_policy_hub.cu @@ -52,15 +52,15 @@ C2H_TEST("DispatchSegmentedReduce::Dispatch: custom policy hub", "[segmented][re using policy_hub_t = my_policy_hub; using dispatch_t = DispatchSegmentedReduce< - input_t*, - output_t*, - const offset_t*, - const offset_t*, - offset_t, - reduction_t, - output_t, - accum_t, - policy_hub_t>; + input_t*, + output_t*, + const offset_t*, + const offset_t*, + offset_t, + reduction_t, + output_t, + accum_t, + policy_hub_t>; size_t temp_size = 0; dispatch_t::Dispatch( diff --git a/cub/test/catch2_test_device_segmented_scan_api.cu b/cub/test/catch2_test_device_segmented_scan_api.cu index 31844aaca19..fc8d2d15afe 100644 --- a/cub/test/catch2_test_device_segmented_scan_api.cu +++ b/cub/test/catch2_test_device_segmented_scan_api.cu @@ -356,7 +356,9 @@ C2H_TEST("cub::DeviceSegmentedScan::ExclusiveSegmentedScan API with two offsets * over boolean values stored as bits in unsigned integer, where addition is bitwise XOR. * Each unsigned integer represents 32-long tuple of GF(2) values */ - auto scan_op = [] __host__ __device__(unsigned v1, unsigned v2) -> unsigned { return v1 ^ v2; }; + auto scan_op = [] __host__ __device__(unsigned v1, unsigned v2) -> unsigned { + return v1 ^ v2; + }; unsigned init_value = 0u; // 128 input elements @@ -432,7 +434,9 @@ C2H_TEST("cub::DeviceSegmentedScan::InclusiveSegmentedScan API with three offset auto d_in = input.begin(); auto d_out = output.begin(); - auto scan_op = [] __host__ __device__(float v1, float v2) noexcept -> float { return cuda::maximum<>{}(v1, v2); }; + auto scan_op = [] __host__ __device__(float v1, float v2) noexcept -> float { + return cuda::maximum<>{}(v1, v2); + }; void* temp_storage = nullptr; size_t temp_storage_bytes; diff --git a/cub/test/catch2_test_device_select_env.cu b/cub/test/catch2_test_device_select_env.cu index ba56728d8fb..7bf54fe8501 100644 --- a/cub/test/catch2_test_device_select_env.cu +++ b/cub/test/catch2_test_device_select_env.cu @@ -1299,7 +1299,7 @@ C2H_TEST("Test UniqueByKeyPolicy properties", "[select_unique_by_key][device]") .load_modifier = cub::CacheLoadModifier::LOAD_DEFAULT, .scan_algorithm = cub::BlockScanAlgorithm::BLOCK_SCAN_WARP_SCANS, .lookback_delay = cub::LookbackDelayPolicy{ - .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 350, .l2_write_latency = 450}}; + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 350, .l2_write_latency = 450}}; # else // _CCCL_STD_VER >= 2020 constexpr auto p2 = p1; # endif // _CCCL_STD_VER >= 2020 @@ -1348,7 +1348,7 @@ C2H_TEST("Test SelectPolicy properties", "[select][device]") .load_modifier = cub::CacheLoadModifier::LOAD_DEFAULT, .scan_algorithm = cub::BlockScanAlgorithm::BLOCK_SCAN_WARP_SCANS, .lookback_delay = cub::LookbackDelayPolicy{ - .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 350, .l2_write_latency = 450}}; + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 350, .l2_write_latency = 450}}; # else // _CCCL_STD_VER >= 2020 constexpr auto p2 = p1; # endif // _CCCL_STD_VER >= 2020 diff --git a/cub/test/catch2_test_device_select_env_api.cu b/cub/test/catch2_test_device_select_env_api.cu index d55fc468d6e..4b18fe1bd55 100644 --- a/cub/test/catch2_test_device_select_env_api.cu +++ b/cub/test/catch2_test_device_select_env_api.cu @@ -493,7 +493,7 @@ struct UniqueByKeyPolicySelector .load_modifier = cub::LOAD_DEFAULT, .scan_algorithm = cub::BLOCK_SCAN_WARP_SCANS, .lookback_delay = cub::LookbackDelayPolicy{ - .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 350, .l2_write_latency = 450}}; + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 350, .l2_write_latency = 450}}; } }; // example-end unique-by-key-policy-selector diff --git a/cub/test/catch2_test_device_three_way_partition_custom_policy_hub.cu b/cub/test/catch2_test_device_three_way_partition_custom_policy_hub.cu index 4cd5e2f9580..23d7e4bf6c9 100644 --- a/cub/test/catch2_test_device_three_way_partition_custom_policy_hub.cu +++ b/cub/test/catch2_test_device_three_way_partition_custom_policy_hub.cu @@ -83,15 +83,15 @@ C2H_TEST("DispatchThreeWayPartitionIf::Dispatch: custom policy hub", "[partition using policy_hub_t = my_policy_hub; using dispatch_t = DispatchThreeWayPartitionIf< - value_t*, - value_t*, - value_t*, - value_t*, - offset_t*, - less_than_zero_t, - equal_zero_t, - offset_t, - policy_hub_t>; + value_t*, + value_t*, + value_t*, + value_t*, + offset_t*, + less_than_zero_t, + equal_zero_t, + offset_t, + policy_hub_t>; size_t temp_size = 0; dispatch_t::Dispatch( diff --git a/cub/test/test_block_radix_rank.cu b/cub/test/test_block_radix_rank.cu index 58f2bb8e954..4f08907e24b 100644 --- a/cub/test/test_block_radix_rank.cu +++ b/cub/test/test_block_radix_rank.cu @@ -287,7 +287,7 @@ void Test() Test(); Test(); - Test(cuda::std::bool_constant < (ThreadsPerBlock % 32) == 0 > {}); + Test(cuda::std::bool_constant<(ThreadsPerBlock % 32) == 0>{}); } int main(int argc, char** argv) diff --git a/cudax/include/cuda/experimental/__copy/copy_optimized.cuh b/cudax/include/cuda/experimental/__copy/copy_optimized.cuh index 4184600340a..997f35decbd 100644 --- a/cudax/include/cuda/experimental/__copy/copy_optimized.cuh +++ b/cudax/include/cuda/experimental/__copy/copy_optimized.cuh @@ -119,15 +119,15 @@ _CCCL_HOST_API void __copy_optimized( const auto __grid_size = ::cuda::ceil_div(__tensor_size, _ExtentT{__block_size}); const auto __config = ::cuda::make_config(::cuda::block_dims<__block_size>(), ::cuda::grid_dims(__grid_size)); const auto& __kernel = ::cuda::experimental::__copy_optimized_kernel< - decltype(__config), - _TpIn, - _TpOut, - _SrcAccessor, - _DstAccessor, - _ExtentT, - _StrideTIn, - _StrideTOut, - _Rank>; + decltype(__config), + _TpIn, + _TpOut, + _SrcAccessor, + _DstAccessor, + _ExtentT, + _StrideTIn, + _StrideTOut, + _Rank>; ::cuda::launch( __stream, diff --git a/cudax/include/cuda/experimental/__cuco/hyperloglog.cuh b/cudax/include/cuda/experimental/__cuco/hyperloglog.cuh index 8f3dafbd626..bd06512a320 100644 --- a/cudax/include/cuda/experimental/__cuco/hyperloglog.cuh +++ b/cudax/include/cuda/experimental/__cuco/hyperloglog.cuh @@ -160,12 +160,12 @@ public: _MemoryResource_&& __memory_resource, precision __precision, const _Policy& __policy = {}) - : __sketch_buffer{__stream, - ::cuda::std::forward<_MemoryResource_>(__memory_resource), - ref_type<>::sketch_bytes( - __precision_in_bounds(__precision, "HyperLogLog precision must be in [4, 18]")) - / sizeof(register_type), - ::cuda::no_init} + : __sketch_buffer{ + __stream, + ::cuda::std::forward<_MemoryResource_>(__memory_resource), + ref_type<>::sketch_bytes(__precision_in_bounds(__precision, "HyperLogLog precision must be in [4, 18]")) + / sizeof(register_type), + ::cuda::no_init} , __ref{::cuda::std::as_writable_bytes(::cuda::std::span{__sketch_buffer.data(), __sketch_buffer.size()}), __policy} { diff --git a/cudax/include/cuda/experimental/__execution/get_completion_signatures.cuh b/cudax/include/cuda/experimental/__execution/get_completion_signatures.cuh index f2a626e3668..7104dde2649 100644 --- a/cudax/include/cuda/experimental/__execution/get_completion_signatures.cuh +++ b/cudax/include/cuda/experimental/__execution/get_completion_signatures.cuh @@ -179,7 +179,7 @@ template # define _CUDAX_PP_EAT_AUTO_auto(_ID) _ID _CCCL_PP_EAT _CCCL_PP_LPAREN # define _CUDAX_PP_EXPAND_AUTO_auto(_ID) auto _ID -# define _CUDAX_LET_COMPLETIONS_ID(...) _CCCL_PP_EXPAND(_CCCL_PP_CAT(_CUDAX_PP_EAT_AUTO_, __VA_ARGS__) _CCCL_PP_RPAREN) +# define _CUDAX_LET_COMPLETIONS_ID(...) _CCCL_PP_EXPAND(_CCCL_PP_CAT(_CUDAX_PP_EAT_AUTO_, __VA_ARGS__) _CCCL_PP_RPAREN) # define _CUDAX_LET_COMPLETIONS(...) \ if constexpr (_CCCL_PP_CAT(_CUDAX_PP_EXPAND_AUTO_, __VA_ARGS__); \ diff --git a/cudax/include/cuda/experimental/__graph/graph_builder_ref.cuh b/cudax/include/cuda/experimental/__graph/graph_builder_ref.cuh index 5727199e116..5e26dda2980 100644 --- a/cudax/include/cuda/experimental/__graph/graph_builder_ref.cuh +++ b/cudax/include/cuda/experimental/__graph/graph_builder_ref.cuh @@ -267,8 +267,9 @@ private: //! \throws cuda::std::cuda_error if `cudaGraphAddChildGraphNode` fails. #if _CCCL_CTK_AT_LEAST(12, 2) template - [[nodiscard]] _CCCL_HOST_API auto - __add_to_graph(cudaGraph_t __parent, ::cuda::std::span __deps) -> graph_node_ref + [[nodiscard]] + _CCCL_HOST_API auto __add_to_graph(cudaGraph_t __parent, ::cuda::std::span __deps) + -> graph_node_ref { ::CUgraphNodeParams __params{}; __params.type = ::CU_GRAPH_NODE_TYPE_GRAPH; diff --git a/cudax/test/execution/test_bulk.cu b/cudax/test/execution/test_bulk.cu index 292ea4215a5..3a002afa005 100644 --- a/cudax/test/execution/test_bulk.cu +++ b/cudax/test/execution/test_bulk.cu @@ -268,7 +268,7 @@ void bulk_can_be_used_with_a_function() auto sndr = ex::just(&counter1) // | ex::bulk(ex::par, n, function); - auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{&counter1}); + auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{&counter1}); ex::start(op); for (int i : counter1) @@ -285,7 +285,7 @@ void bulk_chunked_can_be_used_with_a_function() auto sndr = ex::just(&counter2) // | ex::bulk_chunked(ex::par, n, function_range); - auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{&counter2}); + auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{&counter2}); ex::start(op); for (const auto c : counter2) @@ -302,7 +302,7 @@ void bulk_unchunked_can_be_used_with_a_function() auto sndr = ex::just(&counter3) // | ex::bulk_unchunked(ex::par, n, function); - auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{&counter3}); + auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{&counter3}); ex::start(op); for (const auto c : counter3) @@ -319,7 +319,7 @@ void bulk_can_be_used_with_a_function_object() auto sndr = ex::just() // | ex::bulk(ex::par, n, fn); - auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{}); + auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{}); ex::start(op); for (int i : counter) @@ -336,7 +336,7 @@ void bulk_chunked_can_be_used_with_a_function_object() auto sndr = ex::just() // | ex::bulk_chunked(ex::par, n, fn); - auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{}); + auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{}); ex::start(op); for (const auto c : counter) @@ -353,7 +353,7 @@ void bulk_unchunked_can_be_used_with_a_function_object() auto sndr = ex::just() // | ex::bulk_unchunked(ex::par, n, fn); - auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{}); + auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{}); ex::start(op); for (const auto c : counter) @@ -372,7 +372,7 @@ void bulk_can_be_used_with_a_lambda() | ex::bulk(ex::par, n, [&](int i) { counter[i]++; }); - auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{}); + auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{}); ex::start(op); for (int i : counter) @@ -393,7 +393,7 @@ void bulk_chunked_can_be_used_with_a_lambda() counter[b++]++; } }); - auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{}); + auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{}); ex::start(op); for (int i = 0; i < n; i++) @@ -411,7 +411,7 @@ void bulk_unchunked_can_be_used_with_a_lambda() | ex::bulk_unchunked(ex::par, n, [&](int i) { counter[i]++; }); - auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{}); + auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{}); ex::start(op); for (int i = 0; i < n; i++) @@ -476,7 +476,7 @@ void bulk_forwards_values() (*counter)[i]++; } }); - auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{magic_number, &counter}); + auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{magic_number, &counter}); ex::start(op); for (int i : counter) @@ -501,7 +501,7 @@ void bulk_chunked_forwards_values() } } }); - auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{magic_number}); + auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{magic_number}); ex::start(op); for (const auto c : counter) @@ -523,7 +523,7 @@ void bulk_unchunked_forwards_values() counter[i]++; } }); - auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{magic_number}); + auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{magic_number}); ex::start(op); for (const auto c : counter) @@ -544,7 +544,7 @@ void bulk_forwards_values_that_can_be_taken_by_reference() | ex::bulk(ex::par, n, [&](std::size_t i, ::cuda::std::array& vals) { vals[i] = static_cast(i); }); - auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{vals_expected}); + auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{vals_expected}); ex::start(op); } @@ -561,7 +561,7 @@ void bulk_chunked_forwards_values_that_can_be_taken_by_reference() vals[b] = static_cast(b); } }); - auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{vals_expected}); + auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{vals_expected}); ex::start(op); } @@ -575,7 +575,7 @@ void bulk_unchunked_forwards_values_that_can_be_taken_by_reference() | ex::bulk_unchunked(ex::par, n, [&](std::size_t i, ::cuda::std::array& vals) { vals[i] = static_cast(i); }); - auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{vals_expected}); + auto op = ex::connect(cuda::std::move(sndr), checked_value_receiver{vals_expected}); ex::start(op); } @@ -630,7 +630,7 @@ void bulk_can_throw_and_set_error_will_be_called() | ex::bulk(ex::par, n, [](int) -> int { throw std::logic_error{"err"}; }); - auto op = ex::connect(cuda::std::move(sndr), checked_error_receiver{std::logic_error{"err"}}); + auto op = ex::connect(cuda::std::move(sndr), checked_error_receiver{std::logic_error{"err"}}); ex::start(op); } @@ -642,7 +642,7 @@ void bulk_chunked_can_throw_and_set_error_will_be_called() | ex::bulk_chunked(ex::par, n, [](int, int) -> int { throw std::logic_error{"err"}; }); - auto op = ex::connect(cuda::std::move(sndr), checked_error_receiver{std::logic_error{"err"}}); + auto op = ex::connect(cuda::std::move(sndr), checked_error_receiver{std::logic_error{"err"}}); ex::start(op); } @@ -654,7 +654,7 @@ void bulk_unchunked_can_throw_and_set_error_will_be_called() | ex::bulk_unchunked(ex::par, n, [](int) -> int { throw std::logic_error{"err"}; }); - auto op = ex::connect(cuda::std::move(sndr), checked_error_receiver{std::logic_error{"err"}}); + auto op = ex::connect(cuda::std::move(sndr), checked_error_receiver{std::logic_error{"err"}}); ex::start(op); } #endif // _CCCL_HAS_EXCEPTIONS() && !defined(__CUDA_ARCH__) @@ -668,7 +668,7 @@ void bulk_function_is_not_called_on_error() | ex::bulk(ex::par, n, [&called](int) { called++; }); - auto op = ex::connect(cuda::std::move(sndr), checked_error_receiver{string{"err"}}); + auto op = ex::connect(cuda::std::move(sndr), checked_error_receiver{string{"err"}}); ex::start(op); } @@ -681,7 +681,7 @@ void bulk_chunked_function_is_not_called_on_error() | ex::bulk_chunked(ex::par, n, [&called](int, int) { called++; }); - auto op = ex::connect(cuda::std::move(sndr), checked_error_receiver{string{"err"}}); + auto op = ex::connect(cuda::std::move(sndr), checked_error_receiver{string{"err"}}); ex::start(op); } @@ -694,7 +694,7 @@ void bulk_unchunked_function_is_not_called_on_error() | ex::bulk_unchunked(ex::par, n, [&called](int) { called++; }); - auto op = ex::connect(cuda::std::move(sndr), checked_error_receiver{string{"err"}}); + auto op = ex::connect(cuda::std::move(sndr), checked_error_receiver{string{"err"}}); ex::start(op); } @@ -707,7 +707,7 @@ void bulk_function_in_not_called_on_stop() | ex::bulk(ex::par, n, [&called](int) { called++; }); - auto op = ex::connect(cuda::std::move(sndr), checked_stopped_receiver{}); + auto op = ex::connect(cuda::std::move(sndr), checked_stopped_receiver{}); ex::start(op); } @@ -720,7 +720,7 @@ void bulk_chunked_function_in_not_called_on_stop() | ex::bulk_chunked(ex::par, n, [&called](int, int) { called++; }); - auto op = ex::connect(cuda::std::move(sndr), checked_stopped_receiver{}); + auto op = ex::connect(cuda::std::move(sndr), checked_stopped_receiver{}); ex::start(op); } @@ -733,7 +733,7 @@ void bulk_unchunked_function_in_not_called_on_stop() | ex::bulk_unchunked(ex::par, n, [&called](int) { called++; }); - auto op = ex::connect(cuda::std::move(sndr), checked_stopped_receiver{}); + auto op = ex::connect(cuda::std::move(sndr), checked_stopped_receiver{}); ex::start(op); } diff --git a/cudax/test/execution/test_continues_on.cu b/cudax/test/execution/test_continues_on.cu index 7ab87851f94..b0ff721a6d7 100644 --- a/cudax/test/execution/test_continues_on.cu +++ b/cudax/test/execution/test_continues_on.cu @@ -41,11 +41,11 @@ C2H_TEST("continues_on can be piped", "[adaptors][continues_on]") bool called{false}; auto sched = impulse_scheduler{}; auto snd = ex::just(13) // - | ex::continues_on(sched) // - | ex::then([&](int val) { + | ex::continues_on(sched) // + | ex::then([&](int val) { called = true; return val; - }); + }); static_assert(ex::get_completion_behavior() == ex::completion_behavior::asynchronous); // Start the operation auto op = ex::connect(std::move(snd), checked_value_receiver{13}); diff --git a/cudax/test/execution/test_let_value.cu b/cudax/test/execution/test_let_value.cu index df48014131a..05ddf47b3f6 100644 --- a/cudax/test/execution/test_let_value.cu +++ b/cudax/test/execution/test_let_value.cu @@ -177,7 +177,7 @@ C2H_TEST("let_value can throw, and set_error will be called", "[adaptors][let_va | ex::let_value([](int&) -> decltype(ex::just(0)) { throw std::logic_error{"err"}; }); - auto op = ex::connect(std::move(sndr), checked_error_receiver{std::logic_error{"err"}}); + auto op = ex::connect(std::move(sndr), checked_error_receiver{std::logic_error{"err"}}); ex::start(op); } @@ -187,7 +187,7 @@ C2H_TEST("let_value can be used with just_error", "[adaptors][let_value]") | ex::let_value([]() { return ex::just(17); }); - auto op = ex::connect(std::move(sndr), checked_error_receiver{std::string{"err"}}); + auto op = ex::connect(std::move(sndr), checked_error_receiver{std::string{"err"}}); ex::start(op); } @@ -210,7 +210,7 @@ C2H_TEST("let_value function is not called on error", "[adaptors][let_value]") called = true; return ex::just(x + 5); }); - auto op = ex::connect(std::move(sndr), checked_error_receiver{-1}); + auto op = ex::connect(std::move(sndr), checked_error_receiver{-1}); ex::start(op); CHECK_FALSE(called); } @@ -225,7 +225,7 @@ C2H_TEST("let_value function is not called when cancelled", "[adaptors][let_valu called = true; return ex::just(x + 5); }); - auto op = ex::connect(std::move(sndr), checked_stopped_receiver{}); + auto op = ex::connect(std::move(sndr), checked_stopped_receiver{}); ex::start(op); CHECK_FALSE(called); } @@ -426,8 +426,8 @@ C2H_TEST("let_value works when the function returns a dependent sender", "[adapt { auto sndr = ex::write_env(ex::just() | ex::let_value([] { return ex::read_env(test_query); - }), - ex::prop{test_query, 42}); + }), + ex::prop{test_query, 42}); auto [result] = ex::sync_wait(std::move(sndr)).value(); CHECK(result == 42); } @@ -486,11 +486,11 @@ C2H_TEST("let_value predecessor's domain is accessible via the receiver connecte auto attrs = ex::prop{ex::get_completion_domain, let_value_test_domain2{}}; using Sndr2 = decltype(ex::read_env(ex::get_domain)); - auto sndr = ex::just() // - | ex::write_attrs(attrs) // - | ex::let_value([]() noexcept -> Sndr2 { + auto sndr = ex::just() // + | ex::write_attrs(attrs) // + | ex::let_value([]() noexcept -> Sndr2 { return ex::read_env(ex::get_domain); - }); + }); auto [result] = ex::sync_wait(std::move(sndr)).value(); static_assert(::cuda::std::is_same_v); (void) result; @@ -502,7 +502,7 @@ C2H_TEST("let_value has the correct completion domain", "[adaptors][let_value]") auto attrs = ex::prop{ex::get_completion_domain, let_value_test_domain{}}; auto sndr = ex::just() | ex::let_value([=] { return ex::write_attrs(ex::just(), attrs); - }); + }); auto dom = ex::get_completion_domain(ex::get_env(sndr)); static_assert(::cuda::std::is_same_v); } diff --git a/cudax/test/execution/test_on.cu b/cudax/test/execution/test_on.cu index 99273951639..6a3a7985203 100644 --- a/cudax/test/execution/test_on.cu +++ b/cudax/test/execution/test_on.cu @@ -26,14 +26,14 @@ auto const main_thread_id = ::std::this_thread::get_id(); void simple_start_on_thread_test() { ex::thread_context ctx; - auto sch = ctx.get_scheduler(); - auto sndr = ex::on(sch, ex::just() | ex::then([] { + auto sch = ctx.get_scheduler(); + auto sndr = ex::on(sch, ex::just() | ex::then([] { CHECK(::std::this_thread::get_id() != main_thread_id); - })) - | ex::then([]() -> int { + })) + | ex::then([]() -> int { CHECK(::std::this_thread::get_id() == main_thread_id); return 42; - }); + }); auto [result] = ex::sync_wait(std::move(sndr)).value(); CHECK(result == 42); } @@ -41,14 +41,14 @@ void simple_start_on_thread_test() void simple_continue_on_thread_test() { ex::thread_context ctx; - auto sch = ctx.get_scheduler(); - auto sndr = ex::just() | ex::on(sch, ex::then([] { + auto sch = ctx.get_scheduler(); + auto sndr = ex::just() | ex::on(sch, ex::then([] { CHECK(::std::this_thread::get_id() != main_thread_id); - })) - | ex::then([]() -> int { + })) + | ex::then([]() -> int { CHECK(::std::this_thread::get_id() == main_thread_id); return 42; - }); + }); auto [result] = ex::sync_wait(std::move(sndr)).value(); CHECK(result == 42); } @@ -56,13 +56,13 @@ void simple_continue_on_thread_test() void simple_start_on_stream_test() { cudax::stream str{cuda::device_ref(0)}; - auto sch = cudax::stream_ref{str}; - auto sndr = ex::on(sch, ex::just(42) | ex::then([] __host__ __device__(int i) noexcept -> int { + auto sch = cudax::stream_ref{str}; + auto sndr = ex::on(sch, ex::just(42) | ex::then([] __host__ __device__(int i) noexcept -> int { return _on_device() ? i : -i; - })) - | ex::then([] __host__ __device__(int i) noexcept -> int { + })) + | ex::then([] __host__ __device__(int i) noexcept -> int { return _on_device() ? -1 : i; - }); + }); auto [result] = ex::sync_wait(std::move(sndr)).value(); CHECK(result == 42); } @@ -70,13 +70,13 @@ void simple_start_on_stream_test() void simple_continue_on_stream_test() { cudax::stream str{cuda::device_ref(0)}; - auto sch = cudax::stream_ref{str}; - auto sndr = ex::just(42) | ex::on(sch, ex::then([] __host__ __device__(int i) noexcept -> int { + auto sch = cudax::stream_ref{str}; + auto sndr = ex::just(42) | ex::on(sch, ex::then([] __host__ __device__(int i) noexcept -> int { return _on_device() ? i : -i; - })) - | ex::then([] __host__ __device__(int i) noexcept -> int { + })) + | ex::then([] __host__ __device__(int i) noexcept -> int { return _on_device() ? -1 : i; - }); + }); auto [result] = ex::sync_wait(std::move(sndr)).value(); CHECK(result == 42); } @@ -84,14 +84,14 @@ void simple_continue_on_stream_test() void test_continues_on_updates_env() { ex::thread_context ctx; - auto sch = ctx.get_scheduler(); - auto sndr = ex::just() | ex::on(sch, ex::let_value([] { + auto sch = ctx.get_scheduler(); + auto sndr = ex::just() | ex::on(sch, ex::let_value([] { return ex::read_env(ex::get_scheduler); - })) - | ex::then([](auto sch2) -> int { + })) + | ex::then([](auto sch2) -> int { STATIC_REQUIRE(cuda::std::same_as); return 42; - }); + }); auto [result] = ex::sync_wait(std::move(sndr)).value(); CHECK(result == 42); } diff --git a/cudax/test/execution/test_starts_on.cu b/cudax/test/execution/test_starts_on.cu index f803e9d634c..7f3ea212bed 100644 --- a/cudax/test/execution/test_starts_on.cu +++ b/cudax/test/execution/test_starts_on.cu @@ -45,7 +45,7 @@ C2H_TEST("starts_on can be piped", "[adaptors][starts_on]") | ex::then([](int val) { return val * 2; }); - auto op = ex::connect(std::move(snd), checked_value_receiver{84}); + auto op = ex::connect(std::move(snd), checked_value_receiver{84}); ex::start(op); // The receiver checks if we receive the transformed value } diff --git a/cudax/test/execution/test_task_scheduler.cu b/cudax/test/execution/test_task_scheduler.cu index d125bd85c28..8946b1cfb72 100644 --- a/cudax/test/execution/test_task_scheduler.cu +++ b/cudax/test/execution/test_task_scheduler.cu @@ -37,7 +37,7 @@ C2H_TEST("task_scheduler starts work on the correct execution context", "[schedu ex::task_scheduler sched{ctx.get_scheduler()}; auto sndr = ex::starts_on(sched, ex::just() | ex::then([] { return ::std::this_thread::get_id(); - })); + })); auto [tid] = ex::sync_wait(cuda::std::move(sndr)).value(); CHECK(tid == ctx.get_id()); } diff --git a/cudax/test/execution/test_when_all.cu b/cudax/test/execution/test_when_all.cu index 5257cbcd6bd..0360d59079a 100644 --- a/cudax/test/execution/test_when_all.cu +++ b/cudax/test/execution/test_when_all.cu @@ -82,7 +82,7 @@ C2H_TEST("when_all completes when children complete", "[when_all]") called = true; return a + b + c; }); - auto op = ex::connect(std::move(snd), checked_value_receiver{41}); + auto op = ex::connect(std::move(snd), checked_value_receiver{41}); ex::start(op); // The when_all scheduler will complete only after 3 impulses CHECK_FALSE(called); diff --git a/cudax/test/stf/examples/07-cholesky-redux.cu b/cudax/test/stf/examples/07-cholesky-redux.cu index dc8510f4a2c..024a739d8ee 100644 --- a/cudax/test/stf/examples/07-cholesky-redux.cu +++ b/cudax/test/stf/examples/07-cholesky-redux.cu @@ -259,9 +259,9 @@ void DGEMM( // an accumulation with the add operator auto dep_c = (beta == 1.0) ? C.handle(C_row, C_col).relaxed(redux_op) : C.handle(C_row, C_col).rw(); auto t = ctx.task(exec_place::device(A.get_preferred_devid(C_row, C_col)), - A.handle(A_row, A_col).read(), - B.handle(B_row, B_col).read(), - dep_c); + A.handle(A_row, A_col).read(), + B.handle(B_row, B_col).read(), + dep_c); t.set_symbol("DGEMM"); t->*[transa, transb, alpha, beta](cudaStream_t s, auto sA, auto sB, auto sC) { EXPECT(sC.data_handle() != nullptr); @@ -672,7 +672,9 @@ int main(int argc, char** argv) if (check_result) { - auto rhs_vals = [] __host__ __device__(size_t row, size_t /*unused*/) { return 1.0 * (row + 1); }; + auto rhs_vals = [] __host__ __device__(size_t row, size_t /*unused*/) { + return 1.0 * (row + 1); + }; B_potrs.fill(rhs_vals); Bref_potrs.fill(rhs_vals); } diff --git a/cudax/test/stf/graph/epoch.cu b/cudax/test/stf/graph/epoch.cu index d8ea17dd274..32d6ad11ec5 100644 --- a/cudax/test/stf/graph/epoch.cu +++ b/cudax/test/stf/graph/epoch.cu @@ -36,7 +36,9 @@ int main() for (size_t k = 0; k < NITER; k++) { ctx.parallel_for(blocked_partition(), exec_place::current_device(), lA.shape(), lA.rw()) - ->*[] __host__ __device__(size_t i, slice A) { A(i) = cos(A(i)); }; + ->*[] __host__ __device__(size_t i, slice A) { + A(i) = cos(A(i)); + }; ctx.change_stage(); } diff --git a/libcudacxx/include/cuda/__barrier/barrier_block_scope.h b/libcudacxx/include/cuda/__barrier/barrier_block_scope.h index 834bea3fe0e..cba1c90336b 100644 --- a/libcudacxx/include/cuda/__barrier/barrier_block_scope.h +++ b/libcudacxx/include/cuda/__barrier/barrier_block_scope.h @@ -237,8 +237,8 @@ class barrier : public __bl } #if _CCCL_CUDA_COMPILATION() - [[nodiscard]] _CCCL_DEVICE_API _CCCL_FORCEINLINE bool - __try_wait_sm90(arrival_token __token, ::cuda::std::chrono::nanoseconds __nanosec) const + [[nodiscard]] _CCCL_DEVICE_API + _CCCL_FORCEINLINE bool __try_wait_sm90(arrival_token __token, ::cuda::std::chrono::nanoseconds __nanosec) const { if (::cuda::device::is_object_from(__barrier, ::cuda::device::address_space::shared)) { @@ -338,8 +338,8 @@ class barrier : public __bl } #if _CCCL_CUDA_COMPILATION() - [[nodiscard]] _CCCL_DEVICE_API _CCCL_FORCEINLINE bool - __try_wait_parity_sm90(bool __phase_parity, ::cuda::std::chrono::nanoseconds __nanosec) const + [[nodiscard]] _CCCL_DEVICE_API + _CCCL_FORCEINLINE bool __try_wait_parity_sm90(bool __phase_parity, ::cuda::std::chrono::nanoseconds __nanosec) const { if (::cuda::device::is_object_from(__barrier, ::cuda::device::address_space::shared)) { diff --git a/libcudacxx/include/cuda/__device/arch_id.h b/libcudacxx/include/cuda/__device/arch_id.h index 7d02af3c336..5b4d9843aa3 100644 --- a/libcudacxx/include/cuda/__device/arch_id.h +++ b/libcudacxx/include/cuda/__device/arch_id.h @@ -51,8 +51,8 @@ enum class arch_id : int _OP) " is deprecated and will be deleted in the next major release. Compare cuda::compute_capabilities of the " \ "given " \ "cuda::arch_id instead.") -[[nodiscard]] _CCCL_DEPRECATED_ARCH_ID_COMPARISONS(<) _CCCL_HOST_DEVICE_API constexpr bool -operator<(arch_id __lhs, arch_id __rhs) noexcept +[[nodiscard]] _CCCL_DEPRECATED_ARCH_ID_COMPARISONS(<) _CCCL_HOST_DEVICE_API + constexpr bool operator<(arch_id __lhs, arch_id __rhs) noexcept { return ::cuda::std::to_underlying(__lhs) < ::cuda::std::to_underlying(__rhs); } diff --git a/libcudacxx/include/cuda/__ptx/instructions/generated/clusterlaunchcontrol.h b/libcudacxx/include/cuda/__ptx/instructions/generated/clusterlaunchcontrol.h index 7c62b26f808..1636aab87bc 100644 --- a/libcudacxx/include/cuda/__ptx/instructions/generated/clusterlaunchcontrol.h +++ b/libcudacxx/include/cuda/__ptx/instructions/generated/clusterlaunchcontrol.h @@ -60,7 +60,8 @@ __device__ static inline bool clusterlaunchcontrol_query_cancel_is_canceled( B128 try_cancel_response); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void __cuda_ptx_clusterlaunchcontrol_query_cancel_is_canceled_is_not_supported_before_SM_100__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_clusterlaunchcontrol_query_cancel_is_canceled_is_not_supported_before_SM_100__(); template = true> _CCCL_DEVICE static inline bool clusterlaunchcontrol_query_cancel_is_canceled(_B128 __try_cancel_response) { @@ -95,8 +96,8 @@ __device__ static inline B32 clusterlaunchcontrol_query_cancel_get_first_ctaid_x B128 try_cancel_response); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_clusterlaunchcontrol_query_cancel_get_first_ctaid_x_is_not_supported_before_SM_100__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_clusterlaunchcontrol_query_cancel_get_first_ctaid_x_is_not_supported_before_SM_100__(); template = true, typename _B128, @@ -133,8 +134,8 @@ __device__ static inline B32 clusterlaunchcontrol_query_cancel_get_first_ctaid_y B128 try_cancel_response); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_clusterlaunchcontrol_query_cancel_get_first_ctaid_y_is_not_supported_before_SM_100__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_clusterlaunchcontrol_query_cancel_get_first_ctaid_y_is_not_supported_before_SM_100__(); template = true, typename _B128, @@ -171,8 +172,8 @@ __device__ static inline B32 clusterlaunchcontrol_query_cancel_get_first_ctaid_z B128 try_cancel_response); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_clusterlaunchcontrol_query_cancel_get_first_ctaid_z_is_not_supported_before_SM_100__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_clusterlaunchcontrol_query_cancel_get_first_ctaid_z_is_not_supported_before_SM_100__(); template = true, typename _B128, @@ -210,8 +211,8 @@ __device__ static inline void clusterlaunchcontrol_query_cancel_get_first_ctaid( B128 try_cancel_response); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_clusterlaunchcontrol_query_cancel_get_first_ctaid_is_not_supported_before_SM_100__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_clusterlaunchcontrol_query_cancel_get_first_ctaid_is_not_supported_before_SM_100__(); template = true, typename _B128, diff --git a/libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_tensor_multicast.h b/libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_tensor_multicast.h index 8ec514dd9ee..8fb72e9b04b 100644 --- a/libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_tensor_multicast.h +++ b/libcudacxx/include/cuda/__ptx/instructions/generated/cp_async_bulk_tensor_multicast.h @@ -138,8 +138,8 @@ __device__ static inline void cp_async_bulk_tensor( const uint16_t& ctaMask); */ #if __cccl_ptx_isa >= 800 -extern "C" _CCCL_DEVICE void -__cuda_ptx_cp_async_bulk_tensor_is_only_supported_on_SM_90a_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_cp_async_bulk_tensor_is_only_supported_on_SM_90a_100a_100f_103a_103f_110a_110f__(); template _CCCL_DEVICE static inline void cp_async_bulk_tensor( ::cuda::ptx::space_cluster_t, @@ -260,8 +260,8 @@ __device__ static inline void cp_async_bulk_tensor( const uint16_t& ctaMask); */ #if __cccl_ptx_isa >= 800 -extern "C" _CCCL_DEVICE void -__cuda_ptx_cp_async_bulk_tensor_is_only_supported_on_SM_90a_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_cp_async_bulk_tensor_is_only_supported_on_SM_90a_100a_100f_103a_103f_110a_110f__(); template _CCCL_DEVICE static inline void cp_async_bulk_tensor( ::cuda::ptx::space_cluster_t, @@ -385,8 +385,8 @@ __device__ static inline void cp_async_bulk_tensor( const uint16_t& ctaMask); */ #if __cccl_ptx_isa >= 800 -extern "C" _CCCL_DEVICE void -__cuda_ptx_cp_async_bulk_tensor_is_only_supported_on_SM_90a_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_cp_async_bulk_tensor_is_only_supported_on_SM_90a_100a_100f_103a_103f_110a_110f__(); template _CCCL_DEVICE static inline void cp_async_bulk_tensor( ::cuda::ptx::space_cluster_t, @@ -513,8 +513,8 @@ __device__ static inline void cp_async_bulk_tensor( const uint16_t& ctaMask); */ #if __cccl_ptx_isa >= 800 -extern "C" _CCCL_DEVICE void -__cuda_ptx_cp_async_bulk_tensor_is_only_supported_on_SM_90a_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_cp_async_bulk_tensor_is_only_supported_on_SM_90a_100a_100f_103a_103f_110a_110f__(); template _CCCL_DEVICE static inline void cp_async_bulk_tensor( ::cuda::ptx::space_cluster_t, diff --git a/libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_ld.h b/libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_ld.h index 18b45a7226f..8fd8cb63b4e 100644 --- a/libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_ld.h +++ b/libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_ld.h @@ -35,8 +35,8 @@ __device__ static inline void tcgen05_ld_16x64b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x64b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x64b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_16x64b_pack_16b(_B32 (&__out)[1], ::cuda::std::uint32_t __taddr) { @@ -87,8 +87,8 @@ __device__ static inline void tcgen05_ld_16x64b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x64b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x64b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_16x64b_pack_16b(_B32 (&__out)[2], ::cuda::std::uint32_t __taddr) { @@ -142,8 +142,8 @@ __device__ static inline void tcgen05_ld_16x64b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x64b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x64b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_16x64b_pack_16b(_B32 (&__out)[4], ::cuda::std::uint32_t __taddr) { @@ -204,8 +204,8 @@ __device__ static inline void tcgen05_ld_16x64b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x64b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x64b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_16x64b_pack_16b(_B32 (&__out)[8], ::cuda::std::uint32_t __taddr) { @@ -282,8 +282,8 @@ __device__ static inline void tcgen05_ld_16x64b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x64b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x64b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_16x64b_pack_16b(_B32 (&__out)[16], ::cuda::std::uint32_t __taddr) { @@ -386,8 +386,8 @@ __device__ static inline void tcgen05_ld_16x64b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x64b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x64b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_16x64b_pack_16b(_B32 (&__out)[32], ::cuda::std::uint32_t __taddr) { @@ -541,8 +541,8 @@ __device__ static inline void tcgen05_ld_16x64b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x64b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x64b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_16x64b_pack_16b(_B32 (&__out)[64], ::cuda::std::uint32_t __taddr) { @@ -797,8 +797,8 @@ __device__ static inline void tcgen05_ld_16x64b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x64b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x64b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_16x64b_pack_16b(_B32 (&__out)[128], ::cuda::std::uint32_t __taddr) { @@ -986,8 +986,8 @@ __device__ static inline void tcgen05_ld_16x128b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x128b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x128b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_16x128b_pack_16b(_B32 (&__out)[2], ::cuda::std::uint32_t __taddr) { @@ -1041,8 +1041,8 @@ __device__ static inline void tcgen05_ld_16x128b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x128b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x128b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_16x128b_pack_16b(_B32 (&__out)[4], ::cuda::std::uint32_t __taddr) { @@ -1103,8 +1103,8 @@ __device__ static inline void tcgen05_ld_16x128b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x128b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x128b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_16x128b_pack_16b(_B32 (&__out)[8], ::cuda::std::uint32_t __taddr) { @@ -1181,8 +1181,8 @@ __device__ static inline void tcgen05_ld_16x128b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x128b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x128b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_16x128b_pack_16b(_B32 (&__out)[16], ::cuda::std::uint32_t __taddr) { @@ -1285,8 +1285,8 @@ __device__ static inline void tcgen05_ld_16x128b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x128b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x128b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_16x128b_pack_16b(_B32 (&__out)[32], ::cuda::std::uint32_t __taddr) { @@ -1440,8 +1440,8 @@ __device__ static inline void tcgen05_ld_16x128b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x128b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x128b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_16x128b_pack_16b(_B32 (&__out)[64], ::cuda::std::uint32_t __taddr) { @@ -1696,8 +1696,8 @@ __device__ static inline void tcgen05_ld_16x128b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x128b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x128b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_16x128b_pack_16b(_B32 (&__out)[128], ::cuda::std::uint32_t __taddr) { @@ -1885,8 +1885,8 @@ __device__ static inline void tcgen05_ld_16x256b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x256b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x256b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_16x256b_pack_16b(_B32 (&__out)[4], ::cuda::std::uint32_t __taddr) { @@ -1947,8 +1947,8 @@ __device__ static inline void tcgen05_ld_16x256b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x256b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x256b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_16x256b_pack_16b(_B32 (&__out)[8], ::cuda::std::uint32_t __taddr) { @@ -2025,8 +2025,8 @@ __device__ static inline void tcgen05_ld_16x256b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x256b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x256b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_16x256b_pack_16b(_B32 (&__out)[16], ::cuda::std::uint32_t __taddr) { @@ -2129,8 +2129,8 @@ __device__ static inline void tcgen05_ld_16x256b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x256b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x256b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_16x256b_pack_16b(_B32 (&__out)[32], ::cuda::std::uint32_t __taddr) { @@ -2284,8 +2284,8 @@ __device__ static inline void tcgen05_ld_16x256b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x256b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x256b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_16x256b_pack_16b(_B32 (&__out)[64], ::cuda::std::uint32_t __taddr) { @@ -2540,8 +2540,8 @@ __device__ static inline void tcgen05_ld_16x256b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x256b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x256b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_16x256b_pack_16b(_B32 (&__out)[128], ::cuda::std::uint32_t __taddr) { @@ -2726,8 +2726,8 @@ __device__ static inline void tcgen05_ld_32x32b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_32x32b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_32x32b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_32x32b_pack_16b(_B32 (&__out)[1], ::cuda::std::uint32_t __taddr) { @@ -2778,8 +2778,8 @@ __device__ static inline void tcgen05_ld_32x32b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_32x32b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_32x32b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_32x32b_pack_16b(_B32 (&__out)[2], ::cuda::std::uint32_t __taddr) { @@ -2833,8 +2833,8 @@ __device__ static inline void tcgen05_ld_32x32b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_32x32b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_32x32b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_32x32b_pack_16b(_B32 (&__out)[4], ::cuda::std::uint32_t __taddr) { @@ -2895,8 +2895,8 @@ __device__ static inline void tcgen05_ld_32x32b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_32x32b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_32x32b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_32x32b_pack_16b(_B32 (&__out)[8], ::cuda::std::uint32_t __taddr) { @@ -2973,8 +2973,8 @@ __device__ static inline void tcgen05_ld_32x32b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_32x32b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_32x32b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_32x32b_pack_16b(_B32 (&__out)[16], ::cuda::std::uint32_t __taddr) { @@ -3077,8 +3077,8 @@ __device__ static inline void tcgen05_ld_32x32b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_32x32b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_32x32b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_32x32b_pack_16b(_B32 (&__out)[32], ::cuda::std::uint32_t __taddr) { @@ -3232,8 +3232,8 @@ __device__ static inline void tcgen05_ld_32x32b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_32x32b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_32x32b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_32x32b_pack_16b(_B32 (&__out)[64], ::cuda::std::uint32_t __taddr) { @@ -3488,8 +3488,8 @@ __device__ static inline void tcgen05_ld_32x32b_pack_16b( uint32_t taddr); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_32x32b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_32x32b_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_ld_32x32b_pack_16b(_B32 (&__out)[128], ::cuda::std::uint32_t __taddr) { @@ -3680,8 +3680,8 @@ __device__ static inline void tcgen05_ld_16x32bx2_pack_16b( cuda::ptx::n32_t immHalfSplitoff); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x32bx2_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x32bx2_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true, int _N32> _CCCL_DEVICE static inline void tcgen05_ld_16x32bx2_pack_16b( _B32 (&__out)[1], ::cuda::std::uint32_t __taddr, ::cuda::ptx::n32_t<_N32> __immHalfSplitoff) @@ -3739,8 +3739,8 @@ __device__ static inline void tcgen05_ld_16x32bx2_pack_16b( cuda::ptx::n32_t immHalfSplitoff); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x32bx2_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x32bx2_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true, int _N32> _CCCL_DEVICE static inline void tcgen05_ld_16x32bx2_pack_16b( _B32 (&__out)[2], ::cuda::std::uint32_t __taddr, ::cuda::ptx::n32_t<_N32> __immHalfSplitoff) @@ -3798,8 +3798,8 @@ __device__ static inline void tcgen05_ld_16x32bx2_pack_16b( cuda::ptx::n32_t immHalfSplitoff); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x32bx2_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x32bx2_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true, int _N32> _CCCL_DEVICE static inline void tcgen05_ld_16x32bx2_pack_16b( _B32 (&__out)[4], ::cuda::std::uint32_t __taddr, ::cuda::ptx::n32_t<_N32> __immHalfSplitoff) @@ -3864,8 +3864,8 @@ __device__ static inline void tcgen05_ld_16x32bx2_pack_16b( cuda::ptx::n32_t immHalfSplitoff); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x32bx2_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x32bx2_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true, int _N32> _CCCL_DEVICE static inline void tcgen05_ld_16x32bx2_pack_16b( _B32 (&__out)[8], ::cuda::std::uint32_t __taddr, ::cuda::ptx::n32_t<_N32> __immHalfSplitoff) @@ -3946,8 +3946,8 @@ __device__ static inline void tcgen05_ld_16x32bx2_pack_16b( cuda::ptx::n32_t immHalfSplitoff); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x32bx2_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x32bx2_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true, int _N32> _CCCL_DEVICE static inline void tcgen05_ld_16x32bx2_pack_16b( _B32 (&__out)[16], ::cuda::std::uint32_t __taddr, ::cuda::ptx::n32_t<_N32> __immHalfSplitoff) @@ -4054,8 +4054,8 @@ __device__ static inline void tcgen05_ld_16x32bx2_pack_16b( cuda::ptx::n32_t immHalfSplitoff); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x32bx2_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x32bx2_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true, int _N32> _CCCL_DEVICE static inline void tcgen05_ld_16x32bx2_pack_16b( _B32 (&__out)[32], ::cuda::std::uint32_t __taddr, ::cuda::ptx::n32_t<_N32> __immHalfSplitoff) @@ -4213,8 +4213,8 @@ __device__ static inline void tcgen05_ld_16x32bx2_pack_16b( cuda::ptx::n32_t immHalfSplitoff); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x32bx2_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x32bx2_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true, int _N32> _CCCL_DEVICE static inline void tcgen05_ld_16x32bx2_pack_16b( _B32 (&__out)[64], ::cuda::std::uint32_t __taddr, ::cuda::ptx::n32_t<_N32> __immHalfSplitoff) @@ -4473,8 +4473,8 @@ __device__ static inline void tcgen05_ld_16x32bx2_pack_16b( cuda::ptx::n32_t immHalfSplitoff); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_ld_16x32bx2_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_ld_16x32bx2_pack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true, int _N32> _CCCL_DEVICE static inline void tcgen05_ld_16x32bx2_pack_16b( _B32 (&__out)[128], ::cuda::std::uint32_t __taddr, ::cuda::ptx::n32_t<_N32> __immHalfSplitoff) diff --git a/libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_mma.h b/libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_mma.h index 93d0669ed91..3aa131f2174 100644 --- a/libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_mma.h +++ b/libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_mma.h @@ -105,8 +105,8 @@ __device__ static inline void tcgen05_mma( cuda::ptx::n32_t scale_input_d); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_mma_is_only_supported_on_SM_100a_100f_103a_103f_depending_on_the_variant__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_mma_is_only_supported_on_SM_100a_100f_103a_103f_depending_on_the_variant__(); template _CCCL_DEVICE static inline void tcgen05_mma( ::cuda::ptx::kind_t<_Kind> __kind, @@ -457,8 +457,8 @@ __device__ static inline void tcgen05_mma( cuda::ptx::n32_t scale_input_d); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_mma_is_only_supported_on_SM_100a_100f_103a_103f_depending_on_the_variant__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_mma_is_only_supported_on_SM_100a_100f_103a_103f_depending_on_the_variant__(); template _CCCL_DEVICE static inline void tcgen05_mma( ::cuda::ptx::kind_t<_Kind> __kind, diff --git a/libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_st.h b/libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_st.h index 0066aed1efd..55f007d8f8d 100644 --- a/libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_st.h +++ b/libcudacxx/include/cuda/__ptx/instructions/generated/tcgen05_st.h @@ -38,8 +38,8 @@ __device__ static inline void tcgen05_st_16x64b_unpack_16b( const B32 (&values)[1]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x64b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x64b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x64b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[1]) { @@ -95,8 +95,8 @@ __device__ static inline void tcgen05_st_16x64b_unpack_16b( const B32 (&values)[2]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x64b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x64b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x64b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[2]) { @@ -156,8 +156,8 @@ __device__ static inline void tcgen05_st_16x64b_unpack_16b( const B32 (&values)[4]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x64b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x64b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x64b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[4]) { @@ -223,8 +223,8 @@ __device__ static inline void tcgen05_st_16x64b_unpack_16b( const B32 (&values)[8]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x64b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x64b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x64b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[8]) { @@ -303,8 +303,8 @@ __device__ static inline void tcgen05_st_16x64b_unpack_16b( const B32 (&values)[16]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x64b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x64b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x64b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[16]) { @@ -409,8 +409,8 @@ __device__ static inline void tcgen05_st_16x64b_unpack_16b( const B32 (&values)[32]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x64b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x64b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x64b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[32]) { @@ -566,8 +566,8 @@ __device__ static inline void tcgen05_st_16x64b_unpack_16b( const B32 (&values)[64]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x64b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x64b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x64b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[64]) { @@ -824,8 +824,8 @@ __device__ static inline void tcgen05_st_16x64b_unpack_16b( const B32 (&values)[128]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x64b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x64b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x64b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[128]) { @@ -1016,8 +1016,8 @@ __device__ static inline void tcgen05_st_16x128b_unpack_16b( const B32 (&values)[2]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x128b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x128b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x128b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[2]) { @@ -1077,8 +1077,8 @@ __device__ static inline void tcgen05_st_16x128b_unpack_16b( const B32 (&values)[4]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x128b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x128b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x128b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[4]) { @@ -1144,8 +1144,8 @@ __device__ static inline void tcgen05_st_16x128b_unpack_16b( const B32 (&values)[8]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x128b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x128b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x128b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[8]) { @@ -1224,8 +1224,8 @@ __device__ static inline void tcgen05_st_16x128b_unpack_16b( const B32 (&values)[16]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x128b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x128b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x128b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[16]) { @@ -1330,8 +1330,8 @@ __device__ static inline void tcgen05_st_16x128b_unpack_16b( const B32 (&values)[32]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x128b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x128b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x128b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[32]) { @@ -1487,8 +1487,8 @@ __device__ static inline void tcgen05_st_16x128b_unpack_16b( const B32 (&values)[64]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x128b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x128b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x128b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[64]) { @@ -1745,8 +1745,8 @@ __device__ static inline void tcgen05_st_16x128b_unpack_16b( const B32 (&values)[128]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x128b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x128b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x128b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[128]) { @@ -1939,8 +1939,8 @@ __device__ static inline void tcgen05_st_16x256b_unpack_16b( const B32 (&values)[4]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x256b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x256b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x256b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[4]) { @@ -2006,8 +2006,8 @@ __device__ static inline void tcgen05_st_16x256b_unpack_16b( const B32 (&values)[8]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x256b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x256b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x256b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[8]) { @@ -2086,8 +2086,8 @@ __device__ static inline void tcgen05_st_16x256b_unpack_16b( const B32 (&values)[16]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x256b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x256b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x256b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[16]) { @@ -2192,8 +2192,8 @@ __device__ static inline void tcgen05_st_16x256b_unpack_16b( const B32 (&values)[32]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x256b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x256b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x256b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[32]) { @@ -2349,8 +2349,8 @@ __device__ static inline void tcgen05_st_16x256b_unpack_16b( const B32 (&values)[64]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x256b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x256b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x256b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[64]) { @@ -2607,8 +2607,8 @@ __device__ static inline void tcgen05_st_16x256b_unpack_16b( const B32 (&values)[128]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x256b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x256b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x256b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[128]) { @@ -2797,8 +2797,8 @@ __device__ static inline void tcgen05_st_32x32b_unpack_16b( const B32 (&values)[1]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_32x32b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_32x32b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_32x32b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[1]) { @@ -2854,8 +2854,8 @@ __device__ static inline void tcgen05_st_32x32b_unpack_16b( const B32 (&values)[2]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_32x32b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_32x32b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_32x32b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[2]) { @@ -2915,8 +2915,8 @@ __device__ static inline void tcgen05_st_32x32b_unpack_16b( const B32 (&values)[4]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_32x32b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_32x32b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_32x32b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[4]) { @@ -2982,8 +2982,8 @@ __device__ static inline void tcgen05_st_32x32b_unpack_16b( const B32 (&values)[8]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_32x32b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_32x32b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_32x32b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[8]) { @@ -3062,8 +3062,8 @@ __device__ static inline void tcgen05_st_32x32b_unpack_16b( const B32 (&values)[16]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_32x32b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_32x32b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_32x32b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[16]) { @@ -3168,8 +3168,8 @@ __device__ static inline void tcgen05_st_32x32b_unpack_16b( const B32 (&values)[32]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_32x32b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_32x32b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_32x32b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[32]) { @@ -3325,8 +3325,8 @@ __device__ static inline void tcgen05_st_32x32b_unpack_16b( const B32 (&values)[64]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_32x32b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_32x32b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_32x32b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[64]) { @@ -3583,8 +3583,8 @@ __device__ static inline void tcgen05_st_32x32b_unpack_16b( const B32 (&values)[128]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_32x32b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_32x32b_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_32x32b_unpack_16b(::cuda::std::uint32_t __taddr, const _B32 (&__values)[128]) { @@ -3778,8 +3778,8 @@ __device__ static inline void tcgen05_st_16x32bx2_unpack_16b( const B32 (&values)[1]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x32bx2_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x32bx2_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x32bx2_unpack_16b( ::cuda::std::uint32_t __taddr, ::cuda::ptx::n32_t<_N32> __immHalfSplitoff, const _B32 (&__values)[1]) @@ -3842,8 +3842,8 @@ __device__ static inline void tcgen05_st_16x32bx2_unpack_16b( const B32 (&values)[2]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x32bx2_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x32bx2_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x32bx2_unpack_16b( ::cuda::std::uint32_t __taddr, ::cuda::ptx::n32_t<_N32> __immHalfSplitoff, const _B32 (&__values)[2]) @@ -3909,8 +3909,8 @@ __device__ static inline void tcgen05_st_16x32bx2_unpack_16b( const B32 (&values)[4]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x32bx2_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x32bx2_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x32bx2_unpack_16b( ::cuda::std::uint32_t __taddr, ::cuda::ptx::n32_t<_N32> __immHalfSplitoff, const _B32 (&__values)[4]) @@ -3982,8 +3982,8 @@ __device__ static inline void tcgen05_st_16x32bx2_unpack_16b( const B32 (&values)[8]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x32bx2_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x32bx2_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x32bx2_unpack_16b( ::cuda::std::uint32_t __taddr, ::cuda::ptx::n32_t<_N32> __immHalfSplitoff, const _B32 (&__values)[8]) @@ -4068,8 +4068,8 @@ __device__ static inline void tcgen05_st_16x32bx2_unpack_16b( const B32 (&values)[16]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x32bx2_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x32bx2_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x32bx2_unpack_16b( ::cuda::std::uint32_t __taddr, ::cuda::ptx::n32_t<_N32> __immHalfSplitoff, const _B32 (&__values)[16]) @@ -4180,8 +4180,8 @@ __device__ static inline void tcgen05_st_16x32bx2_unpack_16b( const B32 (&values)[32]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x32bx2_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x32bx2_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x32bx2_unpack_16b( ::cuda::std::uint32_t __taddr, ::cuda::ptx::n32_t<_N32> __immHalfSplitoff, const _B32 (&__values)[32]) @@ -4343,8 +4343,8 @@ __device__ static inline void tcgen05_st_16x32bx2_unpack_16b( const B32 (&values)[64]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x32bx2_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x32bx2_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x32bx2_unpack_16b( ::cuda::std::uint32_t __taddr, ::cuda::ptx::n32_t<_N32> __immHalfSplitoff, const _B32 (&__values)[64]) @@ -4607,8 +4607,8 @@ __device__ static inline void tcgen05_st_16x32bx2_unpack_16b( const B32 (&values)[128]); */ #if __cccl_ptx_isa >= 860 -extern "C" _CCCL_DEVICE void -__cuda_ptx_tcgen05_st_16x32bx2_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); +extern "C" + _CCCL_DEVICE void __cuda_ptx_tcgen05_st_16x32bx2_unpack_16b_is_only_supported_on_SM_100a_100f_103a_103f_110a_110f__(); template = true> _CCCL_DEVICE static inline void tcgen05_st_16x32bx2_unpack_16b( ::cuda::std::uint32_t __taddr, ::cuda::ptx::n32_t<_N32> __immHalfSplitoff, const _B32 (&__values)[128]) diff --git a/libcudacxx/include/cuda/__utility/__basic_any/rtti.h b/libcudacxx/include/cuda/__utility/__basic_any/rtti.h index 67e30ad53df..ad3963653ef 100644 --- a/libcudacxx/include/cuda/__utility/__basic_any/rtti.h +++ b/libcudacxx/include/cuda/__utility/__basic_any/rtti.h @@ -122,7 +122,7 @@ inline constexpr size_t __half_size_t_bits = sizeof(size_t) * CHAR_BIT / 2; // which contains a sub-object of this type. struct __object_metadata { - size_t __size_ : __half_size_t_bits; + size_t __size_ : __half_size_t_bits; size_t __align_ : __half_size_t_bits; ::cuda::std::__type_info_ptr __object_typeid_; ::cuda::std::__type_info_ptr __pointer_typeid_; diff --git a/libcudacxx/include/cuda/std/__algorithm/iterator_operations.h b/libcudacxx/include/cuda/std/__algorithm/iterator_operations.h index 829a95b6272..6ccba3c81de 100644 --- a/libcudacxx/include/cuda/std/__algorithm/iterator_operations.h +++ b/libcudacxx/include/cuda/std/__algorithm/iterator_operations.h @@ -115,8 +115,7 @@ struct _IterOps<_ClassicAlgPolicy> _CCCL_API constexpr static // If the result of dereferencing `_Iter` is a reference type, deduce the result of calling `::cuda::std::move` on // it. Note that the C++03 mode doesn't support `decltype(auto)` as the return type. - __move_t<_Iter> - __iter_move(_Iter&& __i) + __move_t<_Iter> __iter_move(_Iter&& __i) { __validate_iter_reference<_Iter>(); @@ -129,8 +128,7 @@ struct _IterOps<_ClassicAlgPolicy> // If the result of dereferencing `_Iter` is a value type, deduce the return value of this function to also be a // value -- otherwise, after `operator*` returns a temporary, this function would return a dangling reference to // that temporary. Note that the C++03 mode doesn't support `auto` as the return type. - __deref_t<_Iter> - __iter_move(_Iter&& __i) + __deref_t<_Iter> __iter_move(_Iter&& __i) { __validate_iter_reference<_Iter>(); diff --git a/libcudacxx/include/cuda/std/__algorithm/sort.h b/libcudacxx/include/cuda/std/__algorithm/sort.h index 644c4c901a5..d4e3446bc1a 100644 --- a/libcudacxx/include/cuda/std/__algorithm/sort.h +++ b/libcudacxx/include/cuda/std/__algorithm/sort.h @@ -927,9 +927,9 @@ _CCCL_API void __introsort( continue; } // Use bitset partition only if asked for. - auto __ret = _UseBitSetPartition - ? ::cuda::std::__bitset_partition<_AlgPolicy, _RandomAccessIterator, _Compare>(__first, __last, __comp) - : ::cuda::std::__partition_with_equals_on_right<_AlgPolicy, _RandomAccessIterator, _Compare>( + auto __ret = _UseBitSetPartition + ? ::cuda::std::__bitset_partition<_AlgPolicy, _RandomAccessIterator, _Compare>(__first, __last, __comp) + : ::cuda::std::__partition_with_equals_on_right<_AlgPolicy, _RandomAccessIterator, _Compare>( __first, __last, __comp); _RandomAccessIterator __i = __ret.first; // [__first, __i) < *__i and *__i <= [__i+1, __last) diff --git a/libcudacxx/include/cuda/std/__cccl/compiler.h b/libcudacxx/include/cuda/std/__cccl/compiler.h index da41e2f18b0..a2cb1f6d384 100644 --- a/libcudacxx/include/cuda/std/__cccl/compiler.h +++ b/libcudacxx/include/cuda/std/__cccl/compiler.h @@ -27,11 +27,11 @@ // #define MYPRODUCT_<_PROD>() (1, 2) // or // #define MYPRODUCT_<_PROD>() _CCCL_VERSION_INVALID() -#define _CCCL_VERSION_MAJOR_(_MAJOR, _MINOR) _MAJOR -#define _CCCL_VERSION_MAJOR(_PAIR) _CCCL_VERSION_MAJOR_ _PAIR -#define _CCCL_VERSION_INVALID() (-1, -1) -#define _CCCL_MAKE_VERSION(_PREFIX, _PAIR) (_CCCL_PP_EVAL(_CCCL_PP_CAT(_PREFIX, MAKE_VERSION), _CCCL_PP_EXPAND _PAIR)) -#define _CCCL_VERSION_IS_INVALID(_PAIR) (_CCCL_VERSION_MAJOR(_PAIR) == _CCCL_VERSION_MAJOR(_CCCL_VERSION_INVALID())) +#define _CCCL_VERSION_MAJOR_(_MAJOR, _MINOR) _MAJOR +#define _CCCL_VERSION_MAJOR(_PAIR) _CCCL_VERSION_MAJOR_ _PAIR +#define _CCCL_VERSION_INVALID() (-1, -1) +#define _CCCL_MAKE_VERSION(_PREFIX, _PAIR) (_CCCL_PP_EVAL(_CCCL_PP_CAT(_PREFIX, MAKE_VERSION), _CCCL_PP_EXPAND _PAIR)) +#define _CCCL_VERSION_IS_INVALID(_PAIR) (_CCCL_VERSION_MAJOR(_PAIR) == _CCCL_VERSION_MAJOR(_CCCL_VERSION_INVALID())) #define _CCCL_VERSION_COMPARE_1(_PREFIX, _VER) (!_CCCL_VERSION_IS_INVALID(_VER())) #define _CCCL_VERSION_COMPARE_3(_PREFIX, _VER, _OP, _MAJOR) \ (!_CCCL_VERSION_IS_INVALID(_VER()) && (_CCCL_VERSION_MAJOR(_VER()) _OP _MAJOR)) @@ -113,7 +113,7 @@ // The CUDA compiler version shares the implementation with the C++ compiler #define _CCCL_CUDA_COMPILER_MAKE_VERSION(_MAJOR, _MINOR) _CCCL_COMPILER_MAKE_VERSION(_MAJOR, _MINOR) -#define _CCCL_CUDA_COMPILER(...) _CCCL_VERSION_COMPARE(_CCCL_CUDA_COMPILER_, _CCCL_CUDA_COMPILER_##__VA_ARGS__) +#define _CCCL_CUDA_COMPILER(...) _CCCL_VERSION_COMPARE(_CCCL_CUDA_COMPILER_, _CCCL_CUDA_COMPILER_##__VA_ARGS__) #define _CCCL_CUDA_COMPILER_NVCC() _CCCL_VERSION_INVALID() #define _CCCL_CUDA_COMPILER_NVHPC() _CCCL_VERSION_INVALID() diff --git a/libcudacxx/include/cuda/std/__cccl/host_std_lib.h b/libcudacxx/include/cuda/std/__cccl/host_std_lib.h index 634100ea923..2ff52b15099 100644 --- a/libcudacxx/include/cuda/std/__cccl/host_std_lib.h +++ b/libcudacxx/include/cuda/std/__cccl/host_std_lib.h @@ -35,7 +35,7 @@ #endif // ^^^ __has_include() ^^^ #define _CCCL_HOST_STD_LIB_MAKE_VERSION(_MAJOR, _MINOR) ((_MAJOR) * 100 + (_MINOR)) -#define _CCCL_HOST_STD_LIB(...) _CCCL_VERSION_COMPARE(_CCCL_HOST_STD_LIB_, _CCCL_HOST_STD_LIB_##__VA_ARGS__) +#define _CCCL_HOST_STD_LIB(...) _CCCL_VERSION_COMPARE(_CCCL_HOST_STD_LIB_, _CCCL_HOST_STD_LIB_##__VA_ARGS__) #if _CCCL_HOSTED() # if defined(_MSVC_STL_VERSION) diff --git a/libcudacxx/include/cuda/std/__cccl/preprocessor.h b/libcudacxx/include/cuda/std/__cccl/preprocessor.h index a8f6cbbfd96..b5486d1d594 100644 --- a/libcudacxx/include/cuda/std/__cccl/preprocessor.h +++ b/libcudacxx/include/cuda/std/__cccl/preprocessor.h @@ -92,13 +92,13 @@ CCCL_IGNORE_MSVC_TRADITIONAL_PREPROCESSOR_WARNING to suppress this warning. #define _CCCP_PP_CASE_LABEL_(_PREFIX, ...) \ _CCCL_PP_EVAL(_CCCL_PP_CHECK, _CCCL_PP_CAT(_PREFIX##_SWITCH_, _CCCL_PP_FIRST(__VA_ARGS__)), _CCCL_SWITCH_DEFAULT, ) -#define _CCCL_PP_FOR_EACH(_Mp, ...) _CCCL_PP_FOR_EACH_N(_CCCL_PP_COUNT(__VA_ARGS__), _Mp, __VA_ARGS__) -#define _CCCL_PP_FOR_EACH_N(_Np, _Mp, ...) _CCCL_PP_CAT2(_CCCL_PP_FOR_EACH_, _Np)(_Mp, __VA_ARGS__) -#define _CCCL_PP_FOR_EACH_1(_Mp, _1) _Mp(_1) -#define _CCCL_PP_FOR_EACH_2(_Mp, _1, _2) _Mp(_1) _Mp(_2) -#define _CCCL_PP_FOR_EACH_3(_Mp, _1, _2, _3) _Mp(_1) _Mp(_2) _Mp(_3) -#define _CCCL_PP_FOR_EACH_4(_Mp, _1, _2, _3, _4) _Mp(_1) _Mp(_2) _Mp(_3) _Mp(_4) -#define _CCCL_PP_FOR_EACH_5(_Mp, _1, _2, _3, _4, _5) _Mp(_1) _Mp(_2) _Mp(_3) _Mp(_4) _Mp(_5) +#define _CCCL_PP_FOR_EACH(_Mp, ...) _CCCL_PP_FOR_EACH_N(_CCCL_PP_COUNT(__VA_ARGS__), _Mp, __VA_ARGS__) +#define _CCCL_PP_FOR_EACH_N(_Np, _Mp, ...) _CCCL_PP_CAT2(_CCCL_PP_FOR_EACH_, _Np)(_Mp, __VA_ARGS__) +#define _CCCL_PP_FOR_EACH_1(_Mp, _1) _Mp(_1) +#define _CCCL_PP_FOR_EACH_2(_Mp, _1, _2) _Mp(_1) _Mp(_2) +#define _CCCL_PP_FOR_EACH_3(_Mp, _1, _2, _3) _Mp(_1) _Mp(_2) _Mp(_3) +#define _CCCL_PP_FOR_EACH_4(_Mp, _1, _2, _3, _4) _Mp(_1) _Mp(_2) _Mp(_3) _Mp(_4) +#define _CCCL_PP_FOR_EACH_5(_Mp, _1, _2, _3, _4, _5) _Mp(_1) _Mp(_2) _Mp(_3) _Mp(_4) _Mp(_5) #define _CCCL_PP_FOR_EACH_6(_Mp, _1, _2, _3, _4, _5, _6) _Mp(_1) _Mp(_2) _Mp(_3) _Mp(_4) _Mp(_5) _Mp(_6) #define _CCCL_PP_FOR_EACH_7(_Mp, _1, _2, _3, _4, _5, _6, _7) _Mp(_1) _Mp(_2) _Mp(_3) _Mp(_4) _Mp(_5) _Mp(_6) _Mp(_7) #define _CCCL_PP_FOR_EACH_8(_Mp, _1, _2, _3, _4, _5, _6, _7, _8) \ @@ -976,7 +976,7 @@ CCCL_IGNORE_MSVC_TRADITIONAL_PREPROCESSOR_WARNING to suppress this warning. #define _CCCL_PP_REPEAT_REVERSE_AUX3(_N, _M, _S, _F) _CCCL_PP_CAT(_CCCL_PP_REPEAT_REVERSE, _N)(_M, _S, _F) #define _CCCL_PP_REPEAT_REVERSE_AUX(_C, _N, ...) _CCCL_PP_CAT(_CCCL_PP_REPEAT_REVERSE_AUX, _C)(_N, __VA_ARGS__) -#define _CCCL_PP_REPEAT_REVERSE(_N, ...) _CCCL_PP_REPEAT_REVERSE_AUX(_CCCL_PP_COUNT(__VA_ARGS__), _N, __VA_ARGS__) +#define _CCCL_PP_REPEAT_REVERSE(_N, ...) _CCCL_PP_REPEAT_REVERSE_AUX(_CCCL_PP_COUNT(__VA_ARGS__), _N, __VA_ARGS__) #define _CCCL_PP_REPEAT_REVERSE0(_M, _S, _F) #define _CCCL_PP_REPEAT_REVERSE1(_M, _S, _F) _M(_S) diff --git a/libcudacxx/include/cuda/std/__cccl/visibility.h b/libcudacxx/include/cuda/std/__cccl/visibility.h index 3cfba4912a5..b00cdb4f074 100644 --- a/libcudacxx/include/cuda/std/__cccl/visibility.h +++ b/libcudacxx/include/cuda/std/__cccl/visibility.h @@ -134,7 +134,7 @@ # define _CCCL_DEVICE_API _CCCL_DEVICE # define _CCCL_TILE_API _CCCL_TILE #else // ^^^ _CCCL_COMPILER(NVHPC) ^^^ / vvv !_CCCL_COMPILER(NVHPC) vvv -# define _CCCL_API _CCCL_TILE _CCCL_HOST_DEVICE _CCCL_VISIBILITY_HIDDEN _CCCL_EXCLUDE_FROM_EXPLICIT_INSTANTIATION +# define _CCCL_API _CCCL_TILE _CCCL_HOST_DEVICE _CCCL_VISIBILITY_HIDDEN _CCCL_EXCLUDE_FROM_EXPLICIT_INSTANTIATION # define _CCCL_HOST_DEVICE_API _CCCL_HOST_DEVICE _CCCL_VISIBILITY_HIDDEN _CCCL_EXCLUDE_FROM_EXPLICIT_INSTANTIATION # define _CCCL_HOST_API _CCCL_HOST _CCCL_VISIBILITY_HIDDEN _CCCL_EXCLUDE_FROM_EXPLICIT_INSTANTIATION # define _CCCL_DEVICE_API _CCCL_DEVICE _CCCL_VISIBILITY_HIDDEN _CCCL_EXCLUDE_FROM_EXPLICIT_INSTANTIATION diff --git a/libcudacxx/include/cuda/std/__chrono/duration.h b/libcudacxx/include/cuda/std/__chrono/duration.h index 8ab716e68e3..896ed807669 100644 --- a/libcudacxx/include/cuda/std/__chrono/duration.h +++ b/libcudacxx/include/cuda/std/__chrono/duration.h @@ -339,7 +339,8 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT duration #if _CCCL_STD_VER <= 2017 template - [[nodiscard]] _CCCL_API friend constexpr bool operator!=(const duration& __lhs, const duration<_Rep2, _Period2>& __rhs) + [[nodiscard]] + _CCCL_API friend constexpr bool operator!=(const duration& __lhs, const duration<_Rep2, _Period2>& __rhs) { if constexpr (is_same_v>) { diff --git a/libcudacxx/include/cuda/std/__chrono/time_point.h b/libcudacxx/include/cuda/std/__chrono/time_point.h index fac774cd2b4..5bc855f1f9e 100644 --- a/libcudacxx/include/cuda/std/__chrono/time_point.h +++ b/libcudacxx/include/cuda/std/__chrono/time_point.h @@ -164,8 +164,8 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT time_point #if _CCCL_STD_VER <= 2017 template - [[nodiscard]] _CCCL_API friend constexpr bool - operator!=(const time_point& __lhs, const time_point<_Clock, _Duration2>& __rhs) + [[nodiscard]] + _CCCL_API friend constexpr bool operator!=(const time_point& __lhs, const time_point<_Clock, _Duration2>& __rhs) { return __lhs.time_since_epoch() != __rhs.time_since_epoch(); } @@ -181,8 +181,8 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT time_point } #else // ^^^ _LIBCUDACXX_HAS_SPACESHIP_OPERATOR() ^^^ / vvv !_LIBCUDACXX_HAS_SPACESHIP_OPERATOR() vvv template - [[nodiscard]] _CCCL_API friend constexpr bool - operator<(const time_point& __lhs, const time_point<_Clock, _Duration2>& __rhs) + [[nodiscard]] + _CCCL_API friend constexpr bool operator<(const time_point& __lhs, const time_point<_Clock, _Duration2>& __rhs) { return __lhs.time_since_epoch() < __rhs.time_since_epoch(); } diff --git a/libcudacxx/include/cuda/std/__complex/complex.h b/libcudacxx/include/cuda/std/__complex/complex.h index ee959582441..30ad936b7a0 100644 --- a/libcudacxx/include/cuda/std/__complex/complex.h +++ b/libcudacxx/include/cuda/std/__complex/complex.h @@ -299,11 +299,11 @@ template bool __z_inf = ::cuda::std::isinf(__a) || ::cuda::std::isinf(__b); bool __w_inf = ::cuda::std::isinf(__c) || ::cuda::std::isinf(__d); bool __z_nan = !__z_inf - && ((::cuda::std::isnan(__a) && ::cuda::std::isnan(__b)) || (::cuda::std::isnan(__a) && __b == _Tp(0)) - || (__a == _Tp(0) && ::cuda::std::isnan(__b))); - bool __w_nan = !__w_inf - && ((::cuda::std::isnan(__c) && ::cuda::std::isnan(__d)) || (::cuda::std::isnan(__c) && __d == _Tp(0)) - || (__c == _Tp(0) && ::cuda::std::isnan(__d))); + && ((::cuda::std::isnan(__a) && ::cuda::std::isnan(__b)) || (::cuda::std::isnan(__a) && __b == _Tp(0)) + || (__a == _Tp(0) && ::cuda::std::isnan(__b))); + bool __w_nan = !__w_inf + && ((::cuda::std::isnan(__c) && ::cuda::std::isnan(__d)) || (::cuda::std::isnan(__c) && __d == _Tp(0)) + || (__c == _Tp(0) && ::cuda::std::isnan(__d))); if (__z_nan || __w_nan) { return complex<_Tp>(numeric_limits<_Tp>::quiet_NaN(), _Tp(0)); @@ -437,11 +437,11 @@ operator/(const complex<_Tp>& __z, const complex<_Tp>& __w) bool __z_inf = ::cuda::std::isinf(__a) || ::cuda::std::isinf(__b); bool __w_inf = ::cuda::std::isinf(__c) || ::cuda::std::isinf(__d); bool __z_nan = !__z_inf - && ((::cuda::std::isnan(__a) && ::cuda::std::isnan(__b)) || (::cuda::std::isnan(__a) && __b == _Tp(0)) - || (__a == _Tp(0) && ::cuda::std::isnan(__b))); - bool __w_nan = !__w_inf - && ((::cuda::std::isnan(__c) && ::cuda::std::isnan(__d)) || (::cuda::std::isnan(__c) && __d == _Tp(0)) - || (__c == _Tp(0) && ::cuda::std::isnan(__d))); + && ((::cuda::std::isnan(__a) && ::cuda::std::isnan(__b)) || (::cuda::std::isnan(__a) && __b == _Tp(0)) + || (__a == _Tp(0) && ::cuda::std::isnan(__b))); + bool __w_nan = !__w_inf + && ((::cuda::std::isnan(__c) && ::cuda::std::isnan(__d)) || (::cuda::std::isnan(__c) && __d == _Tp(0)) + || (__c == _Tp(0) && ::cuda::std::isnan(__d))); if ((__z_nan || __w_nan) || (__z_inf && __w_inf)) { return complex<_Tp>(numeric_limits<_Tp>::quiet_NaN(), _Tp(0)); diff --git a/libcudacxx/include/cuda/std/__concepts/concept_macros.h b/libcudacxx/include/cuda/std/__concepts/concept_macros.h index c821e022b1c..2f8bc6ddebe 100644 --- a/libcudacxx/include/cuda/std/__concepts/concept_macros.h +++ b/libcudacxx/include/cuda/std/__concepts/concept_macros.h @@ -60,9 +60,9 @@ using __cccl_requires_t = typename __cccl_select<_Bp>::template type<_Tp>; # define _CCCL_TRAILING_REQUIRES(...) ->__VA_ARGS__ _CCCL_TRAILING_REQUIRES_IMPL_ # define _CCCL_CONCEPT concept #else // ^^^ _CCCL_HAS_CONCEPTS() ^^^ / vvv !_CCCL_HAS_CONCEPTS() vvv -# define _CCCL_TEMPLATE(...) template <__VA_ARGS__ -# define _CCCL_REQUIRES(...) , bool __cccl_true_ = true, __cccl_enable_if_t < __VA_ARGS__ && __cccl_true_, int > = 0 > -# define _CCCL_AND &&__cccl_true_, int > = 0, __cccl_enable_if_t < +# define _CCCL_TEMPLATE(...) template <__VA_ARGS__ +# define _CCCL_REQUIRES(...) , bool __cccl_true_ = true, __cccl_enable_if_t < __VA_ARGS__ && __cccl_true_, int > = 0 > +# define _CCCL_AND &&__cccl_true_, int > = 0, __cccl_enable_if_t < # define _CCCL_TRAILING_REQUIRES(...) ->__cccl_requires_t < __VA_ARGS__ _CCCL_TRAILING_REQUIRES_IMPL_ # define _CCCL_TRAILING_REQUIRES_IMPL_(...) , __VA_ARGS__ > # define _CCCL_CONCEPT inline constexpr bool @@ -168,8 +168,8 @@ namespace __cccl_unqualified_cuda_std = ::cuda::std; // NOLINT(misc-unused-alias #define _CCCL_CONCEPT_EAT_TYPENAME_typename(...) __VA_ARGS__ // Converts "[typename]opt TYPE..." to "typename TYPE..." -#define _CCCL_CONCEPT_TRY_ADD_TYPENAME_(...) _CCCL_PP_SWITCH2(_CCCL_CONCEPT_TRY_ADD_TYPENAME, __VA_ARGS__) -#define _CCCL_CONCEPT_TRY_ADD_TYPENAME_SWITCH_typename _CCCL_PP_CASE(_CCCL_SWITCH_TYPENAME) +#define _CCCL_CONCEPT_TRY_ADD_TYPENAME_(...) _CCCL_PP_SWITCH2(_CCCL_CONCEPT_TRY_ADD_TYPENAME, __VA_ARGS__) +#define _CCCL_CONCEPT_TRY_ADD_TYPENAME_SWITCH_typename _CCCL_PP_CASE(_CCCL_SWITCH_TYPENAME) #define _CCCL_CONCEPT_TRY_ADD_TYPENAME_CASE__CCCL_SWITCH_DEFAULT(...) typename __VA_ARGS__ #define _CCCL_CONCEPT_TRY_ADD_TYPENAME_CASE__CCCL_SWITCH_TYPENAME(...) __VA_ARGS__ @@ -202,7 +202,7 @@ namespace __cccl_unqualified_cuda_std = ::cuda::std; // NOLINT(misc-unused-alias // "_CCCL_CONCEPT_FRAGMENT(NAME, ARGS...)(REQS...)" expands into // "concept NAME = requires(ARGS...) { _CCCL_CONCEPT_REQUIREMENT_(REQS)... }" -# define _CCCL_CONCEPT_FRAGMENT(_NAME, ...) concept _NAME = _CCCL_CONCEPT_FRAGMENT_REQUIREMENTS_##__VA_ARGS__ +# define _CCCL_CONCEPT_FRAGMENT(_NAME, ...) concept _NAME = _CCCL_CONCEPT_FRAGMENT_REQUIREMENTS_##__VA_ARGS__ # define _CCCL_CONCEPT_FRAGMENT_REQUIREMENTS_requires(...) requires(__VA_ARGS__) _CCCL_CONCEPT_FRAGMENT_REQUIREMENTS_ # define _CCCL_CONCEPT_FRAGMENT_REQUIREMENTS_(...) {_CCCL_PP_FOR_EACH(_CCCL_CONCEPT_REQUIREMENT_, __VA_ARGS__)} @@ -322,21 +322,21 @@ namespace __cccl_unqualified_cuda_std = ::cuda::std; // NOLINT(misc-unused-alias _CCCL_API inline static auto __cccl_well_formed(__VA_ARGS__) _CCCL_REQUIRES_EXPR_REQUIREMENTS_ // Expands "T1, T2, variadic T3" to ", class T1, class T2, class... T3" -# define _CCCL_REQUIRES_EXPR_TPARAM_DEFNS(...) _CCCL_PP_FOR_EACH(_CCCL_REQUIRES_EXPR_TPARAM_DEFN, __VA_ARGS__) +# define _CCCL_REQUIRES_EXPR_TPARAM_DEFNS(...) _CCCL_PP_FOR_EACH(_CCCL_REQUIRES_EXPR_TPARAM_DEFN, __VA_ARGS__) // Expands "TY" to ", class TY" and "variadic TY" to ", class... TY" -# define _CCCL_REQUIRES_EXPR_TPARAM_DEFN(_TY) , _CCCL_PP_SWITCH2(_CCCL_REQUIRES_EXPR_TPARAM_DEFN, _TY) +# define _CCCL_REQUIRES_EXPR_TPARAM_DEFN(_TY) , _CCCL_PP_SWITCH2(_CCCL_REQUIRES_EXPR_TPARAM_DEFN, _TY) # define _CCCL_REQUIRES_EXPR_TPARAM_DEFN_SWITCH_variadic _CCCL_PP_CASE(_CCCL_SWITCH_VARIADIC) # define _CCCL_REQUIRES_EXPR_TPARAM_DEFN_CASE__CCCL_SWITCH_DEFAULT(_TY) class _TY # define _CCCL_REQUIRES_EXPR_TPARAM_DEFN_CASE__CCCL_SWITCH_VARIADIC(_TY) \ class... _CCCL_PP_CAT(_CCCL_REQUIRES_EXPR_EAT_VARIADIC_, _TY) // Expands "T1, T2, variadic T3" to ", T1, T2, T3..." -# define _CCCL_REQUIRES_EXPR_TPARAM_REFS(...) _CCCL_PP_FOR_EACH(_CCCL_REQUIRES_EXPR_TPARAM_REF, __VA_ARGS__) +# define _CCCL_REQUIRES_EXPR_TPARAM_REFS(...) _CCCL_PP_FOR_EACH(_CCCL_REQUIRES_EXPR_TPARAM_REF, __VA_ARGS__) // Expands "TY" to ", TY" and "variadic TY" to ", TY..." -# define _CCCL_REQUIRES_EXPR_TPARAM_REF(_TY) , _CCCL_PP_SWITCH2(_CCCL_REQUIRES_EXPR_TPARAM_REF, _TY) -# define _CCCL_REQUIRES_EXPR_TPARAM_REF_SWITCH_variadic _CCCL_PP_CASE(_CCCL_SWITCH_VARIADIC) +# define _CCCL_REQUIRES_EXPR_TPARAM_REF(_TY) , _CCCL_PP_SWITCH2(_CCCL_REQUIRES_EXPR_TPARAM_REF, _TY) +# define _CCCL_REQUIRES_EXPR_TPARAM_REF_SWITCH_variadic _CCCL_PP_CASE(_CCCL_SWITCH_VARIADIC) # define _CCCL_REQUIRES_EXPR_TPARAM_REF_CASE__CCCL_SWITCH_DEFAULT(_TY) _TY # define _CCCL_REQUIRES_EXPR_TPARAM_REF_CASE__CCCL_SWITCH_VARIADIC(_TY) \ _CCCL_PP_CAT(_CCCL_REQUIRES_EXPR_EAT_VARIADIC_, _TY)... @@ -358,8 +358,8 @@ namespace __cccl_unqualified_cuda_std = ::cuda::std; // NOLINT(misc-unused-alias _CCCL_PP_EVAL(_CCCL_REQUIRES_EXPR_ID_CONCAT_ALL_IMPL, __VA_ARGS__, , , , , , , , , ) // Expands "TY" to "TY" and "variadic TY" to "TY" -# define _CCCL_REQUIRES_EXPR_ID_IMPL(_TY) , _CCCL_PP_SWITCH2(_CCCL_REQUIRES_EXPR_ID_IMPL, _TY) -# define _CCCL_REQUIRES_EXPR_ID_IMPL_SWITCH_variadic _CCCL_PP_CASE(_CCCL_SWITCH_VARIADIC) +# define _CCCL_REQUIRES_EXPR_ID_IMPL(_TY) , _CCCL_PP_SWITCH2(_CCCL_REQUIRES_EXPR_ID_IMPL, _TY) +# define _CCCL_REQUIRES_EXPR_ID_IMPL_SWITCH_variadic _CCCL_PP_CASE(_CCCL_SWITCH_VARIADIC) # define _CCCL_REQUIRES_EXPR_ID_IMPL_CASE__CCCL_SWITCH_DEFAULT(_TY) _TY # define _CCCL_REQUIRES_EXPR_ID_IMPL_CASE__CCCL_SWITCH_VARIADIC(_TY) \ _CCCL_PP_CAT(_CCCL_REQUIRES_EXPR_EAT_VARIADIC_, _TY) diff --git a/libcudacxx/include/cuda/std/__format/buffer.h b/libcudacxx/include/cuda/std/__format/buffer.h index 11732bed8cb..09e3a862c08 100644 --- a/libcudacxx/include/cuda/std/__format/buffer.h +++ b/libcudacxx/include/cuda/std/__format/buffer.h @@ -458,10 +458,8 @@ class __fmt_direct_iterator_buffer : public __fmt_output_buffer<_CharT> public: _CCCL_API constexpr explicit __fmt_direct_iterator_buffer( _OutIt __out_it, __fmt_max_output_size* __max_output_size = nullptr) noexcept - : __fmt_output_buffer<_CharT>{::cuda::std::__unwrap_iter(__out_it), - __buffer_size, - __prepare_write, - __max_output_size} + : __fmt_output_buffer<_CharT>{ + ::cuda::std::__unwrap_iter(__out_it), __buffer_size, __prepare_write, __max_output_size} , __out_it_{__out_it} {} diff --git a/libcudacxx/include/cuda/std/__iterator/concepts.h b/libcudacxx/include/cuda/std/__iterator/concepts.h index a346d1cbf9c..a20ced92fbc 100644 --- a/libcudacxx/include/cuda/std/__iterator/concepts.h +++ b/libcudacxx/include/cuda/std/__iterator/concepts.h @@ -79,12 +79,12 @@ using iter_common_reference_t = common_reference_t, iter_v // [iterator.concept.writable] template concept indirectly_writable = requires(_Out&& __o, _Tp&& __t) { - *__o = static_cast<_Tp &&>(__t); // not required to be equality-preserving - *static_cast<_Out &&>(__o) = static_cast<_Tp &&>(__t); // not required to be equality-preserving - const_cast &&>(*__o) = static_cast<_Tp &&>(__t); // not required to be - // equality-preserving - const_cast &&>(*static_cast<_Out &&>(__o)) = - static_cast<_Tp &&>(__t); // not required to be equality-preserving + *__o = static_cast<_Tp&&>(__t); // not required to be equality-preserving + *static_cast<_Out&&>(__o) = static_cast<_Tp&&>(__t); // not required to be equality-preserving + const_cast&&>(*__o) = static_cast<_Tp&&>(__t); // not required to be + // equality-preserving + const_cast&&>(*static_cast<_Out&&>(__o)) = + static_cast<_Tp&&>(__t); // not required to be equality-preserving }; // [iterator.concept.winc] @@ -142,7 +142,7 @@ concept input_iterator = input_or_output_iterator<_Ip> && indirectly_readable<_I template concept output_iterator = input_or_output_iterator<_Ip> && indirectly_writable<_Ip, _Tp> && requires(_Ip __it, _Tp&& __t) { - *__it++ = static_cast<_Tp &&>(__t); // not required to be equality-preserving + *__it++ = static_cast<_Tp&&>(__t); // not required to be equality-preserving }; // [iterator.concept.forward] diff --git a/libcudacxx/include/cuda/std/__memory/construct_at.h b/libcudacxx/include/cuda/std/__memory/construct_at.h index 15abcc45f88..9cf98580965 100644 --- a/libcudacxx/include/cuda/std/__memory/construct_at.h +++ b/libcudacxx/include/cuda/std/__memory/construct_at.h @@ -48,7 +48,7 @@ namespace std _CCCL_EXEC_CHECK_DISABLE template ()) _Tp(::cuda::std::declval<_Args>()...))> + class = decltype(::new (::cuda::std::declval()) _Tp(::cuda::std::declval<_Args>()...))> _CCCL_API constexpr _Tp* construct_at(_Tp* __location, _Args&&... __args) { # if defined(_CCCL_BUILTIN_ADDRESSOF) diff --git a/libcudacxx/include/cuda/std/__memory/uninitialized_algorithms.h b/libcudacxx/include/cuda/std/__memory/uninitialized_algorithms.h index 07578d6a98f..443d28be3c3 100644 --- a/libcudacxx/include/cuda/std/__memory/uninitialized_algorithms.h +++ b/libcudacxx/include/cuda/std/__memory/uninitialized_algorithms.h @@ -651,8 +651,8 @@ template ::value_type, class = enable_if_t && is_trivially_move_assignable_v<_Type> && __allocator_has_trivial_move_construct<_Alloc, _Type>>> -_CCCL_API inline _CCCL_CONSTEXPR_CXX20 _Iter2 -__uninitialized_allocator_move_if_noexcept(_Alloc&, _Iter1 __first1, _Iter1 __last1, _Iter2 __first2) +_CCCL_API inline _CCCL_CONSTEXPR_CXX20 +_Iter2 __uninitialized_allocator_move_if_noexcept(_Alloc&, _Iter1 __first1, _Iter1 __last1, _Iter2 __first2) { _CCCL_IF_CONSTEVAL { diff --git a/libcudacxx/include/cuda/std/__memory/unique_ptr.h b/libcudacxx/include/cuda/std/__memory/unique_ptr.h index 477f6246488..26133068d05 100644 --- a/libcudacxx/include/cuda/std/__memory/unique_ptr.h +++ b/libcudacxx/include/cuda/std/__memory/unique_ptr.h @@ -535,8 +535,7 @@ _CCCL_API inline _CCCL_CONSTEXPR_CXX20 bool operator==(const unique_ptr<_T1, _D1 template _CCCL_API inline - bool - operator!=(const unique_ptr<_T1, _D1>& __x, const unique_ptr<_T2, _D2>& __y) + bool operator!=(const unique_ptr<_T1, _D1>& __x, const unique_ptr<_T2, _D2>& __y) { return !(__x == __y); } @@ -545,8 +544,7 @@ _CCCL_API inline template _CCCL_API inline - bool - operator<(const unique_ptr<_T1, _D1>& __x, const unique_ptr<_T2, _D2>& __y) + bool operator<(const unique_ptr<_T1, _D1>& __x, const unique_ptr<_T2, _D2>& __y) { using _P1 = typename unique_ptr<_T1, _D1>::pointer; using _P2 = typename unique_ptr<_T2, _D2>::pointer; @@ -604,8 +602,7 @@ _CCCL_API inline _CCCL_CONSTEXPR_CXX20 bool operator==(const unique_ptr<_T1, _D1 template _CCCL_API inline - bool - operator==(nullptr_t, const unique_ptr<_T1, _D1>& __x) noexcept + bool operator==(nullptr_t, const unique_ptr<_T1, _D1>& __x) noexcept { return !__x; } diff --git a/libcudacxx/include/cuda/std/__pstl/cuda/merge.h b/libcudacxx/include/cuda/std/__pstl/cuda/merge.h index 68f70b27b29..6123a735a92 100644 --- a/libcudacxx/include/cuda/std/__pstl/cuda/merge.h +++ b/libcudacxx/include/cuda/std/__pstl/cuda/merge.h @@ -78,7 +78,7 @@ struct __pstl_dispatch<__pstl_algorithm::__merge, __execution_backend::__cuda> iter_difference_t<_InputIterator1> __count1 = ::cuda::std::distance(__first1, __last1); iter_difference_t<_InputIterator2> __count2 = ::cuda::std::distance(__first2, __last2); auto __ret = __result + static_cast>(__count1) - + static_cast>(__count2); + + static_cast>(__count2); // We pass the policy as an environment to DeviceMerge _CCCL_TRY_CUDA_API( diff --git a/libcudacxx/include/cuda/std/__random/cauchy_distribution.h b/libcudacxx/include/cuda/std/__random/cauchy_distribution.h index 9f520baad79..2a8a1323a1a 100644 --- a/libcudacxx/include/cuda/std/__random/cauchy_distribution.h +++ b/libcudacxx/include/cuda/std/__random/cauchy_distribution.h @@ -135,8 +135,8 @@ class cauchy_distribution return __x.__p_ == __y.__p_; } #if _CCCL_STD_VER <= 2017 - [[nodiscard]] friend _CCCL_API constexpr bool - operator!=(const cauchy_distribution& __x, const cauchy_distribution& __y) noexcept + [[nodiscard]] friend + _CCCL_API constexpr bool operator!=(const cauchy_distribution& __x, const cauchy_distribution& __y) noexcept { return !(__x == __y); } diff --git a/libcudacxx/include/cuda/std/__random/fisher_f_distribution.h b/libcudacxx/include/cuda/std/__random/fisher_f_distribution.h index 64bc87a65f1..d2b49845012 100644 --- a/libcudacxx/include/cuda/std/__random/fisher_f_distribution.h +++ b/libcudacxx/include/cuda/std/__random/fisher_f_distribution.h @@ -138,8 +138,8 @@ class fisher_f_distribution return __x.__p_ == __y.__p_; } #if _CCCL_STD_VER <= 2017 - [[nodiscard]] friend _CCCL_API constexpr bool - operator!=(const fisher_f_distribution& __x, const fisher_f_distribution& __y) noexcept + [[nodiscard]] friend + _CCCL_API constexpr bool operator!=(const fisher_f_distribution& __x, const fisher_f_distribution& __y) noexcept { return !(__x == __y); } diff --git a/libcudacxx/include/cuda/std/__random/gamma_distribution.h b/libcudacxx/include/cuda/std/__random/gamma_distribution.h index cf0a0845a9f..838a0229ff6 100644 --- a/libcudacxx/include/cuda/std/__random/gamma_distribution.h +++ b/libcudacxx/include/cuda/std/__random/gamma_distribution.h @@ -200,8 +200,8 @@ class gamma_distribution return __x.__p_ == __y.__p_; } #if _CCCL_STD_VER <= 2017 - [[nodiscard]] friend _CCCL_API constexpr bool - operator!=(const gamma_distribution& __x, const gamma_distribution& __y) noexcept + [[nodiscard]] friend + _CCCL_API constexpr bool operator!=(const gamma_distribution& __x, const gamma_distribution& __y) noexcept { return !(__x == __y); } diff --git a/libcudacxx/include/cuda/std/__tuple_dir/tuple.h b/libcudacxx/include/cuda/std/__tuple_dir/tuple.h index 881c76371e4..0cfd266b1e7 100644 --- a/libcudacxx/include/cuda/std/__tuple_dir/tuple.h +++ b/libcudacxx/include/cuda/std/__tuple_dir/tuple.h @@ -451,7 +451,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT tuple // [tuple.assign]-15 template , __tuple_types>, enable_if_t<__can_assign<_Trait>, int> = 0> @@ -463,7 +463,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT tuple // [tuple.assign]-18 template , __tuple_types>, enable_if_t<__can_assign<_Trait>, int> = 0> @@ -502,7 +502,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT tuple // [tuple.assign]-39 template >, int> = 0, - __select_assignment _Trait = __tuple_select_tuple_like_assignable_v, __make_tuple_indices_t>, @@ -517,7 +517,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT tuple // [tuple.assign]-42 template >, int> = 0, - __select_assignment _Trait = __tuple_select_tuple_like_assignable_v, __make_tuple_indices_t>, diff --git a/libcudacxx/include/cuda/std/__type_traits/promote.h b/libcudacxx/include/cuda/std/__type_traits/promote.h index 2fc497b2659..68c186adafb 100644 --- a/libcudacxx/include/cuda/std/__type_traits/promote.h +++ b/libcudacxx/include/cuda/std/__type_traits/promote.h @@ -109,7 +109,7 @@ template ::value && __numeric_type<_A2>::value && __numeric_type<_A3>::value - && !__is_mixed_extended_floating_point<_A1, _A2, _A3>::value> + && !__is_mixed_extended_floating_point<_A1, _A2, _A3>::value> class __promote_imp { public: diff --git a/libcudacxx/include/cuda/std/__utility/constant_wrapper.h b/libcudacxx/include/cuda/std/__utility/constant_wrapper.h index 19fe4660c9d..7eee4b9ca39 100644 --- a/libcudacxx/include/cuda/std/__utility/constant_wrapper.h +++ b/libcudacxx/include/cuda/std/__utility/constant_wrapper.h @@ -135,16 +135,16 @@ struct __cw_operators // binary operators _CCCL_TEMPLATE(class _Lp, class _Rp) _CCCL_REQUIRES(__is_constexpr_param_v<_Lp> _CCCL_AND __is_constexpr_param_v<_Rp>) - [[nodiscard]] - _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper<__cw_fixed_value(_Lp::value + _Rp::value)>{}) + [[nodiscard]] _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper< + __cw_fixed_value(_Lp::value + _Rp::value)>{}) operator+(_Lp, _Rp) noexcept { return {}; } _CCCL_TEMPLATE(class _Lp, class _Rp) _CCCL_REQUIRES(__is_constexpr_param_v<_Lp> _CCCL_AND __is_constexpr_param_v<_Rp>) - [[nodiscard]] - _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper<__cw_fixed_value(_Lp::value - _Rp::value)>{}) + [[nodiscard]] _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper< + __cw_fixed_value(_Lp::value - _Rp::value)>{}) operator-(_Lp, _Rp) noexcept { return {}; @@ -159,16 +159,16 @@ struct __cw_operators } _CCCL_TEMPLATE(class _Lp, class _Rp) _CCCL_REQUIRES(__is_constexpr_param_v<_Lp> _CCCL_AND __is_constexpr_param_v<_Rp>) - [[nodiscard]] - _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper<__cw_fixed_value(_Lp::value / _Rp::value)>{}) + [[nodiscard]] _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper< + __cw_fixed_value(_Lp::value / _Rp::value)>{}) operator/(_Lp, _Rp) noexcept { return {}; } _CCCL_TEMPLATE(class _Lp, class _Rp) _CCCL_REQUIRES(__is_constexpr_param_v<_Lp> _CCCL_AND __is_constexpr_param_v<_Rp>) - [[nodiscard]] - _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper<__cw_fixed_value(_Lp::value % _Rp::value)>{}) + [[nodiscard]] _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper< + __cw_fixed_value(_Lp::value % _Rp::value)>{}) operator%(_Lp, _Rp) noexcept { return {}; @@ -176,16 +176,16 @@ struct __cw_operators _CCCL_TEMPLATE(class _Lp, class _Rp) _CCCL_REQUIRES(__is_constexpr_param_v<_Lp> _CCCL_AND __is_constexpr_param_v<_Rp>) - [[nodiscard]] - _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper<__cw_fixed_value(_Lp::value << _Rp::value)>{}) + [[nodiscard]] _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper<__cw_fixed_value(_Lp::value + << _Rp::value)>{}) operator<<(_Lp, _Rp) noexcept { return {}; } _CCCL_TEMPLATE(class _Lp, class _Rp) _CCCL_REQUIRES(__is_constexpr_param_v<_Lp> _CCCL_AND __is_constexpr_param_v<_Rp>) - [[nodiscard]] - _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper<__cw_fixed_value(_Lp::value >> _Rp::value)>{}) + [[nodiscard]] _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper< + __cw_fixed_value(_Lp::value >> _Rp::value)>{}) operator>>(_Lp, _Rp) noexcept { return {}; @@ -200,16 +200,16 @@ struct __cw_operators } _CCCL_TEMPLATE(class _Lp, class _Rp) _CCCL_REQUIRES(__is_constexpr_param_v<_Lp> _CCCL_AND __is_constexpr_param_v<_Rp>) - [[nodiscard]] - _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper<__cw_fixed_value(_Lp::value | _Rp::value)>{}) + [[nodiscard]] _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper<__cw_fixed_value(_Lp::value + | _Rp::value)>{}) operator|(_Lp, _Rp) noexcept { return {}; } _CCCL_TEMPLATE(class _Lp, class _Rp) _CCCL_REQUIRES(__is_constexpr_param_v<_Lp> _CCCL_AND __is_constexpr_param_v<_Rp>) - [[nodiscard]] - _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper<__cw_fixed_value(_Lp::value ^ _Rp::value)>{}) + [[nodiscard]] _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper<__cw_fixed_value(_Lp::value + ^ _Rp::value)>{}) operator^(_Lp, _Rp) noexcept { return {}; @@ -227,8 +227,8 @@ struct __cw_operators _CCCL_TEMPLATE(class _Lp, class _Rp) _CCCL_REQUIRES(__is_constexpr_param_v<_Lp> _CCCL_AND __is_constexpr_param_v<_Rp> _CCCL_AND( !is_constructible_v || !is_constructible_v)) - [[nodiscard]] - _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper<__cw_fixed_value(_Lp::value || _Rp::value)>{}) + [[nodiscard]] _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper<__cw_fixed_value(_Lp::value + || _Rp::value)>{}) operator||(_Lp, _Rp) noexcept { return {}; @@ -247,48 +247,48 @@ struct __cw_operators # endif // _LIBCUDACXX_HAS_SPACESHIP_OPERATOR() _CCCL_TEMPLATE(class _Lp, class _Rp) _CCCL_REQUIRES(__is_constexpr_param_v<_Lp> _CCCL_AND __is_constexpr_param_v<_Rp>) - [[nodiscard]] - _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper<__cw_fixed_value(_Lp::value < _Rp::value)>{}) + [[nodiscard]] _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper< + __cw_fixed_value(_Lp::value < _Rp::value)>{}) operator<(_Lp, _Rp) noexcept { return {}; } _CCCL_TEMPLATE(class _Lp, class _Rp) _CCCL_REQUIRES(__is_constexpr_param_v<_Lp> _CCCL_AND __is_constexpr_param_v<_Rp>) - [[nodiscard]] - _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper<__cw_fixed_value(_Lp::value <= _Rp::value)>{}) + [[nodiscard]] _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper< + __cw_fixed_value(_Lp::value <= _Rp::value)>{}) operator<=(_Lp, _Rp) noexcept { return {}; } _CCCL_TEMPLATE(class _Lp, class _Rp) _CCCL_REQUIRES(__is_constexpr_param_v<_Lp> _CCCL_AND __is_constexpr_param_v<_Rp>) - [[nodiscard]] - _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper<__cw_fixed_value(_Lp::value == _Rp::value)>{}) + [[nodiscard]] _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper< + __cw_fixed_value(_Lp::value == _Rp::value)>{}) operator==(_Lp, _Rp) noexcept { return {}; } _CCCL_TEMPLATE(class _Lp, class _Rp) _CCCL_REQUIRES(__is_constexpr_param_v<_Lp> _CCCL_AND __is_constexpr_param_v<_Rp>) - [[nodiscard]] - _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper<__cw_fixed_value(_Lp::value != _Rp::value)>{}) + [[nodiscard]] _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper< + __cw_fixed_value(_Lp::value != _Rp::value)>{}) operator!=(_Lp, _Rp) noexcept { return {}; } _CCCL_TEMPLATE(class _Lp, class _Rp) _CCCL_REQUIRES(__is_constexpr_param_v<_Lp> _CCCL_AND __is_constexpr_param_v<_Rp>) - [[nodiscard]] - _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper<__cw_fixed_value(_Lp::value > _Rp::value)>{}) + [[nodiscard]] _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper< + __cw_fixed_value(_Lp::value > _Rp::value)>{}) operator>(_Lp, _Rp) noexcept { return {}; } _CCCL_TEMPLATE(class _Lp, class _Rp) _CCCL_REQUIRES(__is_constexpr_param_v<_Lp> _CCCL_AND __is_constexpr_param_v<_Rp>) - [[nodiscard]] - _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper<__cw_fixed_value(_Lp::value >= _Rp::value)>{}) + [[nodiscard]] _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper< + __cw_fixed_value(_Lp::value >= _Rp::value)>{}) operator>=(_Lp, _Rp) noexcept { return {}; @@ -300,8 +300,8 @@ struct __cw_operators _CCCL_TEMPLATE(class _Lp, class _Rp) _CCCL_REQUIRES(__is_constexpr_param_v<_Lp> _CCCL_AND __is_constexpr_param_v<_Rp>) - [[nodiscard]] - _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper<__cw_fixed_value(_Lp::value->*_Rp::value)>{}) + [[nodiscard]] _CCCL_HOST_DEVICE_API friend consteval decltype(__constant_wrapper< + __cw_fixed_value(_Lp::value->*_Rp::value)>{}) operator->*(_Lp, _Rp) noexcept { return {}; @@ -398,8 +398,7 @@ struct __constant_wrapper : __cw_operators operator[](_Args&&...) noexcept # else // ^^^ _CCCL_HAS_STATIC_SUBSCRIPT_OPERATOR() ^^^ / vvv !_CCCL_HAS_STATIC_SUBSCRIPT_OPERATOR() vvv [[nodiscard]] _CCCL_HOST_DEVICE_API consteval __constant_wrapper< - __cw_fixed_value(value[remove_cvref_t<_Args>::value...])> - operator[](_Args&&...) const noexcept + __cw_fixed_value(value[remove_cvref_t<_Args>::value...])> operator[](_Args&&...) const noexcept # endif // ^^^ _CCCL_HAS_STATIC_SUBSCRIPT_OPERATOR() ^^^ { return {}; diff --git a/libcudacxx/include/cuda/std/array b/libcudacxx/include/cuda/std/array index bccd99f88ca..23fff8af833 100644 --- a/libcudacxx/include/cuda/std/array +++ b/libcudacxx/include/cuda/std/array @@ -535,8 +535,8 @@ _CCCL_END_NAMESPACE_CUDA_STD _CCCL_BEGIN_NAMESPACE_STD template -struct _CCCL_TYPE_VISIBILITY_DEFAULT tuple_size<::cuda::std::array<_Tp, _Np>> - : ::cuda::std::integral_constant<::cuda::std::size_t, _Np> +struct _CCCL_TYPE_VISIBILITY_DEFAULT +tuple_size<::cuda::std::array<_Tp, _Np>> : ::cuda::std::integral_constant<::cuda::std::size_t, _Np> {}; template <::cuda::std::size_t _Ip, class _Tp, ::cuda::std::size_t _Np> diff --git a/libcudacxx/include/cuda/std/bitset b/libcudacxx/include/cuda/std/bitset index 7b34347a685..1af35e8914c 100644 --- a/libcudacxx/include/cuda/std/bitset +++ b/libcudacxx/include/cuda/std/bitset @@ -107,7 +107,7 @@ struct __avoid_promotions #define _DEFINE_SHIFT_ASSIGNMENT(__op) \ template \ - _CCCL_API constexpr __avoid_promotions& operator __op##=(_Tp __n) \ + _CCCL_API constexpr __avoid_promotions& operator __op## = (_Tp __n) \ { \ if (__n >= sizeof(_Int) * CHAR_BIT) \ { \ @@ -142,7 +142,7 @@ struct __avoid_promotions #undef _DEFINE_BINARY #define _DEFINE_ASSIGNMENT(__op) \ - _CCCL_API constexpr __avoid_promotions& operator __op##=(__avoid_promotions __rhs) \ + _CCCL_API constexpr __avoid_promotions& operator __op## = (__avoid_promotions __rhs) \ { \ __data = static_cast<_Int>(static_cast<__base>(__data) __op static_cast<__base>(__rhs.__data)); \ return *this; \ diff --git a/libcudacxx/test/libcudacxx/cuda/utilities/tuple/forward_as_tuple_interop.pass.cpp b/libcudacxx/test/libcudacxx/cuda/utilities/tuple/forward_as_tuple_interop.pass.cpp index bda2a0c6402..4ddf0f70633 100644 --- a/libcudacxx/test/libcudacxx/cuda/utilities/tuple/forward_as_tuple_interop.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/utilities/tuple/forward_as_tuple_interop.pass.cpp @@ -34,7 +34,7 @@ constexpr bool test() { using ret = cuda::std::tuple, std::integral_constant>; auto t = cuda::std::tuple_cat(cuda::std::make_tuple(cuda::std::integral_constant()), - cuda::std::make_tuple(std::integral_constant())); + cuda::std::make_tuple(std::integral_constant())); static_assert(cuda::std::is_same::value); assert(cuda::std::get<0>(t) == 42); assert(cuda::std::get<1>(t) == 1337); diff --git a/libcudacxx/test/libcudacxx/force_include.h b/libcudacxx/test/libcudacxx/force_include.h index da97f48dcb9..f7d573b3ca3 100644 --- a/libcudacxx/test/libcudacxx/force_include.h +++ b/libcudacxx/test/libcudacxx/force_include.h @@ -48,8 +48,7 @@ void list_devices() #ifdef __CUDACC_TILE__ __tile__ #endif // __CUDACC_TILE__ - __host__ __device__ int - fake_main(int, char**); + __host__ __device__ int fake_main(int, char**); int cuda_thread_count = 1; int cuda_cluster_size = 1; @@ -59,8 +58,7 @@ __tile_global__ #else // ^^^ __CUDACC_TILE__ ^^^ / vvv !__CUDACC_TILE__ vvv __global__ #endif // !__CUDACC_TILE__ - void - fake_main_kernel(int* ret) + void fake_main_kernel(int* ret) { *ret = fake_main(0, nullptr); } diff --git a/libcudacxx/test/libcudacxx/std/concepts/concepts.compare/concept.equalitycomparable/equality_comparable.compile.pass.cpp b/libcudacxx/test/libcudacxx/std/concepts/concepts.compare/concept.equalitycomparable/equality_comparable.compile.pass.cpp index a11b57c5478..987b8dae6a6 100644 --- a/libcudacxx/test/libcudacxx/std/concepts/concepts.compare/concept.equalitycomparable/equality_comparable.compile.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/concepts/concepts.compare/concept.equalitycomparable/equality_comparable.compile.pass.cpp @@ -51,25 +51,25 @@ static_assert(equality_comparable); static_assert(equality_comparable); static_assert(equality_comparable); static_assert(equality_comparable); -static_assert(equality_comparable < int (S::*)() && noexcept >); +static_assert(equality_comparable); static_assert(equality_comparable); static_assert(equality_comparable); static_assert(equality_comparable); static_assert(equality_comparable); static_assert(equality_comparable); -static_assert(equality_comparable < int (S::*)() const&& noexcept >); +static_assert(equality_comparable); static_assert(equality_comparable); static_assert(equality_comparable); static_assert(equality_comparable); static_assert(equality_comparable); static_assert(equality_comparable); -static_assert(equality_comparable < int (S::*)() volatile && noexcept >); +static_assert(equality_comparable); static_assert(equality_comparable); static_assert(equality_comparable); static_assert(equality_comparable); static_assert(equality_comparable); static_assert(equality_comparable); -static_assert(equality_comparable < int (S::*)() const volatile&& noexcept >); +static_assert(equality_comparable); static_assert(!equality_comparable); } // namespace fundamentals diff --git a/libcudacxx/test/libcudacxx/std/concepts/concepts.compare/concepts.totallyordered/totally_ordered.pass.cpp b/libcudacxx/test/libcudacxx/std/concepts/concepts.compare/concepts.totallyordered/totally_ordered.pass.cpp index 966bf42e73d..374b41bd0e2 100644 --- a/libcudacxx/test/libcudacxx/std/concepts/concepts.compare/concepts.totallyordered/totally_ordered.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/concepts/concepts.compare/concepts.totallyordered/totally_ordered.pass.cpp @@ -92,7 +92,7 @@ static_assert(!totally_ordered); static_assert(!totally_ordered); static_assert(!totally_ordered); static_assert(!totally_ordered); -static_assert(!totally_ordered < int (S::*)() const&& noexcept >); +static_assert(!totally_ordered < int (S::*)() const && noexcept >); static_assert(!totally_ordered); static_assert(!totally_ordered); static_assert(!totally_ordered); @@ -104,7 +104,7 @@ static_assert(!totally_ordered); static_assert(!totally_ordered); static_assert(!totally_ordered); static_assert(!totally_ordered); -static_assert(!totally_ordered < int (S::*)() const volatile&& noexcept >); +static_assert(!totally_ordered < int (S::*)() const volatile && noexcept >); static_assert(!totally_ordered); } // namespace fundamentals diff --git a/libcudacxx/test/libcudacxx/std/concepts/concepts.object/copyable.compile.pass.cpp b/libcudacxx/test/libcudacxx/std/concepts/concepts.object/copyable.compile.pass.cpp index 72741ad78b7..26ec2eaf8b8 100644 --- a/libcudacxx/test/libcudacxx/std/concepts/concepts.object/copyable.compile.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/concepts/concepts.object/copyable.compile.pass.cpp @@ -36,25 +36,25 @@ static_assert(copyable); static_assert(copyable); static_assert(copyable); static_assert(copyable); -static_assert(copyable < int (S::*)() && noexcept >); +static_assert(copyable); static_assert(copyable); static_assert(copyable); static_assert(copyable); static_assert(copyable); static_assert(copyable); -static_assert(copyable < int (S::*)() const&& noexcept >); +static_assert(copyable); static_assert(copyable); static_assert(copyable); static_assert(copyable); static_assert(copyable); static_assert(copyable); -static_assert(copyable < int (S::*)() volatile && noexcept >); +static_assert(copyable); static_assert(copyable); static_assert(copyable); static_assert(copyable); static_assert(copyable); static_assert(copyable); -static_assert(copyable < int (S::*)() const volatile&& noexcept >); +static_assert(copyable); static_assert(copyable); static_assert(copyable); diff --git a/libcudacxx/test/libcudacxx/std/concepts/concepts.object/movable.compile.pass.cpp b/libcudacxx/test/libcudacxx/std/concepts/concepts.object/movable.compile.pass.cpp index 279af5fd078..695915f1dd0 100644 --- a/libcudacxx/test/libcudacxx/std/concepts/concepts.object/movable.compile.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/concepts/concepts.object/movable.compile.pass.cpp @@ -37,25 +37,25 @@ static_assert(movable); static_assert(movable); static_assert(movable); static_assert(movable); -static_assert(movable < int (S::*)() && noexcept >); +static_assert(movable); static_assert(movable); static_assert(movable); static_assert(movable); static_assert(movable); static_assert(movable); -static_assert(movable < int (S::*)() const&& noexcept >); +static_assert(movable); static_assert(movable); static_assert(movable); static_assert(movable); static_assert(movable); static_assert(movable); -static_assert(movable < int (S::*)() volatile && noexcept >); +static_assert(movable); static_assert(movable); static_assert(movable); static_assert(movable); static_assert(movable); static_assert(movable); -static_assert(movable < int (S::*)() const volatile&& noexcept >); +static_assert(movable); static_assert(movable); static_assert(movable); diff --git a/libcudacxx/test/libcudacxx/std/concepts/concepts.object/regular.compile.pass.cpp b/libcudacxx/test/libcudacxx/std/concepts/concepts.object/regular.compile.pass.cpp index 81dce501255..1a18b1ad983 100644 --- a/libcudacxx/test/libcudacxx/std/concepts/concepts.object/regular.compile.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/concepts/concepts.object/regular.compile.pass.cpp @@ -40,25 +40,25 @@ static_assert(regular); static_assert(regular); static_assert(regular); static_assert(regular); -static_assert(regular < int (S::*)() && noexcept >); +static_assert(regular); static_assert(regular); static_assert(regular); static_assert(regular); static_assert(regular); static_assert(regular); -static_assert(regular < int (S::*)() const&& noexcept >); +static_assert(regular); static_assert(regular); static_assert(regular); static_assert(regular); static_assert(regular); static_assert(regular); -static_assert(regular < int (S::*)() volatile && noexcept >); +static_assert(regular); static_assert(regular); static_assert(regular); static_assert(regular); static_assert(regular); static_assert(regular); -static_assert(regular < int (S::*)() const volatile&& noexcept >); +static_assert(regular); union U {}; @@ -69,25 +69,25 @@ static_assert(regular); static_assert(regular); static_assert(regular); static_assert(regular); -static_assert(regular < int (U::*)() && noexcept >); +static_assert(regular); static_assert(regular); static_assert(regular); static_assert(regular); static_assert(regular); static_assert(regular); -static_assert(regular < int (U::*)() const&& noexcept >); +static_assert(regular); static_assert(regular); static_assert(regular); static_assert(regular); static_assert(regular); static_assert(regular); -static_assert(regular < int (U::*)() volatile && noexcept >); +static_assert(regular); static_assert(regular); static_assert(regular); static_assert(regular); static_assert(regular); static_assert(regular); -static_assert(regular < int (U::*)() const volatile&& noexcept >); +static_assert(regular); static_assert(!regular); static_assert(!regular); diff --git a/libcudacxx/test/libcudacxx/std/concepts/concepts.object/semiregular.compile.pass.cpp b/libcudacxx/test/libcudacxx/std/concepts/concepts.object/semiregular.compile.pass.cpp index 0d10a5f79a3..1b4ea217fc3 100644 --- a/libcudacxx/test/libcudacxx/std/concepts/concepts.object/semiregular.compile.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/concepts/concepts.object/semiregular.compile.pass.cpp @@ -36,25 +36,25 @@ static_assert(semiregular); static_assert(semiregular); static_assert(semiregular); static_assert(semiregular); -static_assert(semiregular < int (S::*)() && noexcept >); +static_assert(semiregular); static_assert(semiregular); static_assert(semiregular); static_assert(semiregular); static_assert(semiregular); static_assert(semiregular); -static_assert(semiregular < int (S::*)() const&& noexcept >); +static_assert(semiregular); static_assert(semiregular); static_assert(semiregular); static_assert(semiregular); static_assert(semiregular); static_assert(semiregular); -static_assert(semiregular < int (S::*)() volatile && noexcept >); +static_assert(semiregular); static_assert(semiregular); static_assert(semiregular); static_assert(semiregular); static_assert(semiregular); static_assert(semiregular); -static_assert(semiregular < int (S::*)() const volatile&& noexcept >); +static_assert(semiregular); static_assert(semiregular); static_assert(semiregular); diff --git a/libcudacxx/test/libcudacxx/std/containers/sequences/array/iterators.pass.cpp b/libcudacxx/test/libcudacxx/std/containers/sequences/array/iterators.pass.cpp index d04a2a9621d..db2fc8f4a24 100644 --- a/libcudacxx/test/libcudacxx/std/containers/sequences/array/iterators.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/containers/sequences/array/iterators.pass.cpp @@ -61,8 +61,7 @@ TEST_FUNC #if !TEST_COMPILER(GCC, <, 8) constexpr #endif // !TEST_COMPILER(GCC, <, 8) - bool - tests() + bool tests() { { using C = cuda::std::array; diff --git a/libcudacxx/test/libcudacxx/std/containers/views/mdspan/layout_stride/comparison.pass.cpp b/libcudacxx/test/libcudacxx/std/containers/views/mdspan/layout_stride/comparison.pass.cpp index 7e0b29f58cb..3d66cf49a7b 100644 --- a/libcudacxx/test/libcudacxx/std/containers/views/mdspan/layout_stride/comparison.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/containers/views/mdspan/layout_stride/comparison.pass.cpp @@ -243,7 +243,7 @@ TEST_FUNC constexpr void test_comparison_with() { [[maybe_unused]] constexpr size_t D = cuda::std::dynamic_extent; bool is_left_based = cuda::std::is_same_v - || cuda::std::is_same_v; + || cuda::std::is_same_v; test_comparison_with( true, cuda::std::extents(), cuda::std::array{}, cuda::std::extents()); test_comparison_with( diff --git a/libcudacxx/test/libcudacxx/std/iterators/iterator.requirements/iterator.concepts/iterator.concept.readable/indirectly_readable.compile.pass.cpp b/libcudacxx/test/libcudacxx/std/iterators/iterator.requirements/iterator.concepts/iterator.concept.readable/indirectly_readable.compile.pass.cpp index 4313c00317d..ce2673348b0 100644 --- a/libcudacxx/test/libcudacxx/std/iterators/iterator.requirements/iterator.concepts/iterator.concept.readable/indirectly_readable.compile.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/iterators/iterator.requirements/iterator.concepts/iterator.concept.readable/indirectly_readable.compile.pass.cpp @@ -218,7 +218,7 @@ static_assert(!cuda::std::indirectly_readable); static_assert(!cuda::std::indirectly_readable); static_assert(!cuda::std::indirectly_readable); static_assert(!cuda::std::indirectly_readable); -static_assert(!cuda::std::indirectly_readable < int (S::*)() const&& noexcept >); +static_assert(!cuda::std::indirectly_readable < int (S::*)() const && noexcept >); static_assert(!cuda::std::indirectly_readable); static_assert(!cuda::std::indirectly_readable); static_assert(!cuda::std::indirectly_readable); diff --git a/libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/call.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/call.pass.cpp index 26deca904bb..bd9d993d0d1 100644 --- a/libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/call.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/call.pass.cpp @@ -191,7 +191,7 @@ TEST_FUNC constexpr bool test() // nullary using T = cuda::std::__constant_wrapper<[] { return 42; - }>; + }>; cuda::std::same_as> decltype(auto) result = TEST_CALL(T, ); static_assert(result == 42); #endif // !_CCCL_COMPILER(MSVC) diff --git a/libcudacxx/test/libcudacxx/std/utilities/function.objects/func.invoke/invoke.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/function.objects/func.invoke/invoke.pass.cpp index 8ebfc9a9944..f0dfda115a7 100644 --- a/libcudacxx/test/libcudacxx/std/utilities/function.objects/func.invoke/invoke.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/utilities/function.objects/func.invoke/invoke.pass.cpp @@ -189,7 +189,7 @@ TEST_FUNC void bullet_one_two_tests() test_b12(cl); test_b12(cl); - test_b12(cuda::std::move(cl)); + test_b12(cuda::std::move(cl)); test_b12(cuda::std::move(cl)); test_b12(cuda::std::move(cl)); test_b12(cuda::std::move(cl)); @@ -201,7 +201,7 @@ TEST_FUNC void bullet_one_two_tests() test_b12(cl); test_b12(cl); - test_b12(cuda::std::move(cl)); + test_b12(cuda::std::move(cl)); test_b12(cuda::std::move(cl)); test_b12(cuda::std::move(cl)); test_b12(cuda::std::move(cl)); diff --git a/libcudacxx/test/libcudacxx/std/utilities/optional/optional.object/triviality.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/optional/optional.object/triviality.pass.cpp index 9daa0232647..4e26cc90b62 100644 --- a/libcudacxx/test/libcudacxx/std/utilities/optional/optional.object/triviality.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/utilities/optional/optional.object/triviality.pass.cpp @@ -45,8 +45,8 @@ struct SpecialMemberTest cuda::std::is_trivially_move_constructible_v), "optional is trivially move constructible if T is trivially move constructible"); - static_assert(implies(cuda::std::is_trivially_copy_constructible_v && cuda::std::is_trivially_copy_assignable_v - && cuda::std::is_trivially_destructible_v, + static_assert(implies(cuda::std::is_trivially_copy_constructible_v&& cuda::std::is_trivially_copy_assignable_v&& + cuda::std::is_trivially_destructible_v, cuda::std::is_trivially_copy_assignable_v), "optional is trivially copy assignable if T is " @@ -54,8 +54,8 @@ struct SpecialMemberTest "trivially copy assignable, and " "trivially destructible"); - static_assert(implies(cuda::std::is_trivially_move_constructible_v && cuda::std::is_trivially_move_assignable_v - && cuda::std::is_trivially_destructible_v, + static_assert(implies(cuda::std::is_trivially_move_constructible_v&& cuda::std::is_trivially_move_assignable_v&& + cuda::std::is_trivially_destructible_v, cuda::std::is_trivially_move_assignable_v), "optional is trivially move assignable if T is " diff --git a/libcudacxx/test/libcudacxx/std/utilities/optional/optional.specalg/swap.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/optional/optional.specalg/swap.pass.cpp index 637d7a05aa9..762768864e6 100644 --- a/libcudacxx/test/libcudacxx/std/utilities/optional/optional.specalg/swap.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/utilities/optional/optional.specalg/swap.pass.cpp @@ -185,8 +185,7 @@ struct NonSwappable NonSwappable(NonSwappable const&) = delete; }; TEST_FUNC // what in the world, nvrtc?! - void - swap(NonSwappable&, NonSwappable&) = delete; + void swap(NonSwappable&, NonSwappable&) = delete; TEST_FUNC void test_swap_sfinae() { diff --git a/libcudacxx/test/libcudacxx/std/utilities/template.bitset/bitset.members/to_ullong.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/template.bitset/bitset.members/to_ullong.pass.cpp index ac33102a9c5..108df751475 100644 --- a/libcudacxx/test/libcudacxx/std/utilities/template.bitset/bitset.members/to_ullong.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/utilities/template.bitset/bitset.members/to_ullong.pass.cpp @@ -20,7 +20,7 @@ template TEST_FUNC constexpr void test_to_ullong() { const cuda::std::size_t M = sizeof(unsigned long long) * CHAR_BIT < N ? sizeof(unsigned long long) * CHAR_BIT : N; - const bool is_M_zero = cuda::std::integral_constant < bool, M == 0 > ::value; // avoid compiler warnings + const bool is_M_zero = cuda::std::integral_constant::value; // avoid compiler warnings const cuda::std::size_t X = is_M_zero ? sizeof(unsigned long long) * CHAR_BIT - 1 : sizeof(unsigned long long) * CHAR_BIT - M; const unsigned long long max = is_M_zero ? 0 : (unsigned long long) (-1) >> X; diff --git a/libcudacxx/test/libcudacxx/std/utilities/template.bitset/bitset.members/to_ulong.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/template.bitset/bitset.members/to_ulong.pass.cpp index 8225c0b7f87..dc52746e155 100644 --- a/libcudacxx/test/libcudacxx/std/utilities/template.bitset/bitset.members/to_ulong.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/utilities/template.bitset/bitset.members/to_ulong.pass.cpp @@ -21,7 +21,7 @@ template TEST_FUNC constexpr void test_to_ulong() { const cuda::std::size_t M = sizeof(unsigned long) * CHAR_BIT < N ? sizeof(unsigned long) * CHAR_BIT : N; - const bool is_M_zero = cuda::std::integral_constant < bool, M == 0 > ::value; // avoid compiler warnings + const bool is_M_zero = cuda::std::integral_constant::value; // avoid compiler warnings const cuda::std::size_t X = is_M_zero ? sizeof(unsigned long) * CHAR_BIT - 1 : sizeof(unsigned long) * CHAR_BIT - M; const cuda::std::size_t max = is_M_zero ? 0 : cuda::std::size_t(cuda::std::numeric_limits::max()) >> X; cuda::std::size_t tests[] = { diff --git a/libcudacxx/test/nvtarget/arch_specific/arch_specific.cu b/libcudacxx/test/nvtarget/arch_specific/arch_specific.cu index c806bac123e..47f05478efd 100644 --- a/libcudacxx/test/nvtarget/arch_specific/arch_specific.cu +++ b/libcudacxx/test/nvtarget/arch_specific/arch_specific.cu @@ -37,8 +37,7 @@ #ifdef __CUDACC_TILE__ __tile__ #endif // __CUDACC_TILE__ - __host__ __device__ void - fn() + __host__ __device__ void fn() { #if defined(__CUDA_ARCH__) CHECK_TRUE(NV_IS_EXACTLY_SM_103); diff --git a/libcudacxx/test/nvtarget/dialect/nvtargettest.cu b/libcudacxx/test/nvtarget/dialect/nvtargettest.cu index bbd2c10f54a..1c1ce0bf246 100644 --- a/libcudacxx/test/nvtarget/dialect/nvtargettest.cu +++ b/libcudacxx/test/nvtarget/dialect/nvtargettest.cu @@ -13,6 +13,5 @@ #ifdef __CUDACC_TILE__ __tile__ #endif // __CUDACC_TILE__ - __host__ __device__ void - fn() + __host__ __device__ void fn() {} diff --git a/libcudacxx/test/nvtarget/family_specific/family_specific.cu b/libcudacxx/test/nvtarget/family_specific/family_specific.cu index f953908b872..9898581d587 100644 --- a/libcudacxx/test/nvtarget/family_specific/family_specific.cu +++ b/libcudacxx/test/nvtarget/family_specific/family_specific.cu @@ -37,8 +37,7 @@ #ifdef __CUDACC_TILE__ __tile__ #endif // __CUDACC_TILE__ - __host__ __device__ void - fn() + __host__ __device__ void fn() { #if defined(__CUDA_ARCH__) CHECK_TRUE(NV_IS_EXACTLY_SM_103); diff --git a/libcudacxx/test/support/test_execution_policies.h b/libcudacxx/test/support/test_execution_policies.h index 5995dde05bc..9573fa0f20f 100644 --- a/libcudacxx/test/support/test_execution_policies.h +++ b/libcudacxx/test/support/test_execution_policies.h @@ -17,17 +17,18 @@ #include "test_macros.h" -#define EXECUTION_POLICY_SFINAE_TEST(FUNCTION) \ - template \ - struct sfinae_test_##FUNCTION##_impl : cuda::std::false_type \ - {}; \ - \ - template \ - struct sfinae_test_##FUNCTION##_impl()...))>, \ - Args...> : cuda::std::true_type \ - {}; \ - \ - template \ +#define EXECUTION_POLICY_SFINAE_TEST(FUNCTION) \ + template \ + struct sfinae_test_##FUNCTION##_impl : cuda::std::false_type \ + {}; \ + \ + template \ + struct sfinae_test_##FUNCTION## \ + _impl()...))>, Args...> \ + : cuda::std::true_type \ + {}; \ + \ + template \ inline constexpr bool sfinae_test_##FUNCTION = sfinae_test_##FUNCTION##_impl::value; _CCCL_EXEC_CHECK_DISABLE diff --git a/python/cuda_cccl/benchmarks/compute/histogram/even.py b/python/cuda_cccl/benchmarks/compute/histogram/even.py index 1d11b4af3de..8157f0b32a6 100644 --- a/python/cuda_cccl/benchmarks/compute/histogram/even.py +++ b/python/cuda_cccl/benchmarks/compute/histogram/even.py @@ -17,12 +17,12 @@ sys.path.insert(0, str(Path(__file__).parent.parent)) +import cuda.bench as bench import cupy as cp import numpy as np from utils import FUNDAMENTAL_TYPES as TYPE_MAP from utils import as_cupy_stream, generate_data_with_entropy -import cuda.bench as bench from cuda.compute import make_histogram_even diff --git a/python/cuda_cccl/benchmarks/compute/host/host_benchmark_cases.py b/python/cuda_cccl/benchmarks/compute/host/host_benchmark_cases.py index ceaa51709d6..2af8575a1dc 100644 --- a/python/cuda_cccl/benchmarks/compute/host/host_benchmark_cases.py +++ b/python/cuda_cccl/benchmarks/compute/host/host_benchmark_cases.py @@ -10,11 +10,11 @@ import cupy as cp import numpy as np +from cuda.core import Device import cuda.compute as cc from cuda.compute._cpp_compile import compile_cpp_op_code from cuda.compute.op import RawOp -from cuda.core import Device NOOP_TEMP_STORAGE_BYTES = 1 NUM_ITEMS = 128 diff --git a/python/cuda_cccl/benchmarks/compute/merge_sort/keys.py b/python/cuda_cccl/benchmarks/compute/merge_sort/keys.py index 3d24d98f811..8232d870a2d 100644 --- a/python/cuda_cccl/benchmarks/compute/merge_sort/keys.py +++ b/python/cuda_cccl/benchmarks/compute/merge_sort/keys.py @@ -20,11 +20,11 @@ sys.path.insert(0, str(Path(__file__).parent.parent)) +import cuda.bench as bench import cupy as cp import numpy as np from utils import SIGNED_TYPES, as_cupy_stream, generate_data_with_entropy -import cuda.bench as bench from cuda.compute import OpKind, make_merge_sort diff --git a/python/cuda_cccl/benchmarks/compute/merge_sort/pairs.py b/python/cuda_cccl/benchmarks/compute/merge_sort/pairs.py index efa47e8511f..6dfd48f8965 100644 --- a/python/cuda_cccl/benchmarks/compute/merge_sort/pairs.py +++ b/python/cuda_cccl/benchmarks/compute/merge_sort/pairs.py @@ -19,6 +19,7 @@ sys.path.insert(0, str(Path(__file__).parent.parent)) +import cuda.bench as bench import cupy as cp import numpy as np from utils import ( @@ -28,7 +29,6 @@ generate_data_with_entropy, ) -import cuda.bench as bench from cuda.compute import OpKind, make_merge_sort KEY_TYPE_MAP = SIGNED_TYPES diff --git a/python/cuda_cccl/benchmarks/compute/partition/three_way.py b/python/cuda_cccl/benchmarks/compute/partition/three_way.py index d46423cb60e..d03b543bf16 100644 --- a/python/cuda_cccl/benchmarks/compute/partition/three_way.py +++ b/python/cuda_cccl/benchmarks/compute/partition/three_way.py @@ -23,11 +23,11 @@ sys.path.insert(0, str(Path(__file__).parent.parent)) +import cuda.bench as bench import cupy as cp import numpy as np from utils import FUNDAMENTAL_TYPES, as_cupy_stream, generate_data_with_entropy -import cuda.bench as bench from cuda.compute import make_three_way_partition diff --git a/python/cuda_cccl/benchmarks/compute/radix_sort/keys.py b/python/cuda_cccl/benchmarks/compute/radix_sort/keys.py index dadd231e926..cd6f62a2ebe 100644 --- a/python/cuda_cccl/benchmarks/compute/radix_sort/keys.py +++ b/python/cuda_cccl/benchmarks/compute/radix_sort/keys.py @@ -21,12 +21,12 @@ sys.path.insert(0, str(Path(__file__).parent.parent)) +import cuda.bench as bench import cupy as cp import numpy as np from utils import FUNDAMENTAL_TYPES as TYPE_MAP from utils import as_cupy_stream, generate_data_with_entropy -import cuda.bench as bench from cuda.compute import SortOrder, make_radix_sort diff --git a/python/cuda_cccl/benchmarks/compute/radix_sort/pairs.py b/python/cuda_cccl/benchmarks/compute/radix_sort/pairs.py index 9b604dff8e1..adae677108a 100644 --- a/python/cuda_cccl/benchmarks/compute/radix_sort/pairs.py +++ b/python/cuda_cccl/benchmarks/compute/radix_sort/pairs.py @@ -22,11 +22,11 @@ sys.path.insert(0, str(Path(__file__).parent.parent)) +import cuda.bench as bench import cupy as cp import numpy as np from utils import INTEGRAL_TYPES, as_cupy_stream, generate_data_with_entropy -import cuda.bench as bench from cuda.compute import SortOrder, make_radix_sort KEY_TYPE_MAP = INTEGRAL_TYPES diff --git a/python/cuda_cccl/benchmarks/compute/reduce/custom.py b/python/cuda_cccl/benchmarks/compute/reduce/custom.py index ee0ff65cea1..c9b76f46a15 100644 --- a/python/cuda_cccl/benchmarks/compute/reduce/custom.py +++ b/python/cuda_cccl/benchmarks/compute/reduce/custom.py @@ -18,12 +18,12 @@ sys.path.insert(0, str(Path(__file__).parent.parent)) +import cuda.bench as bench import cupy as cp import numpy as np from utils import SIGNED_TYPES as TYPE_MAP from utils import as_cupy_stream, generate_data_with_entropy -import cuda.bench as bench from cuda.compute import make_reduce_into diff --git a/python/cuda_cccl/benchmarks/compute/reduce/min.py b/python/cuda_cccl/benchmarks/compute/reduce/min.py index 82c51520af8..12a2f648ce8 100644 --- a/python/cuda_cccl/benchmarks/compute/reduce/min.py +++ b/python/cuda_cccl/benchmarks/compute/reduce/min.py @@ -18,12 +18,12 @@ sys.path.insert(0, str(Path(__file__).parent.parent)) +import cuda.bench as bench import cupy as cp import numpy as np from utils import FUNDAMENTAL_TYPES as TYPE_MAP from utils import as_cupy_stream, generate_data_with_entropy -import cuda.bench as bench from cuda.compute import OpKind, make_reduce_into diff --git a/python/cuda_cccl/benchmarks/compute/reduce/nondeterministic.py b/python/cuda_cccl/benchmarks/compute/reduce/nondeterministic.py index 6e9d2e8f1cb..d2e34624c13 100644 --- a/python/cuda_cccl/benchmarks/compute/reduce/nondeterministic.py +++ b/python/cuda_cccl/benchmarks/compute/reduce/nondeterministic.py @@ -19,12 +19,12 @@ sys.path.insert(0, str(Path(__file__).parent.parent)) +import cuda.bench as bench import cupy as cp import numpy as np from utils import ALL_TYPES as _ALL_TYPES from utils import as_cupy_stream, generate_data_with_entropy -import cuda.bench as bench from cuda.compute import Determinism, OpKind, make_reduce_into TYPE_MAP = {k: _ALL_TYPES[k] for k in ("I32", "I64", "F32", "F64")} diff --git a/python/cuda_cccl/benchmarks/compute/reduce/sum.py b/python/cuda_cccl/benchmarks/compute/reduce/sum.py index 765af40b23a..c10436ec474 100644 --- a/python/cuda_cccl/benchmarks/compute/reduce/sum.py +++ b/python/cuda_cccl/benchmarks/compute/reduce/sum.py @@ -17,12 +17,12 @@ sys.path.insert(0, str(Path(__file__).parent.parent)) +import cuda.bench as bench import cupy as cp import numpy as np from utils import SIGNED_TYPES as TYPE_MAP from utils import as_cupy_stream, generate_data_with_entropy -import cuda.bench as bench from cuda.compute import OpKind, make_reduce_into diff --git a/python/cuda_cccl/benchmarks/compute/scan/exclusive/custom.py b/python/cuda_cccl/benchmarks/compute/scan/exclusive/custom.py index 703869a644c..e6b511583ef 100644 --- a/python/cuda_cccl/benchmarks/compute/scan/exclusive/custom.py +++ b/python/cuda_cccl/benchmarks/compute/scan/exclusive/custom.py @@ -19,12 +19,12 @@ sys.path.insert(0, str(Path(__file__).parent.parent.parent)) +import cuda.bench as bench import cupy as cp import numpy as np from utils import SIGNED_TYPES as TYPE_MAP from utils import as_cupy_stream, generate_data_with_entropy -import cuda.bench as bench from cuda.compute import make_exclusive_scan diff --git a/python/cuda_cccl/benchmarks/compute/scan/exclusive/sum.py b/python/cuda_cccl/benchmarks/compute/scan/exclusive/sum.py index 7a1d5130abc..8fea9a9a28e 100644 --- a/python/cuda_cccl/benchmarks/compute/scan/exclusive/sum.py +++ b/python/cuda_cccl/benchmarks/compute/scan/exclusive/sum.py @@ -18,12 +18,12 @@ sys.path.insert(0, str(Path(__file__).parent.parent.parent)) +import cuda.bench as bench import cupy as cp import numpy as np from utils import SIGNED_TYPES as TYPE_MAP from utils import as_cupy_stream, generate_data_with_entropy -import cuda.bench as bench from cuda.compute import OpKind, make_exclusive_scan diff --git a/python/cuda_cccl/benchmarks/compute/segmented_reduce/variable_sum.py b/python/cuda_cccl/benchmarks/compute/segmented_reduce/variable_sum.py index 9d27dfc738c..98065dc51c6 100644 --- a/python/cuda_cccl/benchmarks/compute/segmented_reduce/variable_sum.py +++ b/python/cuda_cccl/benchmarks/compute/segmented_reduce/variable_sum.py @@ -18,6 +18,7 @@ sys.path.insert(0, str(Path(__file__).parent.parent)) +import cuda.bench as bench import cupy as cp import numpy as np from utils import ( @@ -27,7 +28,6 @@ generate_uniform_segment_offsets, ) -import cuda.bench as bench from cuda.compute import OpKind, make_segmented_reduce TYPE_MAP = {k: ALL_TYPES[k] for k in ("I32", "I64", "F32", "F64")} diff --git a/python/cuda_cccl/benchmarks/compute/segmented_sort/keys.py b/python/cuda_cccl/benchmarks/compute/segmented_sort/keys.py index b66c935f622..ec177635086 100644 --- a/python/cuda_cccl/benchmarks/compute/segmented_sort/keys.py +++ b/python/cuda_cccl/benchmarks/compute/segmented_sort/keys.py @@ -20,6 +20,7 @@ sys.path.insert(0, str(Path(__file__).parent.parent)) +import cuda.bench as bench import cupy as cp import numpy as np from utils import ( @@ -32,7 +33,6 @@ generate_uniform_segment_offsets, ) -import cuda.bench as bench from cuda.compute import SortOrder, make_segmented_sort diff --git a/python/cuda_cccl/benchmarks/compute/select/if.py b/python/cuda_cccl/benchmarks/compute/select/if.py index 64fb9836418..1cdf220c4d2 100644 --- a/python/cuda_cccl/benchmarks/compute/select/if.py +++ b/python/cuda_cccl/benchmarks/compute/select/if.py @@ -22,6 +22,7 @@ sys.path.insert(0, str(Path(__file__).parent.parent)) +import cuda.bench as bench import cupy as cp import numpy as np from utils import ( @@ -34,7 +35,6 @@ FUNDAMENTAL_TYPES as TYPE_MAP, ) -import cuda.bench as bench from cuda.compute import make_select # Entropy values from C++ benchmark diff --git a/python/cuda_cccl/benchmarks/compute/select/unique_by_key.py b/python/cuda_cccl/benchmarks/compute/select/unique_by_key.py index 8de29055f20..b2d284ec22a 100644 --- a/python/cuda_cccl/benchmarks/compute/select/unique_by_key.py +++ b/python/cuda_cccl/benchmarks/compute/select/unique_by_key.py @@ -21,11 +21,11 @@ sys.path.insert(0, str(Path(__file__).parent.parent)) +import cuda.bench as bench import cupy as cp import numpy as np from utils import INTEGRAL_TYPES, SIGNED_TYPES, as_cupy_stream, generate_key_segments -import cuda.bench as bench from cuda.compute import OpKind, make_unique_by_key KEY_TYPE_MAP = INTEGRAL_TYPES diff --git a/python/cuda_cccl/benchmarks/compute/transform/babelstream.py b/python/cuda_cccl/benchmarks/compute/transform/babelstream.py index 5a256af875c..7e4ecfbdcf3 100644 --- a/python/cuda_cccl/benchmarks/compute/transform/babelstream.py +++ b/python/cuda_cccl/benchmarks/compute/transform/babelstream.py @@ -17,11 +17,11 @@ sys.path.insert(0, str(Path(__file__).parent.parent)) +import cuda.bench as bench import cupy as cp from utils import ALL_TYPES as _ALL_TYPES from utils import as_cupy_stream -import cuda.bench as bench import cuda.compute from cuda.compute import ZipIterator diff --git a/python/cuda_cccl/benchmarks/compute/transform/complex_cmp.py b/python/cuda_cccl/benchmarks/compute/transform/complex_cmp.py index 5721e7576a1..0b9dbee35b5 100644 --- a/python/cuda_cccl/benchmarks/compute/transform/complex_cmp.py +++ b/python/cuda_cccl/benchmarks/compute/transform/complex_cmp.py @@ -21,11 +21,11 @@ import math +import cuda.bench as bench import cupy as cp import numpy as np from utils import as_cupy_stream, generate_data_with_entropy -import cuda.bench as bench from cuda.compute import make_binary_transform _COMPLEX_EPS = np.finfo(np.float32).eps diff --git a/python/cuda_cccl/benchmarks/compute/transform/fib.py b/python/cuda_cccl/benchmarks/compute/transform/fib.py index 7ab77300fde..bceb4b0aab4 100644 --- a/python/cuda_cccl/benchmarks/compute/transform/fib.py +++ b/python/cuda_cccl/benchmarks/compute/transform/fib.py @@ -20,11 +20,11 @@ sys.path.insert(0, str(Path(__file__).parent.parent)) +import cuda.bench as bench import cupy as cp import numpy as np from utils import as_cupy_stream, generate_data_with_entropy -import cuda.bench as bench from cuda.compute import make_unary_transform diff --git a/python/cuda_cccl/benchmarks/compute/transform/fill.py b/python/cuda_cccl/benchmarks/compute/transform/fill.py index db735027488..48699950d50 100644 --- a/python/cuda_cccl/benchmarks/compute/transform/fill.py +++ b/python/cuda_cccl/benchmarks/compute/transform/fill.py @@ -16,11 +16,11 @@ sys.path.insert(0, str(Path(__file__).parent.parent)) +import cuda.bench as bench import cupy as cp from utils import INTEGRAL_TYPES as TYPE_MAP from utils import as_cupy_stream -import cuda.bench as bench import cuda.compute from cuda.compute import ConstantIterator, OpKind diff --git a/python/cuda_cccl/benchmarks/compute/transform/grayscale.py b/python/cuda_cccl/benchmarks/compute/transform/grayscale.py index e589e8c06e6..1b5856f6eac 100644 --- a/python/cuda_cccl/benchmarks/compute/transform/grayscale.py +++ b/python/cuda_cccl/benchmarks/compute/transform/grayscale.py @@ -19,11 +19,11 @@ sys.path.insert(0, str(Path(__file__).parent.parent)) +import cuda.bench as bench import cupy as cp from utils import FLOAT_TYPES as TYPE_MAP from utils import as_cupy_stream, generate_data_with_entropy -import cuda.bench as bench from cuda.compute import gpu_struct, make_unary_transform diff --git a/python/cuda_cccl/benchmarks/compute/transform/heavy.py b/python/cuda_cccl/benchmarks/compute/transform/heavy.py index 27ba1781922..be5af05efdd 100644 --- a/python/cuda_cccl/benchmarks/compute/transform/heavy.py +++ b/python/cuda_cccl/benchmarks/compute/transform/heavy.py @@ -15,13 +15,13 @@ sys.path.insert(0, str(Path(__file__).parent.parent)) +import cuda.bench as bench import cupy as cp import numba import numpy as np from numba import cuda as lang from utils import as_cupy_stream, generate_data_with_entropy -import cuda.bench as bench import cuda.compute diff --git a/python/cuda_cccl/benchmarks/compute/transform_reduce/sum.py b/python/cuda_cccl/benchmarks/compute/transform_reduce/sum.py index 6e04d3eb3b3..fd0976adbaa 100644 --- a/python/cuda_cccl/benchmarks/compute/transform_reduce/sum.py +++ b/python/cuda_cccl/benchmarks/compute/transform_reduce/sum.py @@ -18,12 +18,12 @@ sys.path.insert(0, str(Path(__file__).parent.parent)) +import cuda.bench as bench import cupy as cp import numpy as np from utils import SIGNED_TYPES as TYPE_MAP from utils import as_cupy_stream, generate_data_with_entropy -import cuda.bench as bench from cuda.compute import OpKind, TransformIterator, make_reduce_into diff --git a/python/cuda_cccl/benchmarks/compute/utils.py b/python/cuda_cccl/benchmarks/compute/utils.py index 3b163db4065..8814a08714a 100644 --- a/python/cuda_cccl/benchmarks/compute/utils.py +++ b/python/cuda_cccl/benchmarks/compute/utils.py @@ -1,8 +1,7 @@ +import cuda.bench as bench import cupy as cp import numpy as np -import cuda.bench as bench - ALL_TYPES = { "I8": np.int8, "I16": np.int16, diff --git a/python/cuda_cccl/benchmarks/coop/_experimental/bench_warp_reduce.py b/python/cuda_cccl/benchmarks/coop/_experimental/bench_warp_reduce.py index 25fa84a19e8..97e34d1a436 100644 --- a/python/cuda_cccl/benchmarks/coop/_experimental/bench_warp_reduce.py +++ b/python/cuda_cccl/benchmarks/coop/_experimental/bench_warp_reduce.py @@ -5,12 +5,11 @@ import os import sys +import cuda.bench as bench import numba import numpy as np from numba import cuda -import cuda.bench as bench - sys.path.insert(0, os.path.abspath(os.path.dirname(__file__))) from device_side_benchmark import ( # isort: skip # type: ignore[import-not-found] # noqa: E402 make_unrolled_kernel, diff --git a/python/cuda_cccl/benchmarks/coop/_experimental/device_side_benchmark.py b/python/cuda_cccl/benchmarks/coop/_experimental/device_side_benchmark.py index 64cadd95f5f..4d4a5048bda 100644 --- a/python/cuda_cccl/benchmarks/coop/_experimental/device_side_benchmark.py +++ b/python/cuda_cccl/benchmarks/coop/_experimental/device_side_benchmark.py @@ -4,13 +4,13 @@ import math +import cuda.bindings.driver as driver +from cuda.core import Device from llvmlite import ir from numba import cuda, types from numba.core.extending import intrinsic -import cuda.bindings.driver as driver import cuda.coop._experimental as coop -from cuda.core import Device @intrinsic diff --git a/python/cuda_cccl/cuda/compute/_bindings.py b/python/cuda_cccl/cuda/compute/_bindings.py index 57d21c6cd6c..541cc41f238 100644 --- a/python/cuda_cccl/cuda/compute/_bindings.py +++ b/python/cuda_cccl/cuda/compute/_bindings.py @@ -27,11 +27,12 @@ import importlib import os -from cuda.cccl._cuda_version_utils import detect_cuda_version, get_recommended_extra from cuda.pathfinder import ( # type: ignore[import-not-found] load_nvidia_dynamic_lib, ) +from cuda.cccl._cuda_version_utils import detect_cuda_version, get_recommended_extra + def _load_cuda_libraries(): # Load appropriate libraries for the detected CUDA version diff --git a/python/cuda_cccl/cuda/compute/_cpp_compile.py b/python/cuda_cccl/cuda/compute/_cpp_compile.py index 87916ca7bb2..bc21e127965 100644 --- a/python/cuda_cccl/cuda/compute/_cpp_compile.py +++ b/python/cuda_cccl/cuda/compute/_cpp_compile.py @@ -10,9 +10,10 @@ import functools -from cuda.cccl import get_include_paths from cuda.core import Device, Program, ProgramOptions +from cuda.cccl import get_include_paths + from ._bindings import TypeEnum from ._device_code import DeviceCode diff --git a/python/cuda_cccl/tests/_utils/device_array.py b/python/cuda_cccl/tests/_utils/device_array.py index b2496173a28..b4124ee68d2 100644 --- a/python/cuda_cccl/tests/_utils/device_array.py +++ b/python/cuda_cccl/tests/_utils/device_array.py @@ -9,9 +9,8 @@ from collections.abc import Iterable import numpy as np -from numpy.typing import DTypeLike - from cuda.core import Buffer, Device, Stream +from numpy.typing import DTypeLike def get_compute_capability() -> tuple[int, int]: diff --git a/python/cuda_cccl/tests/compute/conftest.py b/python/cuda_cccl/tests/compute/conftest.py index 57c73ed7de8..701735686b6 100644 --- a/python/cuda_cccl/tests/compute/conftest.py +++ b/python/cuda_cccl/tests/compute/conftest.py @@ -3,7 +3,6 @@ import numpy as np import pytest - from cuda.core import Device, Stream try: diff --git a/python/cuda_cccl/tests/compute/examples/raw_op/cpp_stateful.py b/python/cuda_cccl/tests/compute/examples/raw_op/cpp_stateful.py index 970498e332e..4da7d530576 100644 --- a/python/cuda_cccl/tests/compute/examples/raw_op/cpp_stateful.py +++ b/python/cuda_cccl/tests/compute/examples/raw_op/cpp_stateful.py @@ -15,10 +15,10 @@ import cupy as cp import numpy as np +from cuda.core import Device, Program, ProgramOptions import cuda.compute from cuda.compute.op import RawOp -from cuda.core import Device, Program, ProgramOptions def get_arch(): diff --git a/python/cuda_cccl/tests/compute/examples/raw_op/cpp_stateless.py b/python/cuda_cccl/tests/compute/examples/raw_op/cpp_stateless.py index 2db4e09e2d5..a15035338e9 100644 --- a/python/cuda_cccl/tests/compute/examples/raw_op/cpp_stateless.py +++ b/python/cuda_cccl/tests/compute/examples/raw_op/cpp_stateless.py @@ -12,10 +12,10 @@ import cupy as cp import numpy as np +from cuda.core import Device, Program, ProgramOptions import cuda.compute from cuda.compute.op import RawOp -from cuda.core import Device, Program, ProgramOptions def get_arch(): diff --git a/python/cuda_cccl/tests/compute/test_multi_cc_serialization.py b/python/cuda_cccl/tests/compute/test_multi_cc_serialization.py index 8c326d305df..27b40eb3989 100644 --- a/python/cuda_cccl/tests/compute/test_multi_cc_serialization.py +++ b/python/cuda_cccl/tests/compute/test_multi_cc_serialization.py @@ -19,6 +19,7 @@ import numpy as np import pytest from _utils.device_array import DeviceArray +from cuda.core import Device from cuda.compute import ( ProxyArray, @@ -37,7 +38,6 @@ from cuda.compute._target_cc import target_cc from cuda.compute.iterators import CountingIterator, TransformIterator from cuda.compute.types import from_numpy_dtype -from cuda.core import Device try: from cuda.compute._build_info import USING_V2 diff --git a/python/cuda_cccl/tests/compute/test_raw_op.py b/python/cuda_cccl/tests/compute/test_raw_op.py index b8544658a01..d2f7ca3615b 100644 --- a/python/cuda_cccl/tests/compute/test_raw_op.py +++ b/python/cuda_cccl/tests/compute/test_raw_op.py @@ -9,12 +9,12 @@ import numpy as np import pytest from _utils.device_array import DeviceArray, get_compute_capability +from cuda.core import Program, ProgramOptions import cuda.compute from cuda.compute import types from cuda.compute._cpp_compile import _get_include_paths from cuda.compute.op import RawOp -from cuda.core import Program, ProgramOptions # Mark all tests in this module as no_numba pytestmark = pytest.mark.no_numba diff --git a/thrust/testing/cuda/transform_iterator.cu b/thrust/testing/cuda/transform_iterator.cu index 2973b5d569f..551d718f766 100644 --- a/thrust/testing/cuda/transform_iterator.cu +++ b/thrust/testing/cuda/transform_iterator.cu @@ -8,7 +8,9 @@ // see also: https://github.com/NVIDIA/cccl/issues/3541 void TestTransformWithLambda() { - auto l = [] __host__ __device__(int v) { return v < 4; }; + auto l = [] __host__ __device__(int v) { + return v < 4; + }; thrust::host_vector A{1, 2, 3, 4, 5, 6, 7}; ASSERT_EQUAL(thrust::any_of(A.begin(), A.end(), l), true); diff --git a/thrust/testing/dereference.cu b/thrust/testing/dereference.cu index 9a65b138243..03288f99127 100644 --- a/thrust/testing/dereference.cu +++ b/thrust/testing/dereference.cu @@ -13,8 +13,7 @@ template #if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA __global__ #endif - void - simple_copy_on_device(Iterator1 first1, Iterator1 last1, Iterator2 first2) + void simple_copy_on_device(Iterator1 first1, Iterator1 last1, Iterator2 first2) { while (first1 != last1) { diff --git a/thrust/testing/unittest/special_types.h b/thrust/testing/unittest/special_types.h index 8b10a101480..87697632346 100644 --- a/thrust/testing/unittest/special_types.h +++ b/thrust/testing/unittest/special_types.h @@ -29,8 +29,7 @@ struct FixedVector #if _CCCL_COMPILER(NVHPC) __attribute__((noinline)) #endif - FixedVector - operator+(const FixedVector& bs) const + FixedVector operator+(const FixedVector& bs) const { FixedVector output; for (unsigned int i = 0; i < N; i++) diff --git a/thrust/testing/unittest/testframework.h b/thrust/testing/unittest/testframework.h index 04c9b57a070..6e7facc7c21 100644 --- a/thrust/testing/unittest/testframework.h +++ b/thrust/testing/unittest/testframework.h @@ -149,11 +149,11 @@ class custom_numeric #define CONCAT(X, Y) X##Y -#define DEFINE_OPERATOR(op) \ - _CCCL_HOST_DEVICE constexpr custom_numeric& operator CONCAT(op, =)(const custom_numeric & other) \ - { \ - fill(value[0] op other.value[0]); \ - return *this; \ +#define DEFINE_OPERATOR(op) \ + _CCCL_HOST_DEVICE constexpr custom_numeric& operator CONCAT(op, =)(const custom_numeric& other) \ + { \ + fill(value[0] op other.value[0]); \ + return *this; \ } DEFINE_OPERATOR(+) diff --git a/thrust/testing/zip_iterator.cu b/thrust/testing/zip_iterator.cu index 40c11cfe302..1758566b27c 100644 --- a/thrust/testing/zip_iterator.cu +++ b/thrust/testing/zip_iterator.cu @@ -107,8 +107,8 @@ void TestZipIteratorTraits() using traits = cuda::std::iterator_traits; using value_type = cuda::std::tuple; using reference = thrust::detail::tuple_of_iterator_references< - int&, - cuda::__transform_input_output_proxy, cuda::std::plus<>, short*>>; + int&, + cuda::__transform_input_output_proxy, cuda::std::plus<>, short*>>; static_assert(cuda::std::is_same_v); static_assert(cuda::std::is_same_v); diff --git a/thrust/thrust/iterator/constant_iterator.h b/thrust/thrust/iterator/constant_iterator.h index eeee0a4e2cf..a11ab31ec10 100644 --- a/thrust/thrust/iterator/constant_iterator.h +++ b/thrust/thrust/iterator/constant_iterator.h @@ -103,8 +103,8 @@ struct make_constant_iterator_base * \endverbatim */ template -class CCCL_DEPRECATED_BECAUSE("Use cuda::constant_iterator instead") constant_iterator - : public detail::make_constant_iterator_base::type +class CCCL_DEPRECATED_BECAUSE("Use cuda::constant_iterator instead") +constant_iterator : public detail::make_constant_iterator_base::type { //! \cond friend class iterator_core_access; diff --git a/thrust/thrust/iterator/iterator_facade.h b/thrust/thrust/iterator/iterator_facade.h index 9661e0e0236..c48beeb2e78 100644 --- a/thrust/thrust/iterator/iterator_facade.h +++ b/thrust/thrust/iterator/iterator_facade.h @@ -460,13 +460,11 @@ template -inline - _CCCL_HOST_DEVICE - // XXX it might be nice to implement this at some point - // typename enable_if_interoperable::type // exposition - bool - operator==(iterator_facade const& lhs, - iterator_facade const& rhs) +inline _CCCL_HOST_DEVICE +// XXX it might be nice to implement this at some point +// typename enable_if_interoperable::type // exposition +bool operator==(iterator_facade const& lhs, + iterator_facade const& rhs) { return iterator_core_access ::equal(*static_cast(&lhs), *static_cast(&rhs)); } @@ -483,13 +481,11 @@ template -inline - _CCCL_HOST_DEVICE - // XXX it might be nice to implement this at some point - // typename enable_if_interoperable::type // exposition - bool - operator!=(iterator_facade const& lhs, - iterator_facade const& rhs) +inline _CCCL_HOST_DEVICE +// XXX it might be nice to implement this at some point +// typename enable_if_interoperable::type // exposition +bool operator!=(iterator_facade const& lhs, + iterator_facade const& rhs) { return !iterator_core_access ::equal(*static_cast(&lhs), *static_cast(&rhs)); } @@ -506,13 +502,11 @@ template -inline - _CCCL_HOST_DEVICE - // XXX it might be nice to implement this at some point - // typename enable_if_interoperable::type // exposition - bool - operator<(iterator_facade const& lhs, - iterator_facade const& rhs) +inline _CCCL_HOST_DEVICE +// XXX it might be nice to implement this at some point +// typename enable_if_interoperable::type // exposition +bool operator<(iterator_facade const& lhs, + iterator_facade const& rhs) { return 0 > iterator_core_access ::distance_from(*static_cast(&lhs), *static_cast(&rhs)); @@ -530,13 +524,11 @@ template -inline - _CCCL_HOST_DEVICE - // XXX it might be nice to implement this at some point - // typename enable_if_interoperable::type // exposition - bool - operator>(iterator_facade const& lhs, - iterator_facade const& rhs) +inline _CCCL_HOST_DEVICE +// XXX it might be nice to implement this at some point +// typename enable_if_interoperable::type // exposition +bool operator>(iterator_facade const& lhs, + iterator_facade const& rhs) { return 0 < iterator_core_access ::distance_from(*static_cast(&lhs), *static_cast(&rhs)); @@ -554,13 +546,11 @@ template -inline - _CCCL_HOST_DEVICE - // XXX it might be nice to implement this at some point - // typename enable_if_interoperable::type // exposition - bool - operator<=(iterator_facade const& lhs, - iterator_facade const& rhs) +inline _CCCL_HOST_DEVICE +// XXX it might be nice to implement this at some point +// typename enable_if_interoperable::type // exposition +bool operator<=(iterator_facade const& lhs, + iterator_facade const& rhs) { return 0 >= iterator_core_access ::distance_from(*static_cast(&lhs), *static_cast(&rhs)); @@ -578,13 +568,11 @@ template -inline - _CCCL_HOST_DEVICE - // XXX it might be nice to implement this at some point - // typename enable_if_interoperable::type // exposition - bool - operator>=(iterator_facade const& lhs, - iterator_facade const& rhs) +inline _CCCL_HOST_DEVICE +// XXX it might be nice to implement this at some point +// typename enable_if_interoperable::type // exposition +bool operator>=(iterator_facade const& lhs, + iterator_facade const& rhs) { return 0 <= iterator_core_access ::distance_from(*static_cast(&lhs), *static_cast(&rhs)); diff --git a/thrust/thrust/iterator/permutation_iterator.h b/thrust/thrust/iterator/permutation_iterator.h index ce9160f11ee..363355d492d 100644 --- a/thrust/thrust/iterator/permutation_iterator.h +++ b/thrust/thrust/iterator/permutation_iterator.h @@ -115,8 +115,8 @@ struct make_permutation_iterator_base * \endverbatim */ template -class _CCCL_DECLSPEC_EMPTY_BASES permutation_iterator - : public detail::make_permutation_iterator_base::type +class _CCCL_DECLSPEC_EMPTY_BASES +permutation_iterator : public detail::make_permutation_iterator_base::type { //! \cond diff --git a/thrust/thrust/iterator/tabulate_output_iterator.h b/thrust/thrust/iterator/tabulate_output_iterator.h index bc2b4d3e26b..7dd6645f1a8 100644 --- a/thrust/thrust/iterator/tabulate_output_iterator.h +++ b/thrust/thrust/iterator/tabulate_output_iterator.h @@ -104,8 +104,8 @@ inline constexpr bool is_proxy_reference_v -class CCCL_DEPRECATED_BECAUSE("Use cuda::tabulate_output_iterator instead") tabulate_output_iterator - : public detail::make_tabulate_output_iterator_base +class CCCL_DEPRECATED_BECAUSE("Use cuda::tabulate_output_iterator instead") +tabulate_output_iterator : public detail::make_tabulate_output_iterator_base { public: //! \cond diff --git a/thrust/thrust/system/cuda/detail/util.h b/thrust/thrust/system/cuda/detail/util.h index b72fdbc7829..173dde68984 100644 --- a/thrust/thrust/system/cuda/detail/util.h +++ b/thrust/thrust/system/cuda/detail/util.h @@ -113,8 +113,8 @@ _CCCL_HOST_DEVICE cudaError_t synchronize_optional(Policy& policy) #if _CCCL_HOSTED() template -_CCCL_HOST_API _CCCL_FORCEINLINE cudaError_t -trivial_copy_from_device(Type* dst, Type const* src, size_t count, cudaStream_t stream) +_CCCL_HOST_API +_CCCL_FORCEINLINE cudaError_t trivial_copy_from_device(Type* dst, Type const* src, size_t count, cudaStream_t stream) { cudaError status = cudaSuccess; if (count == 0) From 1c5bce9b8a02f557641476e53955dd505e1c33b1 Mon Sep 17 00:00:00 2001 From: "pre-commit-ci[bot]" <66853113+pre-commit-ci[bot]@users.noreply.github.com> Date: Wed, 15 Jul 2026 15:29:13 +0000 Subject: [PATCH 3/3] [pre-commit.ci] auto code formatting --- .../device/dispatch/kernels/kernel_reduce_deterministic.cuh | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/cub/cub/device/dispatch/kernels/kernel_reduce_deterministic.cuh b/cub/cub/device/dispatch/kernels/kernel_reduce_deterministic.cuh index cb5b0891037..22513cde4f0 100644 --- a/cub/cub/device/dispatch/kernels/kernel_reduce_deterministic.cuh +++ b/cub/cub/device/dispatch/kernels/kernel_reduce_deterministic.cuh @@ -50,9 +50,9 @@ struct deterministic_num_items { using element_t = it_value_t; using type = ::cuda::std::conditional_t< - ::cuda::std::is_signed_v && sizeof(element_t) == 4, - int, - ::cuda::std::conditional_t<::cuda::std::is_signed_v, ::cuda::std::int64_t, ::cuda::std::uint64_t>>; + ::cuda::std::is_signed_v && sizeof(element_t) == 4, + int, + ::cuda::std::conditional_t<::cuda::std::is_signed_v, ::cuda::std::int64_t, ::cuda::std::uint64_t>>; }; template