Skip to content

refactor(gpu): finish the cuda::std::complex migration (Cpr/Mod/cuOuter) and delete cucomplex_arithmetic#1019

Merged
ianmccul merged 5 commits into
masterfrom
refactor/gpu-cuda-std-complex
Jul 10, 2026
Merged

refactor(gpu): finish the cuda::std::complex migration (Cpr/Mod/cuOuter) and delete cucomplex_arithmetic#1019
ianmccul merged 5 commits into
masterfrom
refactor/gpu-cuda-std-complex

Conversation

@yingjerkao

Copy link
Copy Markdown
Collaborator

Summary

Completes #1004's migration of the GPU element-wise ops to cuda::std::complex and deletes the hand-rolled utils/cucomplex_arithmetic.{hpp,cu}.

#1004 introduced the cuda::std::complex infrastructure (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 legacy cuComplex path, which still depended on the mixed-type operator==/operator* overloads in cucomplex_arithmetic. Those overloads are also the ones flagged as wrong-by-promotion in the #995 / #994 discussion (cuFloatComplex * double -> cuFloatComplex, dropping precision after #858 made ComplexFloat*Double -> ComplexDouble); they can't be fixed in isolation because the generated GPU dispatch tables rely on the truncating behavior. Routing everything through the cuda::std::complex typed 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)

  1. Cpr → new cuCpr_dispatch (mirrors cuMul_dispatch): out(bool) = (TO(lhs) == TO(rhs)), TO = type_promote_gpu(TL,TR). Casting to TO first is required because cuda::std::complex has no mixed-type operators (complex<float> == double is ill-formed) and reproduces the old make_cuDoubleComplex(v,0) widening. 13 Cpr.cpp call sites rewired off cuAri_ii.
  2. ModcuMod_dispatch: per-category (% for integral, fmod/fmodf for floating, complex rejected at host — matches the CPU ModOp<T>). 25 Mod.cpp call sites rewired.
  3. cuOutercuda::std::complex types + TO(a)*TO(b) cast kernel; the 4 hand-added cuGer (cuBLAS ger) fast paths preserved; boostraps/cy_type.py cuDTYPES_FULL updated for consistency.
  4. Cleanup → delete cucomplex_arithmetic.{hpp,cu} + the dead legacy cuCpr_internal kernels + the 120 dead type==4 branches in the cuArithmetic_internal forwarders (Cpr has no in-place variant, so type==4 is 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.a compiles + 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).
  • In-place iadd/isub/imul/idiv: 8/8 passed — regression check that stripping type==4 from the shared cuArithmetic_internal forwarders didn't break the live in-place path (which still uses cuAri_ii).

cuOuter has no GPU runtime test in-tree, but it compiles + links.

Notes / follow-ups (not in this PR)

  • The legacy cuAri_ii → cuArithmetic_internal → cuAdd/Mul/Sub/Div/Mod_internal layer 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.
  • Test-build config gap (from Use cuda::std::complex for typed GPU complex dispatch #1004): gpu_test_main's host-compiled test .cpp files include Type.hpp<cuda/std/complex> but the tests/gpu CMake target lacks the CCCL include dir (-I.../include/cccl) that the cytnx library target has, so the GPU test binary doesn't build out-of-the-box under CUDA 13. Worth fixing in tests/gpu/CMakeLists.txt.

🤖 Generated with Claude Code

yingjerkao and others added 4 commits July 9, 2026 10:11
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>

@gemini-code-assist gemini-code-assist Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Comment thread src/backend/linalg_internal_gpu/cuCpr_dispatch.cu
Comment thread src/backend/linalg_internal_gpu/cuMod_dispatch.cu
@yingjerkao
yingjerkao requested a review from IvanaGyro July 9, 2026 03:20
@yingjerkao

Copy link
Copy Markdown
Collaborator Author

Re: @gemini-code-assist review — both "critical" findings are false positives

Gemini flagged the non-contiguous branches in cuCpr_dispatch.cu and cuMod_dispatch.cu as host writes to device-only pointers (m_accu_shape / m_old_accu_shapeL / m_old_accu_shapeR) that would segfault. They won't:

  • cuCalloc_gpu / cuMalloc_gpu allocate managed (unified) memory via cudaMallocManaged (src/backend/utils_internal_gpu/cuAlloc_gpu.cu:9-19), not device-only memory — so writing those arrays from the host is valid; the pages migrate to the device on the kernel launch.
  • This is the identical pattern used by the already-merged cuAdd_dispatch / cuMul_dispatch from Use cuda::std::complex for typed GPU complex dispatch #1004, which I mirrored deliberately.
  • Verified on real hardware (CUDA 13.0, RTX 4070 Ti): *cpr* (80/80) and *mod* (49/49) GPU tests pass, including the non-contiguous multi-dim shapes ({16,16,16,16}, {64,64,64}, …) that exercise exactly this shape.size() > 0 branch — a host write to device-only memory would have crashed there.

No changes needed for either finding. (Replied inline on both threads as well.)

@gemini-code-assist

Copy link
Copy Markdown
Contributor

Thank you for the detailed explanation and for verifying the behavior on hardware. Your points regarding the use of managed memory (cudaMallocManaged) and the consistency with existing patterns in cuAdd_dispatch / cuMul_dispatch are well-taken. Since the memory is managed, the host-to-device migration is handled correctly by the CUDA runtime upon kernel launch, and your testing confirms that the non-contiguous shape branches function as expected without memory access violations. I have noted your clarification.

@codecov

codecov Bot commented Jul 9, 2026

Copy link
Copy Markdown

Codecov Report

✅ All modified and coverable lines are covered by tests.
✅ Project coverage is 55.62%. Comparing base (3268385) to head (fffbde5).
⚠️ Report is 1 commits behind head on master.
✅ All tests successful. No failed tests found.

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           
Flag Coverage Δ
cpp 55.55% <ø> (ø)
python 61.84% <ø> (ø)

Flags with carried forward coverage won't be shown. Click here to find out more.

Components Coverage Δ
C++ backend 53.30% <ø> (ø)
Python bindings 70.35% <ø> (ø)
Python package 61.84% <ø> (ø)
Files with missing lines Coverage Δ
src/linalg/Cpr.cpp 14.18% <ø> (ø)
src/linalg/Mod.cpp 72.70% <ø> (ø)

Continue to review full report in Codecov by Harness.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update 3268385...fffbde5. Read the comment docs.

🚀 New features to boost your workflow:
  • 📦 JS Bundle Analysis: Save yourself from yourself by tracking and limiting bundle sizes in JS merges.

@IvanaGyro IvanaGyro left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The problems on cpr also on mod. I reviewed on my phone. The screen is so small that I may misread something.

Do the problems I addressed comes from the original implementation?

// 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) {

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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;

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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,

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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) {

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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]);

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

If len is always 1 here, this is a trivial action, and constconst_kernel is not needed.

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

@ianmccul

ianmccul commented Jul 9, 2026

Copy link
Copy Markdown
Collaborator

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 TraceLayout structure in src/backend/linalg_internal_gpu/cuTrace_internal.cu .

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.

@ianmccul

ianmccul commented Jul 9, 2026

Copy link
Copy Markdown
Collaborator

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 IvanaGyro left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I see the kernel implementatipn in this PR is just a copy from the old code. We can fix them in future.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants