Repository navigation
Add host and device only type tests for unique_ptr - #11371
theo-ai-lab wants to merge 3 commits into
Conversation
PR NVIDIA#4732 suppressed the execution space checks on six unique_ptr functions so that they can be used with types whose members are callable from only one execution space. Nothing exercised that, so the suppression could regress silently. These two tests cover exactly those six functions: default_delete<T>::operator() default_delete<T[]>::operator() make_unique<T>() make_unique<T[]>() make_unique_for_overwrite<T>() make_unique_for_overwrite<T[]>() The deleters are reached through scope exit, both for pointers the unique_ptr allocated itself and for pointers handed to it, so each deleter overload runs against a single object and an array. They follow the cuda/utilities/<type>/{host,device}_only_types.pass.cpp family introduced in NVIDIA#3578, which is where every test using host_device_types.h lives. Unlike those siblings there is no tile variant and no tile branch in the device test: unique_ptr allocates, and dynamic allocation is not available in tile mode, which is why the existing make_unique tests are all UNSUPPORTED on force-tile. Closes NVIDIA#4748
|
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Path: .coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (2)
🚧 Files skipped from review as they are similar to previous changes (1)
Included review availability: Your plan provides up to 12 included reviews per hour; 11 remain after this review. 📝 SummarySummary by CodeRabbit
WalkthroughAdded host-only and device-only tests for Changesunique_ptr host-device test coverage
Assessment Against Linked Issues
Priority: ⬇️ Low Change: Other Merge Risk: ⚪ Minimal · up to This test-only change adds focused host and device coverage, with no remaining concrete merge-blocking risk. Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 1
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 531672f7-a727-4832-8a5c-ad12ace8c288
📒 Files selected for processing (2)
libcudacxx/test/libcudacxx/cuda/utilities/unique_ptr/device_only_types.pass.cpplibcudacxx/test/libcudacxx/cuda/utilities/unique_ptr/host_only_types.pass.cpp
Included review availability: Your plan provides up to 12 included reviews per hour; 11 remain after this review.
miscco
left a comment
miscco
left a comment
There was a problem hiding this comment.
Looks good to me, thanks a lot
|
/ok to test a1575fd |
😬 CI Workflow Results🟥 Finished in 56m 41s: Pass: 55%/140 | Total: 1d 00h | Max: 56m 41s | Hits: 99%/216445See results here. AI failure analysis1. libcudacxx unique_ptr host-only test violates CUDA execution space · 59 jobsExplanation: All 59 libcudacxx jobs have this same earliest failure in the newly added host-only unique_ptr test. A host-only constructor is instantiated through an unidentified host/device context; the logs lack an instantiation backtrace, so the precise unique_ptr operation or test expression responsible must be isolated narrowly. Evidence: Copy this prompt into a coding agentJobs:
2. RAPIDS cuVS and cuDF reference missing rmm::cuda_stream_view · 3 jobsExplanation: Both cuVS jobs and the cuDF job build current RAPIDS main branches against the same current RMM dependency, whose visible namespace no longer supplies the referenced stream type. This is downstream dependency-version skew rather than a failure caused by the two-file PR diff; the logs do not identify the exact RMM removal commit. Evidence: Copy this prompt into a coding agentJobs: 3. cuOpt calls missing cuda::stream_ref::synchronize · 1 jobExplanation: cuOpt main calls `synchronize()` repeatedly on `cuda::stream_ref`, while the checked-out CCCL API provides `sync()` instead. The first barrier compilation fails at all seven such calls before the integration wrapper reports the cuOpt failure. Evidence: Copy this prompt into a coding agentJobs: |
nvcc lowers new T[n] inside make_unique<T[]> and make_unique_for_overwrite<T[]>
to an unnamed __host__ __device__ element construction loop, and
_CCCL_EXEC_CHECK_DISABLE on the enclosing function does not extend to that
generated helper:
host_only_types.pass.cpp(101): error #20011-D: calling a __host__ function
("host_only_type::host_only_type(int)") from a __host__ __device__
function("<unknown>") is not allowed
Exactly two such sites, exactly two errors. The three single object factories
use a direct new-expression and are fine.
The host test now reaches default_delete<T[]>::operator() through arrays the
unique_ptr takes ownership of, including via reset, so both deleter overloads
stay covered. device_only_types.pass.cpp is unchanged and already passes: its
helpers are __device__, so no cross execution space call arises.
|
Thanks for the review and for running CI. It found a genuine gap, and I think it is worth your opinion.
Both come from the two array factories. The device test is unaffected because its helpers are So The question: this looks like a real hole in the #4732 coverage rather than a test problem — Also, for the record on the analyzer's other suggestion: Could I get a fresh |
This branch has not been deployed
2 participants
Description
closes #4748
Adds
host_only_types.pass.cppanddevice_only_types.pass.cppforunique_ptr, usingtest/support/host_device_types.hand following thecuda/utilities/{optional,tuple,expected}/pattern introduced in #3578.What is covered
The six functions #4732 marked
_CCCL_EXEC_CHECK_DISABLE:default_delete<T>::operator()(T*)unique_ptr<T>default_delete<T[]>::operator()<U>(U*)(convertible-array member template)unique_ptr<T[]>make_unique<T>(Args&&...)→__unique_singlemake_unique<T>(size_t)→__unique_array_unknown_boundmake_unique_for_overwrite<T>()→__unique_singlemake_unique_for_overwrite<T>(size_t)→__unique_array_unknown_boundEach deleter overload is exercised both for a pointer the
unique_ptrallocated itself andfor one handed to it.
Gap:
mainnow has nine markers, not sixunique_ptr.honmaincarries nine_CCCL_EXEC_CHECK_DISABLEmarkers. The three addedafter #4732 are not covered here:
unique_ptr<T, D>::reset(pointer)unique_ptr<T[], D>::reset(_Pp)unique_ptr<T[], D>::reset(nullptr_t)I scoped this PR to the #4732 set that the issue refers to. Happy to extend it to the three
resetoverloads in this PR if you would prefer them covered together.Placement — would like your call
I put the tests under
cuda/utilities/unique_ptr/, following precedent: all 31*_only_types.pass.cppfiles and all includers ofhost_device_types.hlive undercuda/,and #3578 created new
cuda/utilities/<type>/directories for exactly this kind of test eventhough those types already had large
std/homes.Two things point the other way, so I may have chosen wrong:
std/utilities/memory/smartptr/unique.ptr/unique.ptr.dltr/..agent/skills/cccl-test/references/libcudacxx.mdsays to put CUDA-specific API testsunder
cuda/..."unless an adjacentstd/...directory is clearly the established homefor the functionality" — and with 54 files,
std/.../unique.ptr/arguably is.Either location is a
git mvaway; tell me which you want and I will move it. The directoryspelling is open too (
cuda/utilities/unique_ptr/vs a deepercuda/utilities/memory/smartptr/unique.ptr/mirror).Three deliberate deviations from the sibling tests
// UNSUPPORTED: force-tile, notenable-tile.unique_ptrallocates, and dynamicallocation is unavailable in tile mode. This matches the existing
make_unique/make_unique_for_overwritetests, which all useforce-tileonly.tile_only_types.pass.cpp. The seven sibling tile tests are gated// REQUIRES: enable-tile, so a tile variant would run only where allocation cannot work.#include <cuda/std/__memory_>rather than a non-internal header.<cuda/std/memory>does not declare
make_uniqueormake_unique_for_overwrite;__memory_is what all 54existing
unique.ptrtests include.I also left off optional's
// UNSUPPORTED: clang-14— that exclusion is aboutcuda::std::construct_at, andunique_ptr.hnever routes through it (all four factories useplain
new), which matchestuple/pair/variant/unexpectednot carrying it either.Verification
Verified locally:
host_only_types.pass.cppcompiles and runs green at-std=c++17,c++20andc++23.new[]/delete[]in the array-deleter paths).pre-commit run --files <both files>exits 0,clang-formatincluded.device_only_types.pass.cppis not verified. There is nonvcc, no CUDA toolkit and noNVIDIA GPU on the machine I wrote it on, so
_CCCL_DEVICE_COMPILATION()is false and theentire test body preprocesses away — it has never been compiled as device code, let alone
run. It needs CI to be meaningfully exercised. Two things I therefore could not check: whether
NVRTC needs an additional
// UNSUPPORTED: nvrtc, and whether device-sidenewbehaves asexpected for these types.
Checklist