refactor(linalg): typed GPU dispatch for unary ops Abs/Exp/Pow/Conj (#1003 step 11)#1094
refactor(linalg): typed GPU dispatch for unary ops Abs/Exp/Pow/Conj (#1003 step 11)#1094yingjerkao wants to merge 5 commits into
Conversation
![gemini-code-assist[bot]](https://avatars.githubusercontent.com/in/956858?s=60&v=4){kind=small}
There was a problem hiding this comment.
Code Review
This pull request introduces a typed GPU dispatch for the elementwise Abs operation (cuAbs_dispatch.cu), replacing the legacy lookup table, and updates Abs.cpp and Abs_.cpp to use it. The review feedback highlights several key improvements: renaming CuAbsValue to cu_abs_value to follow the snake_case style guide rule, using cuda::std::abs to correctly handle -0.0 for floating-point types, casting blockIdx.x to cytnx_uint64 to prevent integer overflow in the GPU kernel, passing the scalar Nelem by value instead of const& per the style guide, and simplifying the block count calculation.
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.
| template <typename TIn> | ||
| __device__ inline auto CuAbsValue(const TIn &x) { | ||
| if constexpr (std::is_same_v<TIn, cytnx_cuda_complex128> || | ||
| std::is_same_v<TIn, cytnx_cuda_complex64>) { | ||
| return cuda::std::abs(x); // magnitude: complex<T> -> T | ||
| } else if constexpr (std::is_signed_v<TIn>) { | ||
| return x < TIn(0) ? static_cast<TIn>(-x) : x; // fabs / abs for signed real & integer | ||
| } else { | ||
| return x; // unsigned / bool: abs is the identity | ||
| } | ||
| } |
There was a problem hiding this comment.
The helper function CuAbsValue uses PascalCase, which violates the repository style guide's requirement that new identifiers use snake_case. Additionally, the ternary operator x < TIn(0) ? static_cast<TIn>(-x) : x does not handle -0.0 correctly for floating-point types (since -0.0 < 0.0 is false in IEEE 754, returning -0.0 instead of 0.0).
We should rename it to cu_abs_value and use cuda::std::abs for floating-point types to ensure correctness.
template <typename TIn>
__device__ inline auto cu_abs_value(const TIn &x) {
if constexpr (std::is_same_v<TIn, cytnx_cuda_complex128> ||
std::is_same_v<TIn, cytnx_cuda_complex64>) {
return cuda::std::abs(x); // magnitude: complex<T> -> T
} else if constexpr (std::is_floating_point_v<TIn>) {
return cuda::std::abs(x); // handles -0.0 correctly
} else if constexpr (std::is_signed_v<TIn>) {
return x < TIn(0) ? static_cast<TIn>(-x) : x; // abs for signed integer
} else {
return x; // unsigned / bool: abs is the identity
}
}
References
- New identifiers use snake_case; a leading underscore is prohibited in new code. (link)
There was a problem hiding this comment.
Why not pass all type except !std::is_signed_v<TIn> to cuda::std::abs?
| template <typename TOut, typename TIn> | ||
| __global__ void cuAbs_dispatch_kernel(TOut *out, const TIn *in, const cytnx_uint64 Nelem) { | ||
| const cytnx_uint64 idx = blockIdx.x * blockDim.x + threadIdx.x; | ||
| if (idx < Nelem) out[idx] = static_cast<TOut>(CuAbsValue(in[idx])); | ||
| } |
There was a problem hiding this comment.
In the global index calculation, blockIdx.x * blockDim.x is performed using 32-bit unsigned integer arithmetic. If the total number of elements Nelem is large (greater than blockIdx.x to cytnx_uint64 before multiplication prevents this overflow.
Also, we should update the call to use the renamed cu_abs_value helper.
template <typename TOut, typename TIn>
__global__ void cuAbs_dispatch_kernel(TOut *out, const TIn *in, const cytnx_uint64 Nelem) {
const cytnx_uint64 idx = static_cast<cytnx_uint64>(blockIdx.x) * blockDim.x + threadIdx.x;
if (idx < Nelem) out[idx] = static_cast<TOut>(cu_abs_value(in[idx]));
}
| template <typename TIn> | ||
| void cuAbs_dispatch_typed(boost::intrusive_ptr<Storage_base> &out, | ||
| const boost::intrusive_ptr<Storage_base> &in, | ||
| const cytnx_uint64 &Nelem) { |
There was a problem hiding this comment.
The parameter Nelem is passed as a const cytnx_uint64 &. According to the repository style guide, small/scalar types like cytnx_uint64 should be passed by value rather than by const& for better performance (riding in registers).
template <typename TIn>
void cuAbs_dispatch_typed(boost::intrusive_ptr<Storage_base> &out,
const boost::intrusive_ptr<Storage_base> &in,
cytnx_uint64 Nelem) {
References
- Pass small/scalar types by value, not by const&. For built-ins, enums, cytnx_uint64, Scalar, and even complex, by-value is never slower in Cytnx and is often faster. (link)
| cytnx_uint32 NBlocks = Nelem / 512; | ||
| if (Nelem % 512) NBlocks += 1; |
There was a problem hiding this comment.
The calculation of NBlocks can be simplified and made more idiomatic using integer division rounding up: (Nelem + 511) / 512. This avoids the modulo operation and conditional branch, and we can mark NBlocks as const.
const cytnx_uint32 NBlocks = (Nelem + 511) / 512;
| void cuAbs_dispatch(boost::intrusive_ptr<Storage_base> &out, | ||
| const boost::intrusive_ptr<Storage_base> &in, const cytnx_uint64 &Nelem) { |
There was a problem hiding this comment.
The parameter Nelem is passed as a const cytnx_uint64 &. According to the repository style guide, small/scalar types like cytnx_uint64 should be passed by value rather than by const&.
void cuAbs_dispatch(boost::intrusive_ptr<Storage_base> &out,
const boost::intrusive_ptr<Storage_base> &in, cytnx_uint64 Nelem) {
References
- Pass small/scalar types by value, not by const&. For built-ins, enums, cytnx_uint64, Scalar, and even complex, by-value is never slower in Cytnx and is often faster. (link)
| void cuAbs_dispatch(boost::intrusive_ptr<Storage_base> &out, | ||
| const boost::intrusive_ptr<Storage_base> &in, const cytnx_uint64 &Nelem); |
There was a problem hiding this comment.
The parameter Nelem is passed as a const cytnx_uint64 &. According to the repository style guide, small/scalar types like cytnx_uint64 should be passed by value rather than by const&.
| void cuAbs_dispatch(boost::intrusive_ptr<Storage_base> &out, | |
| const boost::intrusive_ptr<Storage_base> &in, const cytnx_uint64 &Nelem); | |
| void cuAbs_dispatch(boost::intrusive_ptr<Storage_base> &out, | |
| const boost::intrusive_ptr<Storage_base> &in, cytnx_uint64 Nelem); |
References
- Pass small/scalar types by value, not by const&. For built-ins, enums, cytnx_uint64, Scalar, and even complex, by-value is never slower in Cytnx and is often faster. (link)
![chatgpt-codex-connector[bot]](https://avatars.githubusercontent.com/in/1144995?s=60&v=4){kind=small}
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 9d0ca02d62
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
| std::is_same_v<TIn, cytnx_cuda_complex64>) { | ||
| return cuda::std::abs(x); // magnitude: complex<T> -> T | ||
| } else if constexpr (std::is_signed_v<TIn>) { | ||
| return x < TIn(0) ? static_cast<TIn>(-x) : x; // fabs / abs for signed real & integer |
There was a problem hiding this comment.
When TIn is cytnx_float or cytnx_double and a GPU tensor contains -0.0 (or a negative NaN), this signed-type branch does not implement fabs semantics: -0.0 < 0 is false, so Abs returns a value with the negative sign bit. The old CUDA kernels used fabs/fabsf and the CPU path still uses std::fabs, so CUDA Abs now differs from CPU/standard abs for those floating values; handle floating types with fabs/cuda::std::abs before the integer branch.
Useful? React with 👍 / 👎.
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## master #1094 +/- ##
==========================================
+ Coverage 72.71% 73.25% +0.53%
==========================================
Files 226 225 -1
Lines 28151 27802 -349
Branches 71 71
==========================================
- Hits 20471 20365 -106
+ Misses 7659 7416 -243
Partials 21 21
Flags with carried forward coverage won't be shown. Click here to find out more.
... and 5 files with indirect coverage changes Continue to review full report in Codecov by Harness.
🚀 New features to boost your workflow:
|
yingjerkao
changed the title
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 88a4d24899
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
| if constexpr (std::is_same_v<T, cytnx_cuda_complex128>) { | ||
| return cuda::std::pow(x, p); | ||
| } else if constexpr (std::is_same_v<T, cytnx_cuda_complex64>) { | ||
| return cuda::std::pow(x, static_cast<float>(p)); |
There was a problem hiding this comment.
For GPU ComplexFloat tensors with a double exponent that is not exactly representable as float (for example p = 1.0 + 2^-25 on large-magnitude values), this cast drops exponent precision before evaluating the complex power. The CPU Pow_internal_cf path calls std::pow(complex<float>, double) and the previous CUDA kernel used the double p in its magnitude/angle formula, so CUDA Pow/Pow_ can now diverge from CPU for those inputs; keep p as double until the final ComplexFloat result cast.
Useful? React with 👍 / 👎.
Problem: GPU Abs (out-of-place Abs.cpp and in-place Abs_.cpp) dispatched through the legacy `lii.cuAbs_ii[dtype]` lookup table, whose kernels use the old CUDA C complex types (cuDoubleComplex / cuCabs). The CPU path already uses typed std::visit dispatch (AbsInternalImpl). Part of #1003 step 11 (unary ops). Fix: add a typed `cuAbs_dispatch(out, in, Nelem)` that switches on the input dtype and runs a templated kernel, encoding the operation-specific output-dtype rule Abs(complex) -> real (ComplexDouble -> Double, ComplexFloat -> Float; every other dtype -> itself) and using cuda::std::abs on the cuda::std::complex GPU scalar types. Both Abs.cpp and Abs_.cpp (which already allocate `out` with the Abs output dtype) now call it instead of the lookup table. The legacy cuAbs_ii table + cuAbs_internal.cu are now unreferenced by any live path; removing them (and their remaining cuDoubleComplex usage) is a follow-up, mirroring how the binary-arithmetic table removal (#1003 step 12) followed the dispatch migration. Testing: on RTX 4070 Ti / CUDA 13, all 32 GPU Abs tests pass (gpu_tensor_abs_* and gpu_tensor_abs_inplace_*, every dtype x shape, comparing CUDA Abs against CPU Abs for dtype + value). clang-format-14 clean. Advances #1003 step 11 (Abs; Exp/Pow/Conj to follow the same pattern). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Completes the GPU unary migration started for Abs. Exp, Pow and Conj still dispatched through the legacy lii.cuExp_ii / cuPow_ii / cuConj_inplace_ii lookup tables and their CUDA C complex kernels. Add typed cuExp_dispatch / cuPow_dispatch / cuConj_inplace_dispatch: - Exp and Pow are dtype-preserving (the front end pre-casts integer/bool inputs to Double and passes the buffer as both in and out), so they only ever see ComplexDouble/ComplexFloat/Double/Float; they use cuda::std::exp / cuda::std::pow on the cuda::std::complex GPU scalar types (Pow carries the double exponent). - Conj is in place and complex-only (real Conj is a no-op), using cuda::std::conj. All ten call sites (Exp/Expf/Exp_/Expf_, Pow x2/Pow_ x2, Conj/Conj_) now call the typed dispatch instead of the lookup tables. Adds tests/gpu/linalg_test/UnaryExpPowConj_test.cpp: these ops had no GPU tests. It checks GPU Exp/Pow/Conj (out-of-place and in-place) against the CPU result for dtype and value; Exp/Pow use small deterministic inputs to avoid overflow, Conj uses genuine random complex data. The legacy cuExp_ii / cuPow_ii / cuConj_inplace_ii tables + their *_internal.cu are now unreferenced by any live path; removing them is a follow-up (as with Abs). Testing: on RTX 4070 Ti / CUDA 13, the 3 new GpuUnary tests pass (and the 32 existing Abs tests still pass). clang-format-14 clean. Advances #1003 step 11 (unary ops Abs/Exp/Pow/Conj now on typed dispatch). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
… 11)
With Abs/Exp/Pow/Conj migrated to typed dispatch (previous commits), the legacy
cuAbs_ii / cuExp_ii / cuPow_ii / cuConj_inplace_ii lookup tables are unreachable.
Remove them and the CUDA C complex kernels behind them:
- drop the four cu*_ii members + their population from
linalg_internal_interface.{hpp,cpp};
- delete cuAbs_internal.cu / cuExp_internal.cu / cuPow_internal.cu (whose kernels
used cuDoubleComplex / cuCabs) and slim their headers to just the *_dispatch
declaration;
- drop the three removed .cu from the GPU CMakeLists.
cuConj_inplace_internal.cu is kept: cuSvd_internal.cu / cuGeSvd_internal.cu still
call cuConj_inplace_internal_cd/cf directly (only its dead table is removed). The
per-op typedefs (Expfunc_oii/Powfunc_oii/Conjinplacefunc_oii) are also kept -- the
CPU Exp_ii/Pow_ii/Conj_inplace_ii tables still use them.
Testing: on RTX 4070 Ti / CUDA 13, gpu_test_main builds and links clean, and the
32 Abs + 3 Exp/Pow/Conj GPU tests still pass. clang-format-14 clean.
Completes the #1003 step-11 unary migration (typed dispatch + dead-table removal).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
ff6fd31 to
9fb7230
Compare
|
Rebased onto current master (
Re-verified on RTX 4070 Ti (CUDA 13, sm_89): |
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 9fb7230255
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
| for (auto dt : kFloatish) { | ||
| Tensor t = SmallInput(dt); | ||
| SCOPED_TRACE("Exp dtype " + std::to_string(dt)); | ||
| EXPECT_TRUE(SameAsCpu("Exp", la::Exp(t), la::Exp(t.to(cytnx::Device.cpu)))); |
There was a problem hiding this comment.
In the GPU test context I inspected, this new test uses the CPU Cytnx Exp implementation as the oracle, and the same CPU-path comparison pattern is repeated for Pow and Conj. The repo guidance in CLAUDE.md says arithmetic tests must check independent expected values rather than comparing one Cytnx implementation path against another, because a shared semantic bug can pass on both CPU and GPU; please use literal/std-library expected values for these unary cases instead.
Useful? React with 👍 / 👎.
…#1003 step 11) Review follow-up (#1094): - cuAbs: route every non-unsigned type through cuda::std::abs (IvanaGyro). This fixes floating Abs(-0.0)/Abs(negative NaN), which the old `x < 0 ? -x : x` returned with the sign bit set (-0.0 < 0 is false); rename CuAbsValue -> cu_abs_value (snake_case). - All four unary kernels: compute the global index in 64 bits (static_cast<cytnx_uint64>(blockIdx.x) * blockDim.x + ...) to avoid a 32-bit overflow past 2^32 elements; simplify NBlocks to (Nelem + 511) / 512; pass Nelem (and Pow's exponent p) by value, not const&. - cuPow ComplexFloat: keep the double exponent (compute in complex<double>, then narrow) instead of casting p to float first, matching the CPU Pow_internal_cf path for exponents not representable as float. - UnaryExpPowConj_test: add GpuUnary.matches_std_library (Exp/Pow/Conj checked against std::exp/std::pow/std::conj, independent of the Cytnx CPU path per CLAUDE.md) and GpuUnary.abs_negative_zero. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
|
Thanks — all addressed in the latest commit (verified: 37/37 unary + Abs GPU tests pass on an RTX 4070 Ti):
|
|
Review status (bot threads): the review-feedback commit (cd40c06) already addresses almost all of these against the current head:
Remaining: the GPU unary tests compare against the CPU path rather than independent literals (@chatgpt-codex-connector). GPU-vs-CPU is a reasonable cross-check (distinct kernels), but adding a couple of hand-computed anchors (e.g. |
|
To use Codex here, create an environment for this repo. |
…#1003 review) Addresses @chatgpt-codex-connector's review of #1094: GpuUnary.matches_std_library already checked Exp/Pow/Conj against independent std-library values for Double and ComplexDouble, but the narrower Float/ComplexFloat dtypes were only cross-checked against the Cytnx CPU path. Extend matches_std_library to anchor Float (Exp, Pow) and ComplexFloat (Exp, Pow, Conj) against std::exp/std::pow/std::conj directly, so a shared CPU/GPU semantic bug in the float kernels cannot hide behind a GPU-vs-CPU compare. Verified: gpu_test_main (CUDA 13, RTX 4070 Ti) GpuUnary.* all pass. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
|
Added the independent anchors for the narrower float dtypes in d4280ac: |
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: d4280ac54b
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
| # cudaMemcpyTruncation.hpp | ||
|
|
||
| cuAbs_internal.cu | ||
| cuAbs_dispatch.cu |
There was a problem hiding this comment.
When the CUDA library is built through developer_tools/Makefile with GPU_Enable=1, this source rename is not mirrored there: rg shows the Makefile still adds cuAbs_internal.o, cuPow_internal.o, and cuExp_internal.o with rules pointing at the deleted .cu files, and it never compiles the new cu*dispatch.cu objects such as cuConj_inplace_dispatch.cu. That build path will either stop on missing sources or later link undefined references from the updated src/linalg/{Abs,Exp,Pow,Conj}*.cpp callers; please update the Makefile object list/rules along with the CMake source list.
Useful? React with 👍 / 👎.
|
is |
|
Codex review notes: The numerical fixes here look sound, but I think this still needs changes before merge.
template <typename TOut, typename TIn, typename Op>
__global__ void unary_kernel(TOut* out, const TIn* in, std::size_t size, Op op) {
const std::size_t i =
static_cast<std::size_t>(blockIdx.x) * blockDim.x + threadIdx.x;
if (i < size) out[i] = op(in[i]);
}with
T* ptr = reinterpret_cast<T*>(storage->data());
There are also minor new-style violations such as The red Finally, this branch is now 60 commits behind current |
3 participants
Problem
#1003 step 11 (unary ops). The GPU unary elementwise ops —
Abs,Exp,Pow,Conj— dispatched through legacy per-dtype lookup tables (lii.cuAbs_ii/cuExp_ii/cuPow_ii/cuConj_inplace_ii) whose kernels use the old CUDA C complex types (cuDoubleComplex,cuCabs, …), while the CPU paths already use typedstd::visitdispatch.Fix
Typed dispatch (
cuda::std::complexGPU scalar types), replacing the lookup tables at every call site:cuAbs_dispatch— operation-specific output ruleAbs(complex) → real(ComplexDouble→Double,ComplexFloat→Float; else identity);cuda::std::abs.cuExp_dispatch/cuPow_dispatch— dtype-preserving; the front end pre-casts integer/bool inputs toDoubleand passes the buffer as bothinandout, so onlyComplexDouble/ComplexFloat/Double/Floatare seen;cuda::std::exp/cuda::std::pow(Pow carries thedoubleexponent).cuConj_inplace_dispatch— in place, complex-only (realConjis a no-op);cuda::std::conj.Call sites rewired:
Abs/Abs_,Exp/Expf/Exp_/Expf_,Pow×2/Pow_×2,Conj/Conj_.Dead-table removal (final commit): the now-unreachable
cuAbs_ii/cuExp_ii/cuPow_ii/cuConj_inplace_iitables are dropped fromlinalg_internal_interface, andcuAbs_internal.cu/cuExp_internal.cu/cuPow_internal.cu(thecuDoubleComplex/cuCabskernels) are deleted.cuConj_inplace_internal.cuis kept —cuSvd_internal.cu/cuGeSvd_internal.custill callcuConj_inplace_internal_cd/cfdirectly (only its dead table is removed).Testing
On RTX 4070 Ti / CUDA 13:
tests/gpu/linalg_test/UnaryExpPowConj_test.cpp(these ops had no GPU tests): GPUExp/Pow/Conjout-of-place + in-place vs CPU — 3 tests pass.gpu_test_mainbuilds/links clean and all 35 still pass.clang-format-14 clean. Completes #1003 step 11. Independent of the sibling #1003 branches (#1091–#1093).
🤖 Generated with Claude Code