libcu++: atomics SASS FileCheck suite for arithmetic atomic operations - #10832
libcu++: atomics SASS FileCheck suite for arithmetic atomic operations#10832griwes wants to merge 18 commits into
Conversation
… volatile codegen tests.
📝 WalkthroughSummary by CodeRabbit
WalkthroughThis change adds configurable atomic and SIMD codegen coverage, expands atomic APIs for cv-qualified and extended floating-point types, and integrates target-specific execution into CMake, build scripts, matrix workflows, and FileCheck validation. ChangesAtomic codegen coverage and CI
Assessment against linked issues
Possibly related PRs
Suggested reviewers: Merge Risk: 🟡 Moderate · up to This PR adds broad SASS coverage and changes codegen-test infrastructure, but the current head can run unsupported 128-bit tests, reject valid test configurations because of duplicate prefixes, or pass while validating the wrong or incomplete instruction patterns. It is not merge-ready until these bounded test correctness issues are fixed or explicitly accepted. 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: 5
🧹 Nitpick comments (4)
libcudacxx/test/cmake/CodegenTest.cmake (2)
253-279: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: a prefix that is only a
PREFIX_COMBINEcomponent is dropped when no combination is active, sofilecheck_prefixescan shrink to whatever remains. If a test ever lists only component prefixes, the resolved list becomes empty and FileCheck fails with an unclear error. Add an explicit guard.set(filecheck_prefixes ${standalone_prefixes} ${active_combinations}) - list(REMOVE_DUPLICATES standalone_prefixes) list(REMOVE_DUPLICATES filecheck_prefixes) + if (NOT filecheck_prefixes) + message( + FATAL_ERROR + "No usable FileCheck prefixes remain for ${test_path} from '${input_prefixes}'" + ) + endif() + list(REMOVE_DUPLICATES standalone_prefixes) list(JOIN standalone_prefixes "," standalone_prefixes)Note that
list(REMOVE_DUPLICATES standalone_prefixes)at Line 273 runs after Line 272 consumed the variable, so it has no effect onfilecheck_prefixes.
562-568: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valuesuggestion: the second loop reuses
has_arch_specific_checksas an output variable and overwrites the accumulated value from Line 529. The accumulated value is no longer needed here, but the shared name invites a future regression. Use a separate output name.libcudacxx_codegen_get_sass_check_prefixes( check_prefixes - has_arch_specific_checks + arch_has_specific_checks "${test_contents}" "${arch}" )libcudacxx/test/atomic_codegen/sass/compare_exchange_types_128_atomic.cu (1)
28-49: 🎯 Functional Correctness | 🔵 Trivial | ⚡ Quick winsuggestion: the checks confirm the lock-based path but never forbid the native 128-bit path. If codegen later selects
ATOM.E.CAS.128for the owning atomic, this test still passes. Add a negative directive to pin the intended path.; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test +; SMXX-NOT: {{.*}}ATOM.E.CAS.128{{.*}} ; SMXX-NOT: {{.*}}MEMBAR.SC.{{.*}}libcudacxx/test/atomic_codegen/sass/arithmetic_floating_types.cu (1)
14-14: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion:
fetch_addandfetch_submap to the same expected SASS opcode in these files, and no check covers the operand negation. Thesubvariant therefore passes even if the negation is missing.arithmetic_apis.cusolves this withFILECHECK_PREFIX_OPplus anADD-NOT/SUB-DAGpair.
libcudacxx/test/atomic_codegen/sass/arithmetic_floating_types.cu#L14-L14: addFILECHECK_PREFIX_OPand a sub-only check on the negation instruction.libcudacxx/test/atomic_codegen/sass/arithmetic_volatile_types_8_16_atomic.cu#L14-L14: add the same op-specific prefix and negation check for theadd/subpair.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 4656a25e-8d61-42cf-a9f9-3d9be67ebb62
📒 Files selected for processing (100)
.github/actions/workflow-build/build-workflow.pyCMakePresets.jsonci/build_libcudacxx.shci/matrix.yamllibcudacxx/include/cuda/__atomic/atomic.hlibcudacxx/include/cuda/std/__atomic/api/common.hlibcudacxx/include/cuda/std/__atomic/api/owned.hlibcudacxx/include/cuda/std/__atomic/api/reference.hlibcudacxx/include/cuda/std/__atomic/functions/cuda_ptx_derived.hlibcudacxx/include/cuda/std/__atomic/types/base.hlibcudacxx/include/cuda/std/__atomic/types/reference.hlibcudacxx/include/cuda/std/__atomic/types/small.hlibcudacxx/include/cuda/std/atomiclibcudacxx/test/CMakeLists.txtlibcudacxx/test/atomic_codegen/CMakeLists.txtlibcudacxx/test/atomic_codegen/atomic_add_non_volatile.culibcudacxx/test/atomic_codegen/atomic_cas_non_volatile.culibcudacxx/test/atomic_codegen/atomic_exch_non_volatile.culibcudacxx/test/atomic_codegen/atomic_load_non_volatile.culibcudacxx/test/atomic_codegen/atomic_store_non_volatile.culibcudacxx/test/atomic_codegen/atomic_sub_non_volatile.culibcudacxx/test/atomic_codegen/sass/CMakeLists.txtlibcudacxx/test/atomic_codegen/sass/arithmetic_apis.culibcudacxx/test/atomic_codegen/sass/arithmetic_floating_types.culibcudacxx/test/atomic_codegen/sass/arithmetic_minmax_apis.culibcudacxx/test/atomic_codegen/sass/arithmetic_minmax_floating_types.culibcudacxx/test/atomic_codegen/sass/arithmetic_minmax_types.culibcudacxx/test/atomic_codegen/sass/arithmetic_minmax_types_8_16_atomic.culibcudacxx/test/atomic_codegen/sass/arithmetic_minmax_types_8_16_atomic_ref.culibcudacxx/test/atomic_codegen/sass/arithmetic_minmax_volatile_apis.culibcudacxx/test/atomic_codegen/sass/arithmetic_minmax_volatile_floating_types.culibcudacxx/test/atomic_codegen/sass/arithmetic_minmax_volatile_types.culibcudacxx/test/atomic_codegen/sass/arithmetic_minmax_volatile_types_8_16_atomic.culibcudacxx/test/atomic_codegen/sass/arithmetic_minmax_volatile_types_8_16_atomic_ref.culibcudacxx/test/atomic_codegen/sass/arithmetic_operators_apis.culibcudacxx/test/atomic_codegen/sass/arithmetic_operators_volatile_apis.culibcudacxx/test/atomic_codegen/sass/arithmetic_pointer_types.culibcudacxx/test/atomic_codegen/sass/arithmetic_types.culibcudacxx/test/atomic_codegen/sass/arithmetic_types_128_atomic_ref.culibcudacxx/test/atomic_codegen/sass/arithmetic_types_8_16_atomic.culibcudacxx/test/atomic_codegen/sass/arithmetic_types_8_16_atomic_ref.culibcudacxx/test/atomic_codegen/sass/arithmetic_volatile_apis.culibcudacxx/test/atomic_codegen/sass/arithmetic_volatile_floating_types.culibcudacxx/test/atomic_codegen/sass/arithmetic_volatile_pointer_types.culibcudacxx/test/atomic_codegen/sass/arithmetic_volatile_types.culibcudacxx/test/atomic_codegen/sass/arithmetic_volatile_types_128_atomic_ref.culibcudacxx/test/atomic_codegen/sass/arithmetic_volatile_types_8_16_atomic.culibcudacxx/test/atomic_codegen/sass/arithmetic_volatile_types_8_16_atomic_ref.culibcudacxx/test/atomic_codegen/sass/atomic_codegen_helpers.hlibcudacxx/test/atomic_codegen/sass/compare_exchange_apis.culibcudacxx/test/atomic_codegen/sass/compare_exchange_single_order.culibcudacxx/test/atomic_codegen/sass/compare_exchange_types.culibcudacxx/test/atomic_codegen/sass/compare_exchange_types_128_atomic.culibcudacxx/test/atomic_codegen/sass/compare_exchange_types_128_atomic_ref.culibcudacxx/test/atomic_codegen/sass/compare_exchange_types_8_16_atomic.culibcudacxx/test/atomic_codegen/sass/compare_exchange_types_8_16_atomic_ref.culibcudacxx/test/atomic_codegen/sass/compare_exchange_volatile_apis.culibcudacxx/test/atomic_codegen/sass/compare_exchange_volatile_types.culibcudacxx/test/atomic_codegen/sass/compare_exchange_volatile_types_128_atomic_ref.culibcudacxx/test/atomic_codegen/sass/compare_exchange_volatile_types_8_16_atomic.culibcudacxx/test/atomic_codegen/sass/compare_exchange_volatile_types_8_16_atomic_ref.culibcudacxx/test/atomic_codegen/sass/exchange_apis.culibcudacxx/test/atomic_codegen/sass/exchange_types.culibcudacxx/test/atomic_codegen/sass/exchange_types_128_atomic.culibcudacxx/test/atomic_codegen/sass/exchange_types_128_atomic_ref.culibcudacxx/test/atomic_codegen/sass/exchange_types_8_16_atomic.culibcudacxx/test/atomic_codegen/sass/exchange_types_8_16_atomic_ref.culibcudacxx/test/atomic_codegen/sass/exchange_volatile_apis.culibcudacxx/test/atomic_codegen/sass/exchange_volatile_types.culibcudacxx/test/atomic_codegen/sass/exchange_volatile_types_128_atomic_ref.culibcudacxx/test/atomic_codegen/sass/exchange_volatile_types_8_16_atomic.culibcudacxx/test/atomic_codegen/sass/exchange_volatile_types_8_16_atomic_ref.culibcudacxx/test/atomic_codegen/sass/load_apis.culibcudacxx/test/atomic_codegen/sass/load_types.culibcudacxx/test/atomic_codegen/sass/load_types_128_atomic.culibcudacxx/test/atomic_codegen/sass/load_types_128_atomic_ref.culibcudacxx/test/atomic_codegen/sass/load_types_8_16_atomic.culibcudacxx/test/atomic_codegen/sass/load_types_8_16_atomic_ref.culibcudacxx/test/atomic_codegen/sass/load_volatile_apis.culibcudacxx/test/atomic_codegen/sass/load_volatile_types.culibcudacxx/test/atomic_codegen/sass/load_volatile_types_128_atomic_ref.culibcudacxx/test/atomic_codegen/sass/load_volatile_types_8_16_atomic.culibcudacxx/test/atomic_codegen/sass/load_volatile_types_8_16_atomic_ref.culibcudacxx/test/atomic_codegen/sass/store_apis.culibcudacxx/test/atomic_codegen/sass/store_types.culibcudacxx/test/atomic_codegen/sass/store_types_128_atomic.culibcudacxx/test/atomic_codegen/sass/store_types_128_atomic_ref.culibcudacxx/test/atomic_codegen/sass/store_types_8_16_atomic.culibcudacxx/test/atomic_codegen/sass/store_types_8_16_atomic_ref.culibcudacxx/test/atomic_codegen/sass/store_volatile_apis.culibcudacxx/test/atomic_codegen/sass/store_volatile_types.culibcudacxx/test/atomic_codegen/sass/store_volatile_types_128_atomic_ref.culibcudacxx/test/atomic_codegen/sass/store_volatile_types_8_16_atomic.culibcudacxx/test/atomic_codegen/sass/store_volatile_types_8_16_atomic_ref.culibcudacxx/test/cmake/CodegenTest.cmakelibcudacxx/test/codegen/dump_and_check.bashlibcudacxx/test/libcudacxx/cuda/atomics/atomic.ext/atomic_fetch_extended_floating_point.pass.cpplibcudacxx/test/libcudacxx/std/atomics/atomics.types.generic/atomic_ref_volatile.pass.cpplibcudacxx/test/simd_codegen/CMakeLists.txtlibcudacxx/test/simd_codegen/load_store/CMakeLists.txt
| ; SMXX-DAG: ld.param.{{b|u}}64 %rd[[#EXPECTED:]], {{.*}}[[FUNCTION]]_param_1{{.*}} | ||
| ; SMXX-DAG: ld.param.{{b|u}}32 %r[[#INPUT:]], {{.*}}[[FUNCTION]]_param_2{{.*}} | ||
| ; SMXX-DAG: cvta.to.global.u64 %rd[[#GOUT:]], %rd[[#EXPECTED]]; | ||
| ; SMXX-DAG: ld.global.{{b|u}}32 %r[[#LOCALEXP:]], [%rd[[#INPUT]]]; |
There was a problem hiding this comment.
🎯 Functional Correctness | 🟠 Major | ⚡ Quick win
important: Match the expected operand address. Line 18 derives the load address from INPUT, which captures the scalar n parameter. The kernel reads the expected value from out, converted in Line 17 into GOUT. Match [%rd[[#GOUT]]] so FileCheck validates the compare-exchange operand data flow.
| // %PARAM% TEMPLATE,SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE api cab=ca,tsb,CTA,block:card=car,tsd,GPU,non_block | ||
| // %PARAM% TYPE,SASS_SIZE,SASS_CALC type f32=float,,FSETP:f64=double,.64,DSETP | ||
| // %PARAM% OP op fetch_min:fetch_max | ||
| // %PARAM% ORDER,FILECHECK_PREFIX_SEQ_CST,FILECHECK_PREFIX_ACQUIRE,FILECHECK_PREFIX_ORDER order relaxed=mor,non_seq_cst,no_acquire,no_membar:acquire=moa,non_seq_cst,acquire,no_membar:release=more,non_seq_cst,no_acquire,release:acq_rel=moar,non_seq_cst,acquire,release:seq_cst=mosc,seq_cst,acquire,seq_cst |
There was a problem hiding this comment.
🎯 Functional Correctness | 🟠 Major | ⚡ Quick win
important: three ORDER parameter lists set FILECHECK_PREFIX_SEQ_CST and FILECHECK_PREFIX_ORDER to the same value seq_cst for the seq_cst variant, so the SEQ_CST prefix is contributed twice. FileCheck rejects a repeated entry in --check-prefixes unless the harness deduplicates.
libcudacxx/test/atomic_codegen/sass/arithmetic_minmax_floating_types.cu#L15-L15: rename theFILECHECK_PREFIX_ORDERvalue forseq_cstto a distinct prefix, or confirm deduplication in the harness.libcudacxx/test/atomic_codegen/sass/arithmetic_minmax_types_8_16_atomic.cu#L15-L15: apply the same rename.libcudacxx/test/atomic_codegen/sass/arithmetic_volatile_types_8_16_atomic_ref.cu#L15-L15: apply the same rename.
📍 Affects 3 files
libcudacxx/test/atomic_codegen/sass/arithmetic_minmax_floating_types.cu#L15-L15(this comment)libcudacxx/test/atomic_codegen/sass/arithmetic_minmax_types_8_16_atomic.cu#L15-L15libcudacxx/test/atomic_codegen/sass/arithmetic_volatile_types_8_16_atomic_ref.cu#L15-L15
| //===----------------------------------------------------------------------===// | ||
|
|
||
| // clang-format off | ||
| // %PARAM% TEMPLATE,SCOPE,SASS_SCOPE,FILECHECK_PREFIX_SCOPE api vcad=vca,tsd,GPU,non_block:vcard=vcar,tsd,GPU,non_block |
There was a problem hiding this comment.
📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick win
important: These volatile SASS tests declare no block-scope variants, so their BLOCK checks are dead and never execute. Either add the corresponding block-scope variants or remove the unused BLOCK directives.
📍 Affects 2 files
libcudacxx/test/atomic_codegen/sass/arithmetic_volatile_floating_types.cu#L12-L12(this comment)libcudacxx/test/atomic_codegen/sass/arithmetic_volatile_types.cu#L12-L12
| ; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test | ||
| ; SMXX-NOT: {{.*}}ATOM.{{.*}} | ||
| ; SMXX: {{.*}}LD.E.64{{(\.SYS)?}} [[ATOM_ADDR:R[0-9]+]], {{.*}} | ||
| ; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} | ||
| ; NON_SC-NOT: {{.*}}CCTL.IVALL{{.*}} | ||
| ; SEQ_CST: {{.*}}MEMBAR.SC.[[SASS_SCOPE]]{{.*}} | ||
| ; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} | ||
| ; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} | ||
| ; NON_SC-NOT: {{.*}}CCTL.IVALL{{.*}} | ||
| ; NON_SC-NOT: {{.*}}MEMBAR.SC.{{.*}} | ||
| ; SMXX-NOT: {{.*}}ATOM.{{.*}} | ||
| ; BLOCK: {{.*}}LD.E.128.STRONG.{{CTA|SM}} R4, {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}} | ||
| ; NON_BLOCK: {{.*}}LD.E.128.STRONG.[[SASS_SCOPE]] R4, {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}} | ||
| ; NON_BLOCK_ACQUIRE-NEXT: {{.*}}CCTL.IVALL{{.*}} | ||
| ; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} | ||
| ; NO_ACQUIRE-NOT: {{.*}}CCTL.IVALL{{.*}} | ||
| ; SMXX-NOT: {{.*}}ATOM.{{.*}} | ||
| ; SMXX: {{.*}}RET.ABS.NODEC{{.*}} |
There was a problem hiding this comment.
🎯 Functional Correctness | 🟠 Major | ⚡ Quick win
important: Gate both 128-bit atomic_ref load tests at SM70 or newer. These operations are unsupported before SM70. The current SMXX checks can run the tests on unsupported targets and fail the codegen suite.
libcudacxx/test/atomic_codegen/sass/load_types_128_atomic_ref.cu#L28-L45: replace the unrestricted architecture checks with an SM70+ gate.libcudacxx/test/atomic_codegen/sass/load_volatile_types_128_atomic_ref.cu#L28-L48: apply the same SM70+ gate.
Based on learnings: “gate load and store operations at SM70+ because they are unsupported before SM70.”
📍 Affects 2 files
libcudacxx/test/atomic_codegen/sass/load_types_128_atomic_ref.cu#L28-L45(this comment)libcudacxx/test/atomic_codegen/sass/load_volatile_types_128_atomic_ref.cu#L28-L48
Source: Learnings
| extern "C" __device__ void atomic_codegen_test(cuda::atomic_ref<volatile TYPE, SCOPE>& atom, TYPE value) | ||
| { | ||
| atom.store(value, ORDER); | ||
| } | ||
|
|
||
| /* | ||
|
|
||
| ; SMXX-LABEL: {{[[:space:]]*}}Function : atomic_codegen_test | ||
| ; SMXX-NOT: {{.*}}ATOM.{{.*}} | ||
| ; SMXX: {{.*}}LD.E.64{{(\.SYS)?}} [[ATOM_ADDR:R[0-9]+]], {{.*}} | ||
| ; SMXX-NOT: {{.*}}ST.E{{.*\[}}[[ATOM_ADDR]]{{(\.64)?(\+0x[0-9a-f]+)?\].*}} | ||
| ; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} | ||
| ; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} | ||
| ; MEMBAR: {{.*}}MEMBAR.[[SASS_MEMBAR]].[[SASS_SCOPE]]{{.*}} | ||
| ; NON_BLOCK_SEQ_CST: {{.*}}CCTL.IVALL{{.*}} | ||
| ; BLOCK-NOT: {{.*}}CCTL.IVALL{{.*}} | ||
| ; NON_SEQ_CST-NOT: {{.*}}CCTL.IVALL{{.*}} | ||
| ; NO_MEMBAR-NOT: {{.*}}MEMBAR.{{.*}} | ||
| ; SMXX-NOT: {{.*}}ST.E{{.*\[}}[[ATOM_ADDR]]{{(\.64)?(\+0x[0-9a-f]+)?\].*}} | ||
| ; SMXX-NOT: {{.*}}ATOM.{{.*}} | ||
| ; BLOCK: {{.*}}ST.E.128.STRONG.{{CTA|SM}} {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}}, R8{{.*}} | ||
| ; NON_BLOCK: {{.*}}ST.E.128.STRONG.[[SASS_SCOPE]] {{.*\[}}[[ATOM_ADDR]]{{(\.64)?\].*}}, R8{{.*}} | ||
| ; SMXX-NOT: {{.*}}ST.E{{.*\[}}[[ATOM_ADDR]]{{(\.64)?(\+0x[0-9a-f]+)?\].*}} | ||
| ; SMXX-NOT: {{.*}}CCTL.IVALL{{.*}} | ||
| ; SMXX-NEXT: {{.*}}RET.ABS.NODEC{{.*}} |
There was a problem hiding this comment.
🎯 Functional Correctness | 🟠 Major | ⚡ Quick win
important: Add an SM70+ gate to this test. libcudacxx_codegen_add_sass_tests registers this file for every selected architecture. This file has only SMXX checks. An SM below 70 cannot support the required 128-bit atomic_ref::store operation. Use the same SM70+ gate as the equivalent non-volatile 128-bit load/store tests.
Based on learnings: 128-bit atomic_ref load and store SASS tests must be gated at SM70+.
Source: Learnings
⏱️ CCCL compile-time benchmark comparison: Public headers compile-time benchResult: 0 regression row(s), 3 improvement row(s) above threshold.
Artifacts: reports and traces Direct file processing
🟢 Direct file processing — Improvements
|
😬 CI Workflow Results🟥 Finished in 2h 55m: Pass: 95%/597 | Total: 21d 18h | Max: 2h 55m | Hits: 29%/2350972See results here. AI failure analysis1. CTK 12.0 MSVC 14.39 C++17 build failure has no captured diagnostic · 1 jobExplanation: The build job failed in step 4, but no job log was collected. The earliest actionable error and developer impact cannot be determined from metadata alone. Root cause: Unknown because `job-94926181885.log` is absent; compiler or build-system diagnostics are required. Suggested next steps: Collect the complete job log and inspect the first error in step 4. Reproduce narrowly with the corresponding CTK 12.0, MSVC 14.39, C++17 libcudacxx build preset. Copy this prompt into a coding agentJobs: 2. CTK 12.0 MSVC 14.39 C++20 build failure has no captured diagnostic · 1 jobExplanation: The build job failed in step 4, but no job log was collected. Its failure mechanism cannot safely be equated with another matrix entry. Root cause: Unknown because `job-94926181894.log` is absent; compiler or build-system diagnostics are required. Suggested next steps: Collect the complete job log and inspect the first error in step 4. Reproduce narrowly with the corresponding CTK 12.0, MSVC 14.39, C++20 libcudacxx build preset. Copy this prompt into a coding agentJobs: 3. CTK 13.3 MSVC 14.44 C++20 build failure has no captured diagnostic · 1 jobExplanation: The build job failed in step 4, but no job log was collected. Its decisive compiler or configuration signature is unavailable. Root cause: Unknown because `job-94926181924.log` is absent. Suggested next steps: Collect the complete job log and reproduce the first step-4 error with CTK 13.3, MSVC 14.44, and C++20. Copy this prompt into a coding agentJobs: 4. CTK 13.3 MSVC 14.50 C++20 build failure has no captured diagnostic · 1 jobExplanation: The build job failed in step 4, but no job log was collected. Metadata does not reveal whether compilation, configuration, or another build phase failed. Root cause: Unknown because `job-94926181929.log` is absent. Suggested next steps: Collect the complete job log and reproduce the first step-4 error with CTK 13.3, MSVC 14.50, and C++20. Copy this prompt into a coding agentJobs: 5. CTK 13.3 MSVC 14.50 C++17 build failure has no captured diagnostic · 1 jobExplanation: The build job failed in step 4, but no output was saved. It cannot reliably be grouped with the C++20 entry without its diagnostic. Root cause: Unknown because `job-94926181930.log` is absent. Suggested next steps: Collect the complete job log and reproduce the first step-4 error with CTK 13.3, MSVC 14.50, and C++17. Copy this prompt into a coding agentJobs: 6. CTK 13.3 MSVC 14.29 C++17 build failure has no captured diagnostic · 1 jobExplanation: The build job failed in step 4, but no job log was collected. The older host compiler may expose a compatibility issue, but no diagnostic confirms that. Root cause: Unknown because `job-94926181936.log` is absent. Suggested next steps: Collect the complete job log and reproduce the first step-4 error with CTK 13.3, MSVC 14.29, and C++17. Copy this prompt into a coding agentJobs: 7. CTK 12.0 MSVC 14.29 C++17 build failure has no captured diagnostic · 1 jobExplanation: The build job failed in step 4, but no job log was collected. Neither the failing target nor the compiler message is available. Root cause: Unknown because `job-94926181977.log` is absent. Suggested next steps: Collect the complete job log and reproduce the first step-4 error with CTK 12.0, MSVC 14.29, and C++17. Copy this prompt into a coding agentJobs: 8. CTK 12.0 Clang 14 C++17 build failure has no captured diagnostic · 1 jobExplanation: The Clang build failed in step 4, but no log was collected. There is no evidence identifying a language, CUDA compatibility, or configuration failure. Root cause: Unknown because `job-94926182651.log` is absent. Suggested next steps: Collect the complete log and reproduce the earliest error using CTK 12.0, Clang 14, and C++17. Copy this prompt into a coding agentJobs: 9. CTK 12.0 Clang 14 C++20 build failure has no captured diagnostic · 1 jobExplanation: The Clang build failed in step 4, but no log was collected. Its failure cannot safely be assumed equivalent to the C++17 entry. Root cause: Unknown because `job-94926182696.log` is absent. Suggested next steps: Collect the complete log and reproduce the earliest error using CTK 12.0, Clang 14, and C++20. Copy this prompt into a coding agentJobs: 10. CTK 12.0 GCC 7 C++17 build failure has no captured diagnostic · 1 jobExplanation: The GCC 7 build failed in step 4, but no output was collected. A compiler-compatibility regression is possible but unconfirmed. Root cause: Unknown because `job-94926183756.log` is absent. Suggested next steps: Collect the complete log and reproduce the first error with CTK 12.0, GCC 7, and C++17. Copy this prompt into a coding agentJobs: 11. CTK 12.0 GCC 12 C++20 build failure has no captured diagnostic · 1 jobExplanation: The GCC build failed in step 4, but no job log was collected. The failing source, target, and diagnostic are unavailable. Root cause: Unknown because `job-94926183762.log` is absent. Suggested next steps: Collect the complete log and reproduce the earliest error using CTK 12.0, GCC 12, and C++20. Copy this prompt into a coding agentJobs: 12. CTK 12.0 GCC 12 C++17 build failure has no captured diagnostic · 1 jobExplanation: The GCC build failed in step 4, but no job log was collected. Its relationship to the C++20 build failure cannot be established. Root cause: Unknown because `job-94926183888.log` is absent. Suggested next steps: Collect the complete log and reproduce the earliest error using CTK 12.0, GCC 12, and C++17. Copy this prompt into a coding agentJobs: 13. CTK 12.0 sm75 atomic SASS check failure has no captured diagnostic · 1 jobExplanation: The atomic code-generation FileCheck job failed in step 4, but no log was saved. The generated instruction mismatch, failed compilation, and invoked test case are unknown. Root cause: Unknown because `job-94926183782.log` is absent. Suggested next steps: Collect the complete log, identify the first failing generated SASS test and FileCheck directive, then rerun only that CTK 12.0/GCC 12/sm75 case. Copy this prompt into a coding agentJobs: 14. CTK 12.0 sm80 atomic SASS check failure has no captured diagnostic · 1 jobExplanation: The atomic code-generation FileCheck job failed in step 4, but no log was saved. No specific SASS pattern or source test can be identified. Root cause: Unknown because `job-94926183818.log` is absent. Suggested next steps: Collect the complete log and rerun only the first failing CTK 12.0/GCC 12/sm80 SASS test. Copy this prompt into a coding agentJobs: 15. CTK 12.9 sm75 atomic SASS check failure has no captured diagnostic · 1 jobExplanation: The atomic SASS FileCheck job failed in step 4, but no job output was collected. The decisive mismatch is unavailable. Root cause: Unknown because `job-94926183860.log` is absent. Suggested next steps: Collect the complete log and rerun only the first failing CTK 12.9/GCC 14/sm75 SASS test. Copy this prompt into a coding agentJobs: 16. CTK 12.9 sm80 atomic SASS check failure has no captured diagnostic · 1 jobExplanation: The atomic SASS FileCheck job failed in step 4, but no job output was collected. It cannot be grouped with another architecture without the mismatched pattern. Root cause: Unknown because `job-94926183812.log` is absent. Suggested next steps: Collect the complete log and rerun only the first failing CTK 12.9/GCC 14/sm80 SASS test. Copy this prompt into a coding agentJobs: 17. CTK 12.9 sm100 atomic SASS check failure has no captured diagnostic · 1 jobExplanation: The atomic SASS FileCheck job failed in step 4, but no job output was collected. The failing instruction expectation is unknown. Root cause: Unknown because `job-94926183923.log` is absent. Suggested next steps: Collect the complete log and rerun only the first failing CTK 12.9/GCC 14/sm100 SASS test. Copy this prompt into a coding agentJobs: 18. CTK 12.9 sm120 atomic SASS check failure has no captured diagnostic · 1 jobExplanation: The atomic SASS FileCheck job failed in step 4, but no log was collected. The failure could occur during architecture support checks, compilation, disassembly, or FileCheck. Root cause: Unknown because `job-94926183861.log` is absent. Suggested next steps: Collect the complete log and reproduce the earliest CTK 12.9/GCC 14/sm120 failure. Copy this prompt into a coding agentJobs: 19. CTK 13.0 sm75 atomic SASS check failure has no captured diagnostic · 1 jobExplanation: The atomic SASS FileCheck job failed in step 4, but no log was saved. No particular test or generated-code mismatch can be identified. Root cause: Unknown because `job-94926183910.log` is absent. Suggested next steps: Collect the complete log and rerun only the first failing CTK 13.0/GCC 15/sm75 SASS case. Copy this prompt into a coding agentJobs: 20. CTK 13.0 sm80 atomic SASS check failure has no captured diagnostic · 1 jobExplanation: The atomic SASS FileCheck job failed in step 4, but no log was saved. Its decisive generated-code signature is unavailable. Root cause: Unknown because `job-94926183898.log` is absent. Suggested next steps: Collect the complete log and rerun only the first failing CTK 13.0/GCC 15/sm80 SASS case. Copy this prompt into a coding agentJobs: 21. CTK 13.0 sm100 atomic SASS check failure has no captured diagnostic · 1 jobExplanation: The atomic SASS FileCheck job failed in step 4, but no log was saved. The failing SASS test and mismatch are unknown. Root cause: Unknown because `job-94926183912.log` is absent. Suggested next steps: Collect the complete log and rerun only the first failing CTK 13.0/GCC 15/sm100 SASS case. Copy this prompt into a coding agentJobs: 22. CTK 13.0 sm120 atomic SASS check failure has no captured diagnostic · 1 jobExplanation: The atomic SASS FileCheck job failed in step 4, but no output was collected. Architecture setup, compilation, and FileCheck failures cannot be distinguished. Root cause: Unknown because `job-94926183979.log` is absent. Suggested next steps: Collect the complete log and reproduce the earliest CTK 13.0/GCC 15/sm120 error. Copy this prompt into a coding agentJobs: 23. CTK 13.3 sm75 atomic SASS check failure has no captured diagnostic · 1 jobExplanation: The atomic SASS FileCheck job failed in step 4, but no job output was collected. The failing test and SASS expectation are unavailable. Root cause: Unknown because `job-94926191341.log` is absent. Suggested next steps: Collect the complete log and rerun only the first failing CTK 13.3/GCC 15/sm75 SASS case. Copy this prompt into a coding agentJobs: 24. CTK 13.3 sm80 atomic SASS check failure has no captured diagnostic · 1 jobExplanation: The atomic SASS FileCheck job failed in step 4, but no job output was collected. Its failure signature cannot be compared with the other architecture entries. Root cause: Unknown because `job-94926191396.log` is absent. Suggested next steps: Collect the complete log and rerun only the first failing CTK 13.3/GCC 15/sm80 SASS case. Copy this prompt into a coding agentJobs: 25. CTK 13.3 sm100 atomic SASS check failure has no captured diagnostic · 1 jobExplanation: The atomic SASS FileCheck job failed in step 4, but no job output was collected. No decisive mismatch or failing source is available. Root cause: Unknown because `job-94926191361.log` is absent. Suggested next steps: Collect the complete log and rerun only the first failing CTK 13.3/GCC 15/sm100 SASS case. Copy this prompt into a coding agentJobs: 26. CTK 13.3 sm120 atomic SASS check failure has no captured diagnostic · 1 jobExplanation: The atomic SASS FileCheck job failed in step 4, but no job output was collected. The metadata cannot identify whether architecture configuration or a test assertion failed. Root cause: Unknown because `job-94926191343.log` is absent. Suggested next steps: Collect the complete log and reproduce the earliest CTK 13.3/GCC 15/sm120 error. Copy this prompt into a coding agentJobs: 27. CTK 13.3 Clang 21 T4 runtime test failure has no captured diagnostic · 1 jobExplanation: The libcudacxx test job failed in step 4, but no log was collected. The failing test name, assertion, crash, or infrastructure error is unavailable. Root cause: Unknown because `job-94936921218.log` is absent; runtime test output is required. Suggested next steps: Collect the complete log and rerun only the earliest failing test with CTK 13.3, Clang 21, C++20, and a T4 GPU. Copy this prompt into a coding agentJobs: |
1 participant
Description
Resolves #10666.
This PR introduces the third SASS FileCheck suite for the atomic codegen, which includes a wide coverage of volatile and non-volatile codegen results for arithmetic operations: fetch_add, fetch_sub, increment/decrement operators, fetch_min, and fetch_max.
This PR temporarily also includes the commits from #10722, to facilitate CI testing. The commits relevant to this PR are the four last ones, starting with "libcu++: add SASS tests for atomic arithmetic operations.".
This PR temporarily also includes the commits from #10828.
Checklist