fix(linalg): compute GPU Kron product in the promoted output dtype

[pcchen]

What

Fix a GPU build break in cuKron_kernel: compute the Kron product in the promoted output dtype instead of relying on the raw operator* result type.

Why

cuKron_kernel did out[i] = Lin[x] * Rin[y]. The cucomplex operator* overloads follow the left operand's precision, so cuFloatComplex * double -> cuFloatComplex. After #858 made mixed complex/real pairs promote by precision (ComplexFloat + Double -> ComplexDouble), the output type is cuDoubleComplex, and assigning the cuFloatComplex product to it fails to compile:

src/backend/linalg_internal_gpu/cuKron_internal.cuh:49: error: no operator "=" matches these operands
  operand types are: cuDoubleComplex = cuFloatComplex

This broke the CUDA build for the (cuFloatComplex, cytnx_double) instantiation (and its mirror). CI didn't catch it because it does not build the GPU path; the CPU Kron_general handles the conversion and compiled fine.

How

Introduce a cuKron_mul<TO> helper that promotes both operands to TO before multiplying (cuCmul/cuCmulf for complex TO, static_cast<TO> otherwise), mirroring the existing cuDiv_dispatch.cu pattern. This makes the result assignable for every promoted (TO, TL, TR) triple.

Testing

Built the full cytnx target with CUDA + cuTENSOR + cuQuantum (CUDA 12.9, sm_70 / V100) — compiles and links. Ran Kron(ComplexFloat, Double) on the GPU: output dtype is ComplexDouble and the values match the CPU reference and hand-computed expected results.

Note: this class of regression would be caught by a minimal GPU compile job in CI (no GPU hardware needed to compile the templates). Worth considering as a follow-up.

Fixes #994

🤖 Generated with Claude Code

Posted by Claude Code on behalf of @pcchen

[gemini-code-assist]

Code Review

This pull request introduces helper functions (cuKron_to_cd, cuKron_to_cf, and cuKron_mul) in cuKron_internal.cuh to correctly handle type promotion and multiplication of mixed complex/real types on CUDA devices, updating cuKron_kernel to use cuKron_mul. Feedback is provided to use the standard CUDA utility function cuComplexDoubleToFloat instead of manually reconstructing the float complex number in cuKron_to_cf for better consistency.

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.

[gemini-code-assist]

For consistency with cuComplexFloatToDouble used in cuKron_to_cd, you can use the standard CUDA utility function cuComplexDoubleToFloat from <cuComplex.h> instead of manually reconstructing the float complex number.

    __device__ inline cuComplex cuKron_to_cf(const cuDoubleComplex &v) {
      return cuComplexDoubleToFloat(v);
    }

[codecov]

Codecov Report

✅ All modified and coverable lines are covered by tests.
✅ Project coverage is 30.88%. Comparing base (d02cb29) to head (eb0fe1f).
✅ All tests successful. No failed tests found.

Additional details and impacted files

@@           Coverage Diff           @@
##           master     #995   +/-   ##
=======================================
  Coverage   30.88%   30.88%           
=======================================
  Files         229      229           
  Lines       34758    34758           
  Branches    14409    14409           
=======================================
  Hits        10734    10734           
  Misses      16720    16720           
  Partials     7304     7304           

Flag	Coverage Δ
cpp	30.39% <ø> (ø)
python	59.41% <ø> (ø)

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

Components	Coverage Δ
C++ backend	32.21% <ø> (ø)
Python bindings	17.28% <ø> (ø)
Python package	59.41% <ø> (ø)

---

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 d02cb29...eb0fe1f. 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.

[pcchen]

Priority note: master's GPU build currently does not compile without this — Kron.cpp fails to build for the (cuFloatComplex, cytnx_double) instantiation after #858 changed mixed complex/real promotion (see #994). So this effectively blocks building/testing any GPU code on master, including the GPU paths in the scalar-in-place cluster (#980 / #992 / #937). Suggest prioritizing it.

Independently verified: built the full cytnx target with CUDA (CUDA 12.9, sm_70 / Tesla V100) and ran Kron(ComplexFloat, Double) on-device — output dtype is ComplexDouble and values match the CPU reference. Also confirmed the pre-fix build fails at Kron.cpp.o and this change makes it compile + link + run.

Note that CI has no CUDA runner, so none of this is exercised by CI — worth keeping in mind for the whole GPU cluster.

Posted by Claude Code on behalf of @pcchen

[ianmccul]

From discussion in #999, the actual root cause here is the Cytnx declarations in include/utils/cucomplex_arithmetic.hpp, definitions in src/utils/cucomplex_arithmetic.cu, that are incorrect with respect to any sane promotion rules:

 __host__ __device__
cuFloatComplex operator*(const cuFloatComplex &ln, const cytnx_double &rn) {
  return cuCmulf(make_cuFloatComplex(rn, 0), ln);
}

__host__ __device__
cuFloatComplex operator*(const cytnx_double &rn, const cuFloatComplex &ln) {
  return cuCmulf(make_cuFloatComplex(rn, 0), ln);
}

It is better to fix these, rather than patch Kron only, since they are used elsewhere; cuOuter, iMul, etc.

Even better would be to replace the old cuXXXX types with cuda::std::complex, which works with CUDA >= 11.3.

[yingjerkao]

Superseded — closing. master already computes the Kron product in the promoted output dtype via out[i] = TO(Lin[x]) * TO(Rin[y]) (landed in #1004, commit 7bbd814, as part of the cuda::std::complex migration), and #1019 removed the cucomplex_arithmetic overloads this change worked around. The build break this fixed (#994/#999) is resolved on master, so this standalone fix is no longer needed. Thanks!