Repository navigation
Fix uninitialized read in uninitialized_copy_n - #5811
Merged
Merged
Conversation
We need to tighten the check to use copy_n only when the copy ctor also has no side effect. Fixes: NVIDIA#5809
Contributor
🥳 CI Workflow Results🟩 Finished in 3h 26m: Pass: 100%/154 | Total: 5d 17h | Max: 3h 19m | Hits: 77%/177577See results here. |
miscco
approved these changes
Sep 9, 2025
Naraenda
pushed a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 26, 2026
…initialized_copy_n (NVIDIA/cccl#5811) Upstream: NVIDIA/cccl@ee141c6 - Clean port (system/cuda): uninitialized_copy_n only delegates to copy_n when construction AND assignment are trivial. - HIP counterpart: no change: the copy_n fast path in uninitialized_copy_n was ported to system/cuda only (performance-only, item 4b2201184c4); hip uninitialized_copy still always placement-constructs, so the bug does not exist there - test/ counterpart: none (no upstream test change) - Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>
amd-jmahovsky
added a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 28, 2026
…ialized_copy delegates to copy_n Related upstream commit: 4b2201184c4 Use `thrust::copy` in `thrust::uninitialized_copy[_n]` in CUDA system when possible (NVIDIA/cccl#5181) Related upstream commit: ee141c6b175 Fix uninitialized read in uninitialized_copy_n (NVIDIA/cccl#5811) Related local port commit: 9d7099b Related local port commit: c91e8ff Both upstream commits were ported to system/cuda only, with the HIP counterpart marked "no change: performance-only". Bring system/hip in line: - hip_rocprim::uninitialized_copy_n now delegates to hip_rocprim::copy_n when the output value type is trivially constructible from, and trivially assignable from, the input reference (the #5811 condition). For trivially relocatable device-to-device copies this turns the placement-new parallel_for kernel into a hipMemcpyAsync. - copy_n is forward-declared rather than included, as upstream does, to avoid the circular include via internal/copy_cross_system.h. - Includes follow upstream's granular libhipcxx headers (__iterator/distance.h, __iterator/incrementable_traits.h, and the is_trivially_{constructible,assignable} traits actually used). Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>
Naraenda
pushed a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 29, 2026
…initialized_copy_n (NVIDIA/cccl#5811) Upstream: NVIDIA/cccl@ee141c6 - Clean port (system/cuda): uninitialized_copy_n only delegates to copy_n when construction AND assignment are trivial. - HIP counterpart: no change: the copy_n fast path in uninitialized_copy_n was ported to system/cuda only (performance-only, item 4b2201184c4); hip uninitialized_copy still always placement-constructs, so the bug does not exist there - test/ counterpart: none (no upstream test change) - Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>
Naraenda
pushed a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 29, 2026
…initialized_copy_n (NVIDIA/cccl#5811) Upstream: NVIDIA/cccl@ee141c6 - Clean port (system/cuda): uninitialized_copy_n only delegates to copy_n when construction AND assignment are trivial. - HIP counterpart: no change: the copy_n fast path in uninitialized_copy_n was ported to system/cuda only (performance-only, item 4b2201184c4); hip uninitialized_copy still always placement-constructs, so the bug does not exist there - test/ counterpart: none (no upstream test change) - Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>
Naraenda
pushed a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 29, 2026
…ialized_copy delegates to copy_n Related upstream commit: 4b2201184c4 Use `thrust::copy` in `thrust::uninitialized_copy[_n]` in CUDA system when possible (NVIDIA/cccl#5181) Related upstream commit: ee141c6b175 Fix uninitialized read in uninitialized_copy_n (NVIDIA/cccl#5811) Related local port commit: 9d7099b Related local port commit: c91e8ff Both upstream commits were ported to system/cuda only, with the HIP counterpart marked "no change: performance-only". Bring system/hip in line: - hip_rocprim::uninitialized_copy_n now delegates to hip_rocprim::copy_n when the output value type is trivially constructible from, and trivially assignable from, the input reference (the #5811 condition). For trivially relocatable device-to-device copies this turns the placement-new parallel_for kernel into a hipMemcpyAsync. - copy_n is forward-declared rather than included, as upstream does, to avoid the circular include via internal/copy_cross_system.h. - Includes follow upstream's granular libhipcxx headers (__iterator/distance.h, __iterator/incrementable_traits.h, and the is_trivially_{constructible,assignable} traits actually used). Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>
Naraenda
pushed a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 30, 2026
…initialized_copy_n (NVIDIA/cccl#5811) Upstream: NVIDIA/cccl@ee141c6 - Clean port (system/cuda): uninitialized_copy_n only delegates to copy_n when construction AND assignment are trivial. - HIP counterpart: no change: the copy_n fast path in uninitialized_copy_n was ported to system/cuda only (performance-only, item 4b2201184c4); hip uninitialized_copy still always placement-constructs, so the bug does not exist there - test/ counterpart: none (no upstream test change) - Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>
Beanavil
pushed a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 30, 2026
…initialized_copy_n (NVIDIA/cccl#5811) Upstream: NVIDIA/cccl@ee141c6 - Clean port (system/cuda): uninitialized_copy_n only delegates to copy_n when construction AND assignment are trivial. - HIP counterpart: no change: the copy_n fast path in uninitialized_copy_n was ported to system/cuda only (performance-only, item 4b2201184c4); hip uninitialized_copy still always placement-constructs, so the bug does not exist there - test/ counterpart: none (no upstream test change) - Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>
amd-jmahovsky
added a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 30, 2026
…ialized_copy delegates to copy_n Related upstream commit: 4b2201184c4 Use `thrust::copy` in `thrust::uninitialized_copy[_n]` in CUDA system when possible (NVIDIA/cccl#5181) Related upstream commit: ee141c6b175 Fix uninitialized read in uninitialized_copy_n (NVIDIA/cccl#5811) Related local port commit: 9d7099b Related local port commit: c91e8ff Both upstream commits were ported to system/cuda only, with the HIP counterpart marked "no change: performance-only". Bring system/hip in line: - hip_rocprim::uninitialized_copy_n now delegates to hip_rocprim::copy_n when the output value type is trivially constructible from, and trivially assignable from, the input reference (the #5811 condition). For trivially relocatable device-to-device copies this turns the placement-new parallel_for kernel into a hipMemcpyAsync. - copy_n is forward-declared rather than included, as upstream does, to avoid the circular include via internal/copy_cross_system.h. - Includes follow upstream's granular libhipcxx headers (__iterator/distance.h, __iterator/incrementable_traits.h, and the is_trivially_{constructible,assignable} traits actually used). Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
We need to tighten the check to use copy_n only when the copy ctor also has no side effect.
Fixes: #5809