diff --git a/.github/actions/workflow-build/build-workflow.py b/.github/actions/workflow-build/build-workflow.py index b2eb13489e6..d2730a4ee44 100755 --- a/.github/actions/workflow-build/build-workflow.py +++ b/.github/actions/workflow-build/build-workflow.py @@ -1267,7 +1267,7 @@ def process_job_array(group_name, array_name, parent_json): ) } - runner_heading = f"🏃‍ Runner counts (total jobs: {total_jobs})" + runner_heading = f"🏃 Runner counts (total jobs: {total_jobs})" runner_counts_table = f"| {'#':^4} | Runner\n" runner_counts_table += "|------|------\n" diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml index d281c6c3523..3928eef7021 100644 --- a/.pre-commit-config.yaml +++ b/.pre-commit-config.yaml @@ -94,6 +94,14 @@ repos: "python/cuda_cccl/cuda/compute/"] pass_filenames: false + - repo: https://github.com/sirosen/texthooks + rev: 0.7.1 + hooks: + - id: fix-smartquotes + - id: fix-spaces + - id: forbid-bidi-controls + - id: fix-ligatures + - repo: local hooks: - id: check-shebang @@ -106,5 +114,15 @@ repos: ^.*libcudacxx/cmake/config\.guess$ ) + - id: unprintable-unicode + name: unprintable-unicode + entry: ci/util/pre-commit/strip_unprintable.bash + language: unsupported_script + types: [text] + exclude: | + (?x)^( + ^.*ci/util/pre-commit/strip_unprintable\.bash$ + ) + default_language_version: python: python3 diff --git a/AGENTS.md b/AGENTS.md index f1131fa3bf2..2141892202c 100644 --- a/AGENTS.md +++ b/AGENTS.md @@ -334,7 +334,7 @@ CCCL's CI is built on GitHub Actions and relies on a dynamically generated job m * Declares build and test jobs for `pull_request`, `nightly`, and `weekly` workflows. * Pull request (PR) runs typically spawn ~250 jobs. * To reduce overhead, you can add an override matrix in `workflows.override`. This limits the PR CI run to a targeted subset of jobs. Overrides are recommended when: - * Changes touch high-dependency areas (e.g. top-level CI/devcontainers, libcudacxx, thrust, CUB). See `ci/inspect_changes.py` for dependency information. + * Changes touch high-dependency areas (e.g. top-level CI/devcontainers, libcudacxx, thrust, CUB). See `ci/inspect_changes.py` for dependency information. * A smaller subset of jobs is enough to validate the change (e.g. infra changes, targeted fixes). * Important rules: * PR merges are blocked while an override matrix is active. diff --git a/ci/util/pre-commit/strip_unprintable.bash b/ci/util/pre-commit/strip_unprintable.bash new file mode 100755 index 00000000000..2e4103ca25c --- /dev/null +++ b/ci/util/pre-commit/strip_unprintable.bash @@ -0,0 +1,100 @@ +#!/usr/bin/env bash +# +# strip_unprintable.bash - Remove invisible / unprintable characters from text files. + +set -euo pipefail + +PROG=$(basename "$0") + +# Canonical definition of what gets removed. One row per range group: +# +# The Perl class and the --help listing are both derived from this table, so +# adding or removing a range only needs to happen here. +ranges() { + printf '%s\t%s\n' \ + '\x{0000}-\x{0008}\x{000B}\x{000C}\x{000E}-\x{001F}\x{007F}' 'C0 controls / DEL (TAB, LF, CR preserved)' \ + '\x{0080}-\x{009F}' 'C1 controls' \ + '\x{00A0}' 'no-break space' \ + '\x{200B}-\x{200F}' 'zero-width space/joiners, bidi marks' \ + '\x{202A}-\x{202E}' 'bidi embedding/override' \ + '\x{2060}-\x{2064}' 'word joiner, invisible operators' \ + '\x{FEFF}' 'BOM / zero-width no-break space' +} + +# Perl character class assembled from column 1 of the ranges table. +BAD_CLASS="[$(ranges | cut -f1 | tr -d '\n')]" +PERL="${PERL:-perl}" + +usage() { + cat <&2 + usage >&2 + exit 2 + ;; + *) + break + ;; + esac +done + +if [[ $# -eq 0 ]]; then + echo "error: no files given" >&2 + usage >&2 + exit 2 +fi + +status=0 +# Both modes process every file in a single perl process to avoid per-file interpreter +# startup overhead. In --check mode, `close ARGV if eof` resets the line counter ($.) at +# each file boundary so reported line numbers are per-file. +if [[ ${check_only} -eq 1 ]]; then + # $cls/$ARGV/$./$& are Perl variables and must stay literal inside the single quotes; + # only BAD_CLASS is interpolated, via the deliberate quote break-out. + # + # shellcheck disable=SC2016 + ${PERL} -CSD -ne ' + BEGIN { $cls = qr/'"${BAD_CLASS}"'/; } + while (/$cls/g) { + printf "%s:%d:%d: U+%04X\n", $ARGV, $., pos() - length($&) + 1, ord($&); + $found = 1; + } + close ARGV if eof; + END { exit($found ? 1 : 0); } + ' "$@" || status=1 + exit "${status}" +fi + +${PERL} -CSD -i -pe 's/'"${BAD_CLASS}"'//g' "$@" diff --git a/ci/windows/build_common.psm1 b/ci/windows/build_common.psm1 index 4d355935016..bb7288b4247 100644 --- a/ci/windows/build_common.psm1 +++ b/ci/windows/build_common.psm1 @@ -1,4 +1,4 @@ -Param( +Param( [Parameter(Mandatory = $false)] [Alias("std")] [ValidateNotNullOrEmpty()] diff --git a/ci/windows/build_cuda_cccl_python.ps1 b/ci/windows/build_cuda_cccl_python.ps1 index 3d6dadd632a..00684ceadb5 100644 --- a/ci/windows/build_cuda_cccl_python.ps1 +++ b/ci/windows/build_cuda_cccl_python.ps1 @@ -28,8 +28,8 @@ When set, only that version is built and the *merge* step is skipped. .PARAMETER Cuda13Image - Optional. The Docker image name used for a nested build of the CUDA 13 - wheel when the outer container defaults to CUDA 12.9. The default value + Optional. The Docker image name used for a nested build of the CUDA 13 + wheel when the outer container defaults to CUDA 12.9. The default value matches the RAPIDS dev‑container image that contains the required toolchain: `rapidsai/devcontainers:26.06-cuda13.0-cl14.44-windows2022`. @@ -39,7 +39,7 @@ Action. .EXAMPLE - # Build a single cuda-cccl wheel for Python 3.13 (consisting of both CUDA + # Build a single cuda-cccl wheel for Python 3.13 (consisting of both CUDA # 12 and 13 versions), and, if in CI, upload the resulting wheel as an # artifact. .\build_cuda_cccl_python.ps1 -PyVersion 3.11 @@ -155,7 +155,7 @@ ${null} = New-Item -ItemType Directory -Path (Join-Path $RepoRoot 'wheelhouse_cu function Invoke-Cuda13NestedBuild { <# .SYNOPSIS - Run the nested Docker build for CUDA 13 when we are already inside a + Run the nested Docker build for CUDA 13 when we are already inside a CUDA 12 builder image. .DESCRIPTION diff --git a/ci/windows/run_cpu_bisect.ps1 b/ci/windows/run_cpu_bisect.ps1 index 40d18306546..911d38d98b8 100644 --- a/ci/windows/run_cpu_bisect.ps1 +++ b/ci/windows/run_cpu_bisect.ps1 @@ -1,4 +1,4 @@ -Param( +Param( [Parameter(ValueFromRemainingArguments = $true)] [string[]]$PassthroughArgs ) diff --git a/ci/windows/run_cpu_target.ps1 b/ci/windows/run_cpu_target.ps1 index 565bf8f288f..db92859514a 100644 --- a/ci/windows/run_cpu_target.ps1 +++ b/ci/windows/run_cpu_target.ps1 @@ -1,4 +1,4 @@ -Param( +Param( [Parameter(ValueFromRemainingArguments = $true)] [string[]]$PassthroughArgs ) diff --git a/ci/windows/run_gpu_bisect.ps1 b/ci/windows/run_gpu_bisect.ps1 index 40d18306546..911d38d98b8 100644 --- a/ci/windows/run_gpu_bisect.ps1 +++ b/ci/windows/run_gpu_bisect.ps1 @@ -1,4 +1,4 @@ -Param( +Param( [Parameter(ValueFromRemainingArguments = $true)] [string[]]$PassthroughArgs ) diff --git a/ci/windows/run_gpu_target.ps1 b/ci/windows/run_gpu_target.ps1 index 565bf8f288f..db92859514a 100644 --- a/ci/windows/run_gpu_target.ps1 +++ b/ci/windows/run_gpu_target.ps1 @@ -1,4 +1,4 @@ -Param( +Param( [Parameter(ValueFromRemainingArguments = $true)] [string[]]$PassthroughArgs ) diff --git a/cudax/include/cuda/experimental/__stf/internal/acquire_release.cuh b/cudax/include/cuda/experimental/__stf/internal/acquire_release.cuh index d3efc1acd10..b1bb31d43bc 100644 --- a/cudax/include/cuda/experimental/__stf/internal/acquire_release.cuh +++ b/cudax/include/cuda/experimental/__stf/internal/acquire_release.cuh @@ -180,7 +180,7 @@ inline event_list task::acquire(backend_ctx_untyped& ctx) result.merge(ctx.get_start_events()); } - // @@@@ TODO@@@@ explain this algorithm, and why we need to go in reversed + // @@@@ TODO@@@@ explain this algorithm, and why we need to go in reversed // order because we skipped equivalent dependencies that were stored // contiguously after sorting. instance_id_t previous_instance_id = instance_id_t::invalid; @@ -193,7 +193,7 @@ inline event_list task::acquire(backend_ctx_untyped& ctx) else { assert(previous_instance_id != instance_id_t::invalid); - // @@@@ TODO @@@@ make a unit test to make sure we have the same instance id for different acquired_data + // @@@@ TODO @@@@ make a unit test to make sure we have the same instance id for different acquired_data // ? fprintf(stderr, "SETTING SKIPPED INSTANCE ID ... %d\n", previous_instance_id); it->set_instance_id(previous_instance_id); } diff --git a/cudax/test/stf/examples/05-stencil-no-copy.cu b/cudax/test/stf/examples/05-stencil-no-copy.cu index babb07852db..acc19de53c3 100644 --- a/cudax/test/stf/examples/05-stencil-no-copy.cu +++ b/cudax/test/stf/examples/05-stencil-no-copy.cu @@ -14,7 +14,7 @@ using namespace cuda::experimental::stf; /* * DATA BLOCKS - * | GHOSTS | DATA | GHOSTS | + * | GHOSTS | DATA | GHOSTS | */ template class data_block diff --git a/cudax/test/stf/examples/05-stencil.cu b/cudax/test/stf/examples/05-stencil.cu index ec9a06b8f7a..a9a6e7ec935 100644 --- a/cudax/test/stf/examples/05-stencil.cu +++ b/cudax/test/stf/examples/05-stencil.cu @@ -16,7 +16,7 @@ static stream_ctx ctx; /* * DATA BLOCKS - * | GHOSTS | DATA | GHOSTS | + * | GHOSTS | DATA | GHOSTS | */ template class data_block diff --git a/cudax/test/stf/stencil/stencil-1D.cu b/cudax/test/stf/stencil/stencil-1D.cu index fc6ae751a9f..0c1b2976dbc 100644 --- a/cudax/test/stf/stencil/stencil-1D.cu +++ b/cudax/test/stf/stencil/stencil-1D.cu @@ -18,7 +18,7 @@ static graph_ctx ctx; /* * DATA BLOCKS - * | GHOSTS | DATA | GHOSTS | + * | GHOSTS | DATA | GHOSTS | */ template class data_block diff --git a/docs/cub/tuning.rst b/docs/cub/tuning.rst index 212b5eecda0..4c0de01cf7c 100644 --- a/docs/cub/tuning.rst +++ b/docs/cub/tuning.rst @@ -542,7 +542,7 @@ since the names of the reported parameters in the variant are derived from these The variant ``ipt_19.tpb_512``, which stands for 19 items per thread (``ipt``) and 512 threads per block (``tpb``), was thus compiled with ``-DTUNE_ITEMS_PER_THREAD=19 -DTUNE_THREADS_PER_BLOCK=512``. The meaning of these values is specific to the benchmark definition, -and we have to check the benchmark’s source code for how they are applied. +and we have to check the benchmark's source code for how they are applied. Equally named tuning parameters may not translate to different benchmarks (please double check). As a user of CUB, such a new set of tuning parameters (i.e. a variant) can then be used to define a policy selector, @@ -561,7 +561,7 @@ as :ref:`sketched above `: } }; -The default tunings defined inside CUB’s source use the same infrastructure +The default tunings defined inside CUB's source use the same infrastructure but should only be changed and extended by the CCCL maintainers. All default tunings are found in the :code:`cub/device/detail/tuning/tuning_*.cuh` headers, organized by algorithm. CUB's policy selectors are highly parameterized on type information and traits of the input arguments to CUB algorithms diff --git a/docs/cudax/stf.rst b/docs/cudax/stf.rst index b919af2ab95..67163c0cda3 100644 --- a/docs/cudax/stf.rst +++ b/docs/cudax/stf.rst @@ -293,7 +293,7 @@ synchronization. Each ``stream_task`` in the ``stream_ctx`` backend represents a task that is associated with an input CUDA stream. Asynchronous work can be submitted in the body of the task using this input stream. Once the ``stream_task`` completes, all work submitted -within the task’s body is assumed to be synchronized with the associated +within the task's body is assumed to be synchronized with the associated stream. Users can query the stream associated to a ``stream_task`` using its @@ -322,7 +322,7 @@ Users can retrieve the graph associated to a ``graph_task`` by using its Logical data ------------ -In traditional computing, “data”, such as a matrix describing a neural +In traditional computing, "data", such as a matrix describing a neural network layer, typically refers to a location in memory with a defined address. However, in mixed CPU/GPU systems, the same conceptual data may exist simultaneously in multiple locations and have multiple addresses @@ -331,7 +331,7 @@ high-bandwidth memory used by GPUs). CUDASTF refers to such conceptual data as *logical data*, an abstract handle for data that may get transparently transferred to or replicated over the different places used by CUDASTF tasks. When user code creates a logical data object from -a user-provided object (e.g. an array of ``double``), they transfer the +a user-provided object (e.g. an array of ``double``), they transfer the ownership of the original data to CUDASTF. As a result, any access to the original data should be performed through the logical data interface, as CUDASTF may transfer the logical data to a CUDA device @@ -411,7 +411,7 @@ Write-back policy When a logical data object is destroyed, the original data instance is updated (unless the logical data was created without a reference value, -e.g. from a shape). The result is only guaranteed to be available on the +e.g. from a shape). The result is only guaranteed to be available on the corresponding data place when after the ``finalize()`` method was called on the context. Likewise, when calling ``finalize()`` a write-back mechanism is automatically issued on all logical data associated to the @@ -428,7 +428,7 @@ Slices To facilitate the use of potentially non-contiguous multi-dimensional arrays, we have introduced a C++ data structure class called ``slice``. -A slice is a partial specialization of C++’s +A slice is a partial specialization of C++'s ``std::mdspan`` (or ``std::experimental::mdspan`` depending on the C++ revision). .. code:: cpp @@ -608,7 +608,7 @@ needs to be updated so it uses a read-write access mode. The object returned by the call ``ctx.task()`` overloads ``operator->*()`` to accept a lambda function on the right-hand side. -This makes it easy for user code to pass the task’s body to the context +This makes it easy for user code to pass the task's body to the context with a syntax akin to a control flow statement. The first argument of the lambda function is a ``cudaStream_t`` that can be used to submit work asynchronously on the selected device within the body of the task. @@ -754,7 +754,7 @@ CUDASTF transparently handles data management (allocations, transfers, …), there can be outstanding asynchronous operations that were not submitted explicitly by the user. Therefore it is not sufficient to use native CUDA synchronization operations because they are not aware of -CUDASTF’s state. Client code must call ``ctx.finalize()`` instead of +CUDASTF's state. Client code must call ``ctx.finalize()`` instead of ``cudaStreamSynchronize()`` or ``cudaDeviceSynchronize()``. - ``ctx.submit()`` initiates the submission of all asynchronous tasks @@ -767,7 +767,7 @@ Usually, creating the task and invoking ``ctx.finalize()`` is sufficient. However, manually calling ``ctx.submit()`` can be beneficial in at least two situations. First, it allows for executing additional unrelated work on the CPU (or another GPU) between submission and -synchronization. Second, when it’s necessary for two contexts to run +synchronization. Second, when it's necessary for two contexts to run concurrently, using the sequence ``ctx1.submit(); ctx2.submit(); ctx1.finalize(); ctx2.finalize();`` achieves this goal (whereas calling @@ -801,7 +801,7 @@ Places CUDASTF uses :ref:`places ` to manage data and execution affinity. Places can represent either *execution places*, which determine where code is executed, or *data places*, specifying the location of data -across the machine’s non-uniform memory. See :ref:`cudax-places` for the +across the machine's non-uniform memory. See :ref:`cudax-places` for the full places API reference, including stream management, memory allocation, grid of places, and partitioning policies. @@ -810,7 +810,7 @@ This section describes how places are used with CUDASTF tasks. Execution places in tasks ^^^^^^^^^^^^^^^^^^^^^^^^^ -A task’s constructor allows choosing an execution place. The example +A task's constructor allows choosing an execution place. The example below creates a logical data variable that describes an integer as a vector of one ``int``. The logical data variable is then updated on device ``0`` and on device ``1`` before being accessed again from the @@ -823,7 +823,7 @@ and tells CUDASTF where the task is expected to execute. the host. Regardless of the *execution place*, it is important to note that the -task’s body (i.e., the contents of the lambda function) corresponds to +task's body (i.e., the contents of the lambda function) corresponds to CPU code that is expected to launch computation asynchronously. When using ``exec_place::device(id)``, CUDASTF will automatically set the current CUDA device to ``id`` when the task is started, and restore the @@ -1577,7 +1577,7 @@ Note the use of the ``mutable`` qualifier because a task accessing a ``const foo`` object might want to read the ``ldata`` field. Submitting a task that use this logical data in read only mode would modify the internal data structures of the logical data, but should probably appear -as a ``const`` operation from user’s perspective. Without this ``mutable`` +as a ``const`` operation from user's perspective. Without this ``mutable`` qualifier, we could not have a ``const`` qualifier on the ``f`` variable in the following code : @@ -1621,7 +1621,7 @@ compilation. ... }; -In most cases, it’s recommended to use the ``auto`` C++ keyword to +In most cases, it's recommended to use the ``auto`` C++ keyword to automatically obtain the correct data types: .. code:: cpp @@ -1823,7 +1823,7 @@ task. Since the token is only used for synchronization purposes, the corresponding argument may be omitted in the lambda function passed as the -task’s implementation. Thus, the above task is equivalent to this code: +task's implementation. Thus, the above task is equivalent to this code: .. code:: cpp @@ -2069,7 +2069,7 @@ would all have the same name. One possible work-around is to let ``ncu`` rename kernels accordingly to ``NVTX`` annotations. To do so, a symbol must be associated to the ``ctx.parallel_for`` and ``ctx.launch`` constructs using the ``set_symbol`` method. In the following example, we -name the generated kernel “updateA” : +name the generated kernel "updateA" : .. code:: cpp diff --git a/docs/libcudacxx/extended_api/memory_model.rst b/docs/libcudacxx/extended_api/memory_model.rst index 9673625d5b0..261e935efcd 100644 --- a/docs/libcudacxx/extended_api/memory_model.rst +++ b/docs/libcudacxx/extended_api/memory_model.rst @@ -211,11 +211,11 @@ leads to a **data race**, and exhibits **undefined behavior**: assert(x == 42); While the memory operations on ``f`` - the store and the loads - are -atomic, the scope of the store operation is “block scope”. Since the +atomic, the scope of the store operation is "block scope". Since the store is performed by Thread 0 of Block 0, it only includes all other threads of Block 0. However, the thread doing the loads is in Block 1, i.e., it is not in a scope included by the store operation performed in -Block 0, causing the store and the load to not be “atomic”, and +Block 0, causing the store and the load to not be "atomic", and introducing a data-race. For more examples see the `PTX memory consistency model litmus diff --git a/docs/libcudacxx/extended_api/synchronization_primitives/atomic_ref.rst b/docs/libcudacxx/extended_api/synchronization_primitives/atomic_ref.rst index 6a884f56214..dc28a9e7be3 100644 --- a/docs/libcudacxx/extended_api/synchronization_primitives/atomic_ref.rst +++ b/docs/libcudacxx/extended_api/synchronization_primitives/atomic_ref.rst @@ -47,7 +47,7 @@ No object or subobject of an object referenced by an ``atomic_­ref`` shall be c For ``cuda::atomic_ref`` and ``cuda::std::atomic_ref`` the type ``T`` must satisfy the following: - ``sizeof(T) <= 16``. - The referenced object must be aligned to its size: ``alignof(T) == sizeof(T)``. - - ``T`` must not have “padding bits”, i.e., T's `object representation `_ + - ``T`` must not have "padding bits", i.e., T's `object representation `_ must not have bits that do not participate in it's value representation. Concurrency Restrictions diff --git a/docs/libcudacxx/extended_api/synchronization_primitives/barrier/barrier_native_handle.rst b/docs/libcudacxx/extended_api/synchronization_primitives/barrier/barrier_native_handle.rst index ab99fae7e49..d36a63f23e2 100644 --- a/docs/libcudacxx/extended_api/synchronization_primitives/barrier/barrier_native_handle.rst +++ b/docs/libcudacxx/extended_api/synchronization_primitives/barrier/barrier_native_handle.rst @@ -22,7 +22,7 @@ If ``bar`` is not in ``__shared__`` memory, the behavior is undefined. Return Value ------------ -A pointer to the PTX “mbarrier” subobject of the :ref:`cuda::barrier ` +A pointer to the PTX "mbarrier" subobject of the :ref:`cuda::barrier ` object. Example diff --git a/docs/libcudacxx/index.rst b/docs/libcudacxx/index.rst index d6bcab993fe..846333d3020 100644 --- a/docs/libcudacxx/index.rst +++ b/docs/libcudacxx/index.rst @@ -27,8 +27,8 @@ Specifically, ``libcu++`` provides: C++ Standard Library Features ----------------------------- -If you are a C++ developer, then you know the C++ Standard Library (`sometimes referred to as “The -STL” `_) +If you are a C++ developer, then you know the C++ Standard Library (`sometimes referred to as "The +STL" `_) as what comes along with your compiler and provides things like ``std::string`` or ``std::vector`` or ``std::atomic``. It provides the fundamental abstractions that C++ developers need to build high quality applications and libraries. @@ -67,7 +67,7 @@ learning curve of learning CUDA. However, there are many aspects of writing high be expressed through purely Standard conforming APIs. For these cases, libcu++ also provides *extensions* of Standard Library utilities. -For example, libcu++ extends ``atomic`` and other synchronization primitives with the notion of a “thread scope” +For example, libcu++ extends ``atomic`` and other synchronization primitives with the notion of a "thread scope" that controls the strength of the memory fence. To use utilities that are extensions to Standard Library features, drop the ``std``: diff --git a/libcudacxx/include/cuda/std/__iterator/iterator_traits.h b/libcudacxx/include/cuda/std/__iterator/iterator_traits.h index 974db1c41ab..8dd0e3665b4 100644 --- a/libcudacxx/include/cuda/std/__iterator/iterator_traits.h +++ b/libcudacxx/include/cuda/std/__iterator/iterator_traits.h @@ -304,7 +304,7 @@ _CCCL_CONCEPT __cpp17_random_access_iterator = _CCCL_REQUIRES_EXPR((_Iter))( } // namespace __iterator_traits_detail // [iterator.traits]#3.1 -// If the qualified-id I​::​pointer is valid and denotes a type, then ​pointer names that type; +// If the qualified-id I::pointer is valid and denotes a type, then pointer names that type; template _CCCL_API auto __iterator_traits_deduce_member_pointer_or_void(int) -> typename _Iter::pointer; // Otherwise, it names void. @@ -317,12 +317,12 @@ using __iterator_traits_member_pointer_or_void = // [iterator.traits]#3.2 // [iterator.traits]#3.2.1 -// If the qualified-id I​::​pointer is valid and denotes a type, then pointer names that type. +// If the qualified-id I::pointer is valid and denotes a type, then pointer names that type. template _CCCL_API auto __iterator_traits_deduce_member_pointer_or_arrow_or_void(int, __priority_tag<1>) -> typename _Iter::pointer; -// Otherwise, if decltype(​declval().operator->()) is well-formed, then pointer names that type. +// Otherwise, if decltype(declval().operator->()) is well-formed, then pointer names that type. template _CCCL_API auto __iterator_traits_deduce_member_pointer_or_arrow_or_void(int, __priority_tag<0>) -> decltype(::cuda::std::declval<_Iter&>().operator->()); diff --git a/libcudacxx/include/cuda/std/__simd/concepts.h b/libcudacxx/include/cuda/std/__simd/concepts.h index 821206d4da1..c9eaa43316c 100644 --- a/libcudacxx/include/cuda/std/__simd/concepts.h +++ b/libcudacxx/include/cuda/std/__simd/concepts.h @@ -118,8 +118,8 @@ inline constexpr bool __is_constexpr_wrapper_value_preserving_v<_From, _ValueTyp // [simd.ctor] implicit value constructor // - From is not an arithmetic type and does not satisfy constexpr-wrapper-like, // - From is an arithmetic type and the conversion from From to value_type is value-preserving -// - From satisfies constexpr-wrapper-like, remove_cvref_t is an arithmetic type, and -// From​::​value is representable by value_type. +// - From satisfies constexpr-wrapper-like, remove_cvref_t is an arithmetic type, and +// From::value is representable by value_type. template > _CCCL_CONCEPT __is_value_ctor_implicit = convertible_to<_Up, _ValueType> diff --git a/libcudacxx/include/cuda/std/__simd/creation.h b/libcudacxx/include/cuda/std/__simd/creation.h index 46a98e7132e..4851c65ca2b 100644 --- a/libcudacxx/include/cuda/std/__simd/creation.h +++ b/libcudacxx/include/cuda/std/__simd/creation.h @@ -73,8 +73,8 @@ using __simd_size_seq = integer_sequence<__simd_size_type, _Ns...>; template <__simd_size_type _Np> using __make_simd_size_seq = make_integer_sequence<__simd_size_type, _Np>; -// "If basic_vec<​typename T​::​​value_type, Abi>​::​size() % T​::​size() is not 0, then -// resize_t​::​size() % T​::​size(), T> is valid and denotes +// "If basic_vec::size() % T::size() is not 0, then +// resize_t::size() % T::size(), T> is valid and denotes // a type." // // Vector: resize_t is valid if __deduce_abi_t is a specialized ABI type @@ -142,7 +142,7 @@ __make_chunk_tuple_element(const _Src& __src) noexcept } // Remainder case: return tuple<_SubVec, ..., _SubVec, _Tail> -// where _Tail is resize_t +// where _Tail is resize_t template [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr auto __make_chunk_tuple(const _Src& __src, __simd_size_seq<_Js...>) noexcept diff --git a/libcudacxx/include/cuda/std/__simd/permute.h b/libcudacxx/include/cuda/std/__simd/permute.h index 9b32f82c9d6..8e7a86f6e86 100644 --- a/libcudacxx/include/cuda/std/__simd/permute.h +++ b/libcudacxx/include/cuda/std/__simd/permute.h @@ -64,7 +64,7 @@ inline constexpr bool __idxmap_result_is_integral_v = : __idxmap_nargs_integral_v, void, __simd_size_type>; //---------------------------------------------------------------------------------------------------------------------- -// gen-fn: idxmap(i, V​::​size()) if that expression is well-formed, and idxmap(i) otherwise. +// gen-fn: idxmap(i, V::size()) if that expression is well-formed, and idxmap(i) otherwise. template inline constexpr bool __idxmap_invocable_two_args_v = @@ -203,7 +203,7 @@ permute(const basic_mask<_Bytes, _Abi>& __v, const basic_vec<_Up, _UAbi>& __indi // [simd.permute.mask] // A data-parallel object where the i-th element is initialized to the result of select-value(i) for all i in the range -// [0, V​::​size()). +// [0, V::size()). template struct __compress_generator @@ -242,7 +242,7 @@ __make_compress_generator(const _Vp& __v, const _Mp& __sel, typename _Vp::value_ } // A data-parallel object where the i-th element is initialized to the result of select-value(i) for all i in the range -// [0, V​::​size()) +// [0, V::size()) template struct __expand_generator diff --git a/libcudacxx/include/cuda/std/__tuple_dir/tuple_constraints.h b/libcudacxx/include/cuda/std/__tuple_dir/tuple_constraints.h index 5782a326748..f33d568bb89 100644 --- a/libcudacxx/include/cuda/std/__tuple_dir/tuple_constraints.h +++ b/libcudacxx/include/cuda/std/__tuple_dir/tuple_constraints.h @@ -276,7 +276,7 @@ __tuple_select_tuple_like_constructible(__tuple_types<_Types...>, __tuple_indice { using ::cuda::std::get; if constexpr (__is_cuda_std_ranges_subrange_v>) - { // [tuple#cnstr]-29.2: remove_cvref_t is not a specialization of ranges​::​subrange, + { // [tuple#cnstr]-29.2: remove_cvref_t is not a specialization of ranges::subrange, return __select_constructor::__none; } else if constexpr (is_same_v<_UTuple, const tuple<_Types...>&> || is_same_v<_UTuple, tuple<_Types...>&&>) @@ -294,9 +294,9 @@ __tuple_select_tuple_like_constructible(__tuple_types<_Types...>, __tuple_indice return __select_constructor::__none; } else if constexpr (!(is_constructible_v<_Types, decltype(get<_Indices>(::cuda::std::declval<_UTuple>()))> && ...)) - { // [tuple.cnstr]-21.2: is_constructible(std​::​forward(u)))>... is true - // [tuple.cnstr]-25.2: is_constructible(std​::​forward(u)))>... is true - // [tuple.cnstr]-29.4: is_constructible(std​::​forward(u)))>... is true + { // [tuple.cnstr]-21.2: is_constructible(std::forward(u)))>... is true + // [tuple.cnstr]-25.2: is_constructible(std::forward(u)))>... is true + // [tuple.cnstr]-29.4: is_constructible(std::forward(u)))>... is true return __select_constructor::__none; } else @@ -316,7 +316,7 @@ __tuple_select_tuple_like_constructible(__tuple_types<_Type>, __tuple_indices<_I { using ::cuda::std::get; if constexpr (__is_cuda_std_ranges_subrange_v>) - { // [tuple#cnstr]-29.2: remove_cvref_t is not a specialization of ranges​::​subrange, + { // [tuple#cnstr]-29.2: remove_cvref_t is not a specialization of ranges::subrange, return __select_constructor::__none; } else if constexpr (is_same_v<_UTuple, const tuple<_Type>&> || is_same_v<_UTuple, tuple<_Type>&&>) @@ -345,9 +345,9 @@ __tuple_select_tuple_like_constructible(__tuple_types<_Type>, __tuple_indices<_I return __select_constructor::__none; } else if constexpr (!is_constructible_v<_Type, decltype(get<_Index>(::cuda::std::declval<_UTuple>()))>) - { // [tuple.cnstr]-21.2: is_constructible(std​::​forward(u)))>... is true - // [tuple.cnstr]-25.2: is_constructible(std​::​forward(u)))>... is true - // [tuple.cnstr]-29.4: is_constructible(std​::​forward(u)))>... is true + { // [tuple.cnstr]-21.2: is_constructible(std::forward(u)))>... is true + // [tuple.cnstr]-25.2: is_constructible(std::forward(u)))>... is true + // [tuple.cnstr]-29.4: is_constructible(std::forward(u)))>... is true return __select_constructor::__none; } else @@ -538,7 +538,7 @@ __tuple_select_tuple_like_assignable(__tuple_types<_Types...>, __tuple_indices<_ return __select_assignment::__none; } else if constexpr (__is_cuda_std_ranges_subrange_v>) - { // [tuple.assign]-39.2: remove_cvref_t is not a specialization of ranges​::​subrange, + { // [tuple.assign]-39.2: remove_cvref_t is not a specialization of ranges::subrange, return __select_assignment::__none; } else if constexpr (!__tuple_like_with_size<_UTuple, sizeof...(_Types)>) @@ -548,7 +548,7 @@ __tuple_select_tuple_like_assignable(__tuple_types<_Types...>, __tuple_indices<_ else if constexpr (_IsConst) { if constexpr ((is_assignable_v(::cuda::std::declval<_UTuple>()))> && ...)) - { // [tuple.assign]-42.4: is_assignable_v(std​::​forward(u)))> is true for + { // [tuple.assign]-42.4: is_assignable_v(std::forward(u)))> is true for // all i return (is_nothrow_assignable_v(::cuda::std::declval<_UTuple>()))> && ...) ? __select_assignment::__is_nothrow @@ -560,7 +560,7 @@ __tuple_select_tuple_like_assignable(__tuple_types<_Types...>, __tuple_indices<_ } } else if constexpr ((is_assignable_v<_Types&, decltype(get<_Indices>(::cuda::std::declval<_UTuple>()))> && ...)) - { // [tuple.assign]-39.4: is_assignable_v(std​::​forward(u)))> is true for all i + { // [tuple.assign]-39.4: is_assignable_v(std::forward(u)))> is true for all i return (is_nothrow_assignable_v<_Types&, decltype(get<_Indices>(::cuda::std::declval<_UTuple>()))> && ...) ? __select_assignment::__is_nothrow : __select_assignment::__may_throw; diff --git a/libcudacxx/test/TODO.TXT b/libcudacxx/test/TODO.TXT index 652a38de75a..814f9907a16 100644 --- a/libcudacxx/test/TODO.TXT +++ b/libcudacxx/test/TODO.TXT @@ -27,19 +27,19 @@ Filesystem Tasks * INCOMPLETE - US 34: Are there attributes of a file that are not an aspect of the file system? * INCOMPLETE - US 35: What synchronization is required to avoid a file system race? * INCOMPLETE - US 36: Symbolic links themselves are attached to a directory via (hard) links - * INCOMPLETE - US 37: The term “redundant current directory (dot) elements” is not defined + * INCOMPLETE - US 37: The term "redundant current directory (dot) elements" is not defined * INCOMPLETE - US 38: Duplicates §17.3.16 * INCOMPLETE - US 39: Remove note: Dot and dot-dot are not directories * INCOMPLETE - US 40: Not all directories have a parent. - * INCOMPLETE - US 41: The term “parent directory” for a (non-directory) file is unusual + * INCOMPLETE - US 41: The term "parent directory" for a (non-directory) file is unusual * INCOMPLETE - US 42: Pathname resolution does not always resolve a symlink * INCOMPLETE - US 43: Concerns about encoded character types * INCOMPLETE - US 44: Definition of path in terms of a string requires leaky abstraction * INCOMPLETE - US 45: Generic format portability compromised by unspecified root-name * INCOMPLETE - US 46: filename can be empty so productions for relative-path are redundant - * INCOMPLETE - US 47: “.” and “..” already match the name production + * INCOMPLETE - US 47: "." and ".." already match the name production * INCOMPLETE - US 48: Multiple separators are often meaningful in a root-name - * INCOMPLETE - US 49: What does “method of conversion method” mean? + * INCOMPLETE - US 49: What does "method of conversion method" mean? * INCOMPLETE - US 50: 27.10.8.1 ¶ 1.4 largely redundant with ¶ 1.3 * INCOMPLETE - US 51: Failing to add / when appending empty string prevents useful apps * INCOMPLETE - US 52: remove_filename() postcondition is not by itself a definition @@ -55,7 +55,7 @@ Filesystem Tasks * INCOMPLETE - US 62: It is important that stem()+extension()==filename() * INCOMPLETE - US 63: lexically_normal() inconsistently treats trailing "/" but not "/.." as directory * INCOMPLETE - US 73, CA 2: root-name is effectively implementation defined - * INCOMPLETE - US 74, CA 3: The term “pathname” is ambiguous in some contexts + * INCOMPLETE - US 74, CA 3: The term "pathname" is ambiguous in some contexts * INCOMPLETE - US 75, CA 4: Extra flag in path constructors is needed * INCOMPLETE - US 76, CA 5: root-name definition is over-specified. * INCOMPLETE - US 77, CA 6: operator/ and other appends not useful if arg has root-name diff --git a/libcudacxx/test/libcudacxx/libcxx/iterators/compressed_movable_box/assign.copy.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/iterators/compressed_movable_box/assign.copy.pass.cpp index 5144b8cceee..4c4bfcb8a02 100644 --- a/libcudacxx/test/libcudacxx/libcxx/iterators/compressed_movable_box/assign.copy.pass.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/iterators/compressed_movable_box/assign.copy.pass.cpp @@ -186,7 +186,7 @@ int main(int, char**) test(); #if TEST_STD_VER >= 2020 # if !TEST_COMPILER(GCC, >=, 14) && !TEST_COMPILER(MSVC) - // GCC: error: destroying ‘b’ outside its lifetime (on the }) + // GCC: error: destroying 'b' outside its lifetime (on the }) // MSVC: error: read of an uninitialized symbol static_assert(test()); # endif // !TEST_COMPILER(GCC, >=, 14) && !TEST_COMPILER(MSVC) diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.nonmodifying/alg.for_each/ranges.for_each.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.nonmodifying/alg.for_each/ranges.for_each.pass.cpp index 77187cddd64..03f428e42eb 100644 --- a/libcudacxx/test/libcudacxx/std/algorithms/alg.nonmodifying/alg.for_each/ranges.for_each.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.nonmodifying/alg.for_each/ranges.for_each.pass.cpp @@ -208,7 +208,7 @@ TEST_FUNC constexpr bool test() int main(int, char**) { test(); -#if !TEST_COMPILER(GCC, <, 10) // accessing value of ‘a’ through a ‘int’ glvalue in a constant expression +#if !TEST_COMPILER(GCC, <, 10) // accessing value of 'a' through a 'int' glvalue in a constant expression static_assert(test()); #endif // !TEST_COMPILER(GCC, <, 10) diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.nonmodifying/alg.for_each/ranges.for_each_n.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.nonmodifying/alg.for_each/ranges.for_each_n.pass.cpp index 5fdc455260d..edacd9d842e 100644 --- a/libcudacxx/test/libcudacxx/std/algorithms/alg.nonmodifying/alg.for_each/ranges.for_each_n.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.nonmodifying/alg.for_each/ranges.for_each_n.pass.cpp @@ -143,7 +143,7 @@ TEST_FUNC constexpr bool test() int main(int, char**) { test(); -#if !TEST_COMPILER(GCC, <, 10) // accessing value of ‘a’ through a ‘int’ glvalue in a constant expression +#if !TEST_COMPILER(GCC, <, 10) // accessing value of 'a' through a 'int' glvalue in a constant expression static_assert(test()); #endif // !TEST_COMPILER(GCC, <, 10) diff --git a/libcudacxx/test/libcudacxx/std/iterators/iterator.primitives/iterator.traits/cxx20_iterator_traits.compile.pass.cpp b/libcudacxx/test/libcudacxx/std/iterators/iterator.primitives/iterator.traits/cxx20_iterator_traits.compile.pass.cpp index 01d5043de31..b273306aff2 100644 --- a/libcudacxx/test/libcudacxx/std/iterators/iterator.primitives/iterator.traits/cxx20_iterator_traits.compile.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/iterators/iterator.primitives/iterator.traits/cxx20_iterator_traits.compile.pass.cpp @@ -194,7 +194,7 @@ struct LegacyInputArrow #if TEST_STD_VER < 2020 TEST_FUNC friend bool operator!=(LegacyInputArrow, LegacyInputArrow); #endif - // Otherwise, if decltype(​declval().operator->()) is well-formed, then pointer names that type. + // Otherwise, if decltype(declval().operator->()) is well-formed, then pointer names that type. TEST_FUNC int* operator->(); TEST_FUNC reference operator*() const; TEST_FUNC LegacyInputArrow& operator++(); @@ -223,7 +223,7 @@ struct LegacyInputPointer { TEST_FUNC operator value_type() const; }; - // If the qualified-id I​::​pointer is valid and denotes a type, then pointer names that type. + // If the qualified-id I::pointer is valid and denotes a type, then pointer names that type. struct pointer {}; @@ -231,7 +231,7 @@ struct LegacyInputPointer #if TEST_STD_VER < 2020 TEST_FUNC friend bool operator!=(LegacyInputPointer, LegacyInputPointer); #endif - // Otherwise, if decltype(​declval().operator->()) is well-formed, then pointer names that type. + // Otherwise, if decltype(declval().operator->()) is well-formed, then pointer names that type. TEST_FUNC int* operator->(); TEST_FUNC reference operator*() const; TEST_FUNC LegacyInputPointer& operator++(); 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 2f73fc29e76..26deca904bb 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 @@ -124,7 +124,7 @@ static_assert(cuda::std::is_nothrow_invocable_v, int>); @@ -228,7 +228,7 @@ TEST_FUNC constexpr bool test() { // gcc < 13 fails this test with error: - // ‘fun_ptr’ is not a valid template argument of type ‘bool (*)(int)’ because ‘fun_ptr’ is not a variable + // 'fun_ptr' is not a valid template argument of type 'bool (*)(int)' because 'fun_ptr' is not a variable #if !_CCCL_COMPILER(GCC, <, 13) // function pointer using T = cuda::std::__constant_wrapper; @@ -239,7 +239,7 @@ TEST_FUNC constexpr bool test() { // gcc < 13 fails this test with error: - // ‘fun_ptr’ is not a valid template argument of type ‘bool (*)(int)’ because ‘fun_ptr’ is not a variable + // 'fun_ptr' is not a valid template argument of type 'bool (*)(int)' because 'fun_ptr' is not a variable #if !_CCCL_COMPILER(GCC, <, 13) // function pointer with constexpr param using T = cuda::std::__constant_wrapper; diff --git a/libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/convert.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/convert.pass.cpp index ef009eaa009..4e613cf3de0 100644 --- a/libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/convert.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/convert.pass.cpp @@ -77,7 +77,7 @@ TEST_FUNC constexpr bool test() { // gcc < 13 fails this test with: - // ‘test()::::_FUN’ is not a valid template argument of type ‘int (*)(int)’ because it is not + // 'test()::::_FUN' is not a valid template argument of type 'int (*)(int)' because it is not // a variable #if !_CCCL_COMPILER(GCC, <, 13) // function pointer conversion diff --git a/libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/cw.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/cw.pass.cpp index a9aa4e00842..4c05955b20e 100644 --- a/libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/cw.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/cw.pass.cpp @@ -52,7 +52,7 @@ TEST_FUNC constexpr bool test() { // gcc < 13 fails this test with error: - // invalid use of non-static data member ‘S::value’ + // invalid use of non-static data member 'S::value' #if !_CCCL_COMPILER(GCC, <, 13) // struct constant constexpr S s{13}; diff --git a/libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/unary_ops.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/unary_ops.pass.cpp index 837601a7031..96e9c011ef3 100644 --- a/libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/unary_ops.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/unary_ops.pass.cpp @@ -244,7 +244,7 @@ TEST_FUNC constexpr bool test() static_assert(result4 == !42); // gcc < 13 fails this test with error: - // the address of ‘cuda::std::__4::__cw_fixed_value{42}’ is not a valid template argument + // the address of 'cuda::std::__4::__cw_fixed_value{42}' is not a valid template argument #if !_CCCL_COMPILER(GCC, <, 13) cuda::std::same_as> decltype(auto) result5 = &cw42; static_assert(result5 == &cw42.value); diff --git a/libcudacxx/test/libcudacxx/std/utilities/time/time.clock/time.clock.system/local_time.types.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/time/time.clock/time.clock.system/local_time.types.pass.cpp index f5c50e290b0..c23957d0579 100644 --- a/libcudacxx/test/libcudacxx/std/utilities/time/time.clock/time.clock.system/local_time.types.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/utilities/time/time.clock/time.clock.system/local_time.types.pass.cpp @@ -17,7 +17,7 @@ // [Example: // sys_seconds{sys_days{1970y/January/1}}.time_since_epoch() is 0s. -// sys_seconds{sys_days{2000y/January/1}}.time_since_epoch() is 946’684’800s, which is 10’957 * 86’400s. +// sys_seconds{sys_days{2000y/January/1}}.time_since_epoch() is 946'684'800s, which is 10'957 * 86'400s. // —end example] #include diff --git a/libcudacxx/test/libcudacxx/std/utilities/time/time.clock/time.clock.system/sys.time.types.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/time/time.clock/time.clock.system/sys.time.types.pass.cpp index 30be523df90..636a8d37be7 100644 --- a/libcudacxx/test/libcudacxx/std/utilities/time/time.clock/time.clock.system/sys.time.types.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/utilities/time/time.clock/time.clock.system/sys.time.types.pass.cpp @@ -16,7 +16,7 @@ // [Example: // sys_seconds{sys_days{1970y/January/1}}.time_since_epoch() is 0s. -// sys_seconds{sys_days{2000y/January/1}}.time_since_epoch() is 946’684’800s, which is 10’957 * 86’400s. +// sys_seconds{sys_days{2000y/January/1}}.time_since_epoch() is 946'684'800s, which is 10'957 * 86'400s. // —end example] #include