refactor(gpu): finish the cuda::std::complex migration (Cpr/Mod/cuOuter) and delete cucomplex_arithmetic#1019
Conversation
Stage 1 of completing #1004's GPU complex migration. Add/Sub/Mul/Div already route through the new cuda::std::complex typed dispatch (type_promote_gpu_t + TO(lhs) op TO(rhs) kernels); Cpr and Mod were the last consumers of the legacy cuAri_ii table (which forwards to the old cuComplex kernels via cuCadd/cuCmul and cucomplex_arithmetic's hand-rolled operators). Add cuCpr_dispatch (mirrors cuMul_dispatch) computing the comparison in the promoted type: out(bool) = (TO(lhs) == TO(rhs)) with TO = type_promote_gpu(TL, TR). Casting both operands to TO first is required because cuda::std::complex has no mixed-type operators (e.g. `complex<float> == double` is ill-formed), and it reproduces the old make_cuDoubleComplex(v, 0) widening for complex-vs- real comparisons. Rewire all 13 Cpr.cpp GPU call sites from lii.cuAri_ii[L][R](..., 4) to cuCpr_dispatch(...). The legacy cuCpr kernels stay for now (still reachable via the shared cuArithmetic_internal forwarder that Mod uses); they are removed once Mod is migrated and cuAri_ii is dropped. nvcc front-end verified: cuCpr_dispatch.cu, Cpr.cpp (GPU + CPU) all compile clean. Not runtime-tested (no GPU CI runner); mirrors the proven cuMul path. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Stage 2. Add cuMod_dispatch (mirrors cuMul_dispatch), computing modulo in the
promoted type TO = type_promote_gpu(TL, TR): integral -> TO(lhs) % TO(rhs),
floating -> fmod/fmodf, matching the CPU ModOp<T>. Complex operands are
rejected at the host dispatch ("[cuMod] Cannot mod complex numbers"), exactly
as the old cuMod_internal_* functions and the CPU path do, so the kernel's
complex branch is an unreachable if-constexpr stub. Rewire all 25 Mod.cpp GPU
call sites (Tensor-Tensor, scalar-Tensor, Tensor-scalar) from
lii.cuAri_ii[L][R](..., 5) to cuMod_dispatch(...).
With Cpr (stage 1) and Mod migrated, nothing calls the legacy cuAri_ii table
anymore; the old cuComplex kernels + cuArithmetic_internal layer become dead
and are removed in the cleanup stage.
nvcc front-end verified: cuMod_dispatch.cu, Mod.cpp (GPU + CPU) compile clean.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Stage 3. Switch cuOuter's generated dispatch instantiations from the cuComplex ABI types (cuDoubleComplex/cuFloatComplex) to cuda::std::complex (cytnx_cuda_complex128/64), and cast both operands to the output type T1 in cuOuter_kernel (T1(val) * T1(ptr)) since cuda::std::complex has no mixed-type operator*. This drops cuOuter's dependency on utils/cucomplex_arithmetic.hpp (removed the include). The 4 hand-added cuGer_internal (cuBLAS ger) fast paths for the cd/cf/d/f diagonals are preserved. Also update the generator's master type list (boostraps/cy_type.py cuDTYPES_FULL) to cuda::std::complex for consistency, so a future regeneration emits the correct types (the checked-in file additionally hand-uses cuGer for 4 BLAS-able cases, which the generator does not emit). nvcc front-end verified: cuOuter_internal.cu compiles clean. cuKron was already fully on cuda::std::complex (from #1004). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…kernels
Stage 4 (cleanup). With Cpr and Mod migrated to their typed dispatch paths and
cuOuter on cuda::std::complex, the hand-rolled cuComplex operator overloads in
utils/cucomplex_arithmetic.{hpp,cu} (operator== / operator* for mixed
cuComplex/real) have exactly one remaining user: the dead legacy
cuCpr_internal kernels, reachable only through the cuArithmetic_internal
forwarders' type==4 branch (Cpr has no in-place variant, and out-of-place Cpr
now uses cuCpr_dispatch, so type==4 is never dispatched).
- Strip the 120 dead `else if (type == 4) -> cuCpr_internal_*` branches from
cuArithmetic_internal.cu (type 4 now falls through to the existing else).
- Delete cuCpr_internal.cu and its old per-type declarations (keep
cuCpr_dispatch); drop it from CMakeLists.
- Remove the now-unused cucomplex_arithmetic include from cuTrace_internal.cu,
utils/utils.hpp, and utils_internal_interface.hpp.
- Delete include/utils/cucomplex_arithmetic.hpp + src/utils/cucomplex_arithmetic.cu
and their CMake entry.
The remaining legacy arithmetic path (cuAri_ii -> cuArithmetic_internal ->
cuAdd/Mul/Sub/Div/Mod_internal, using cuComplex.h builtins cuCadd/cuCmul, not
these operators) stays: it is still used by the in-place ops
iAdd/iSub/iMul/iDiv. That layer can be removed once those are migrated too.
nvcc front-end verified: cuArithmetic_internal, cuMod/cuAdd/cuMul/cuOuter
internals, Cpr.cpp (GPU+CPU) all compile clean; no cucomplex_arithmetic
references remain in the tree.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
There was a problem hiding this comment.
Code Review
This pull request refactors GPU complex arithmetic by removing the custom cucomplex_arithmetic utilities and transitioning to cuda::std::complex (via cytnx_cuda_complex128/64). It introduces cuCpr_dispatch.cu and cuMod_dispatch.cu to handle GPU dispatching for comparison and modulo operations. The reviewer identified critical issues in both dispatch files where the host code attempts to write directly to device-allocated pointers (m_accu_shape, m_old_accu_shapeL, and m_old_accu_shapeR), which would result in segmentation faults. It is recommended to compute these values on the host first and then copy them to the device using cudaMemcpy.
Important
The consumer version of Gemini Code Assist on GitHub is being sunset. Starting June 18, 2026, new organization installations will be blocked, and all code review activity will officially cease on July 17, 2026.
For more details on the timeline and next steps, please review the Help Documentation.
Re: @gemini-code-assist review — both "critical" findings are false positivesGemini flagged the non-contiguous branches in
No changes needed for either finding. (Replied inline on both threads as well.) |
|
Thank you for the detailed explanation and for verifying the behavior on hardware. Your points regarding the use of managed memory ( |
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## master #1019 +/- ##
=======================================
Coverage 55.62% 55.62%
=======================================
Files 230 230
Lines 33099 33099
Branches 77 77
=======================================
Hits 18412 18412
Misses 14662 14662
Partials 25 25
Flags with carried forward coverage won't be shown. Click here to find out more.
Continue to review full report in Codecov by Harness.
🚀 New features to boost your workflow:
|
| // type (imaginary part 0), exactly as the old make_cuDoubleComplex(v, 0) | ||
| // path did. | ||
| template <typename TO, typename TL, typename TR> | ||
| __device__ inline cytnx_bool CuCprDispatchOp(const TL &lhs, const TR &rhs) { |
There was a problem hiding this comment.
I haven't known that cpr means compare until now. cmp, cp, comp are more common in software. We can consider to rename it. Or maybe, "equal" is a more accurate for this function.
| const TL *_Lin = reinterpret_cast<const TL *>(Lin->data()); | ||
| const TR *_Rin = reinterpret_cast<const TR *>(Rin->data()); | ||
|
|
||
| cytnx_uint32 NBlocks = len / 512; |
There was a problem hiding this comment.
512 should be a kConstant with a proper name.
| } | ||
|
|
||
| template <typename TO, typename TL, typename TR> | ||
| __global__ void cuCpr_dispatch_lconst_kernel(cytnx_bool *out, const TL lhs, |
There was a problem hiding this comment.
Do we really need both lconst, and rconst?
Should we follow naming rule while refactoring the whole file?
| const unsigned long long &len, | ||
| const std::vector<cytnx_uint64> &shape, | ||
| const std::vector<cytnx_uint64> &invmapper_L, | ||
| const std::vector<cytnx_uint64> &invmapper_R) { |
There was a problem hiding this comment.
If invmapper is needed here, the input type may be tensor instead of storage
| if (len % 512) NBlocks += 1; | ||
|
|
||
| if (Lin->size() == 1 and Rin->size() == 1) { | ||
| cuCpr_dispatch_constconst_kernel<TO><<<NBlocks, 512>>>(_out, _Lin[0], len, _Rin[0]); |
There was a problem hiding this comment.
If len is always 1 here, this is a trivial action, and constconst_kernel is not needed.
There was a problem hiding this comment.
well, it certainly doesn't need 512 threads... but since _out is also on GPU it is probably better to launch GPU kernel for this, versus copying the parameters to host memory then back to GPU. But I don't understand why this needs to be a special case.
|
Reiterating the comment here since it is a reply to a 'resolved' comment: It is a false positive in some sense - it will actually work without segfaulting, but it is terrible style and very slow. All of these arrays are very short, it would be far better to assemble the data into a small fixed size structure and pass it directly into the kernel as a parameter, similar style as the Writing these tiny arrays through unified memory is particularly bad: the first write will need to copy the entire memory page from GPU to CPU (probably every array ends up on a different page), then copy it back once the GPU kernel needs the memory. |
|
On second thoughts, it is probably not worth fixing the previous comment but just merge this as-is. I'll prepare a PR for #1003 anyway. This PR is still needed for the other fixes to remove the cuFloat types. |
IvanaGyro
left a comment
There was a problem hiding this comment.
I see the kernel implementatipn in this PR is just a copy from the old code. We can fix them in future.
Summary
Completes #1004's migration of the GPU element-wise ops to
cuda::std::complexand deletes the hand-rolledutils/cucomplex_arithmetic.{hpp,cu}.#1004 introduced the
cuda::std::complexinfrastructure (Type_list_gpu,gpu_ptr()/gpu_ptr_as<T>(),type_promote_gpu_t) and migrated Add/Sub/Mul/Div (and cuKron) to a typed dispatch path — but left Cpr, Mod, and cuOuter on the legacycuComplexpath, which still depended on the mixed-typeoperator==/operator*overloads incucomplex_arithmetic. Those overloads are also the ones flagged as wrong-by-promotion in the #995 / #994 discussion (cuFloatComplex * double -> cuFloatComplex, dropping precision after #858 madeComplexFloat*Double -> ComplexDouble); they can't be fixed in isolation because the generated GPU dispatch tables rely on the truncating behavior. Routing everything through thecuda::std::complextyped path (which promotes correctly and casts to the ABI only at cuBLAS/cuSOLVER boundaries) is the real fix Ian suggested on #995.Supersedes the narrow Kron-only patch in #995.
What changed (staged, one op per commit)
cuCpr_dispatch(mirrorscuMul_dispatch):out(bool) = (TO(lhs) == TO(rhs)),TO = type_promote_gpu(TL,TR). Casting toTOfirst is required becausecuda::std::complexhas no mixed-type operators (complex<float> == doubleis ill-formed) and reproduces the oldmake_cuDoubleComplex(v,0)widening. 13Cpr.cppcall sites rewired offcuAri_ii.cuMod_dispatch: per-category (%for integral,fmod/fmodffor floating, complex rejected at host — matches the CPUModOp<T>). 25Mod.cppcall sites rewired.cuda::std::complextypes +TO(a)*TO(b)cast kernel; the 4 hand-addedcuGer(cuBLAS ger) fast paths preserved;boostraps/cy_type.pycuDTYPES_FULLupdated for consistency.cucomplex_arithmetic.{hpp,cu}+ the dead legacycuCpr_internalkernels + the 120 deadtype==4branches in thecuArithmetic_internalforwarders (Cpr has no in-place variant, sotype==4is never dispatched). Net −3,630 lines.Validation
Built + tested locally on real GPU hardware (CUDA 13.0, RTX 4070 Ti, sm_89) — CI has no CUDA runner, so this path is not exercised by CI:
libcytnx.acompiles + device-links + host-links — confirms the deletions/forwarder edits leave no dangling symbols (what a front-end-only check can't catch).*cpr*GPU tests: 80/80 passed (cuCpr_dispatch).*mod*GPU tests: 49/49 passed (cuMod_dispatch).iadd/isub/imul/idiv: 8/8 passed — regression check that strippingtype==4from the sharedcuArithmetic_internalforwarders didn't break the live in-place path (which still usescuAri_ii).cuOuter has no GPU runtime test in-tree, but it compiles + links.
Notes / follow-ups (not in this PR)
cuAri_ii → cuArithmetic_internal → cuAdd/Mul/Sub/Div/Mod_internallayer stays: it's still used by the in-place ops (iAdd/iSub/iMul/iDiv), via cuComplex.h builtins (cuCadd/cuCmul), not the deleted operators. Migrating those four would let the whole legacy layer be deleted.gpu_test_main's host-compiled test.cppfiles includeType.hpp→<cuda/std/complex>but thetests/gpuCMake target lacks the CCCL include dir (-I.../include/cccl) that thecytnxlibrary target has, so the GPU test binary doesn't build out-of-the-box under CUDA 13. Worth fixing intests/gpu/CMakeLists.txt.🤖 Generated with Claude Code