refactor(linalg): pass cuCpr non-contiguous layout by value (#1003) by yingjerkao · Pull Request #1098 · Cytnx-dev/Cytnx

yingjerkao · 2026-07-20T02:44:22Z

Stacked follow-up completing #1003 step 4 for the comparison path.

Depends on #1095 (shared cuNonContigLayout.cuh) and #1097 (which makes the GPU Cpr non-contiguous path live and provides its test coverage). The base of this PR is #1095's branch; the Cpr.cpp / Cpr_test.cpp changes shown in the diff are #1097. Once #1095 and #1097 merge to master, this reduces to just the cuCpr_dispatch.cu change.

What

#1095 gave the arithmetic non-contiguous kernels a by-value layout (no per-call cuMalloc_gpu/memcpy/free) but deliberately skipped cuCpr because its non-contiguous path was then unreachable. #1097 made that path live (GPU Cpr now accepts non-contiguous operands). This PR applies the same by-value treatment to cuCpr_dispatch's non-contiguous kernel: it reuses the shared GpuNonContigLayout / MakeGpuNonContigLayout / ComputeGpuNonContigIndices instead of allocating five device arrays per call (net −46 lines in cuCpr_dispatch.cu).

With this, all GPU elementwise non-contiguous kernels (Add/Sub/Mul/Div and Cpr) are free of per-call device allocation.

Testing

The CprNonContig tests (added in #1097) now exercise the by-value cuCpr kernel: permuted-vs-contiguous in both operand orders (each inverse mapper), GPU-vs-CPU across all dtypes, and a hand-computed identity-pattern literal. On an RTX 4070 Ti (CUDA 13, sm_89): gpu_test_main builds clean, CprNonContig passes, and the full Cpr GPU suite (82 tests) is green. clang-format-14 clean.

Note: GPU results are from the local RTX 4070 Ti — there is no GPU CI runner.

🤖 Generated with Claude Code

…, #988) Problem: the GPU out-of-place Cpr front end rejected a non-contiguous operand with "two tensors must be contiguous" -- the last GPU elementwise op to do so (Add/Sub already accepted it; Mul/Div were enabled in #1096) -- even though the CPU path and cuCpr_dispatch's non-contiguous Bool-output kernel handle it. Fix: Cpr.cpp now passes the layout (shape + inverse mappers) to cuCpr_dispatch in the non-contiguous branch, mirroring Add/Sub/Mul/Div, instead of erroring. cuCpr already had a working non-contiguous kernel; it was simply unreachable. Behavior change: GPU Cpr now accepts a non-contiguous operand and returns the correct Bool result (matching CPU) rather than throwing. Testing: new CprNonContig tests in Cpr_test.cpp -- permuted-vs-contiguous in both operand orders (each inverse mapper exercised), GPU-vs-CPU across all dtypes AND a hand-computed identity-pattern literal (equal iff i==j, a mix a wrong gather would break). Both fail on the pre-fix code (which threw). On an RTX 4070 Ti (CUDA 13, sm_89): gpu_test_main builds clean, the new tests pass, and the full Cpr GPU suite is green. clang-format-14 clean. With #1096 this closes the GPU non-contiguous gap (#988) for all five elementwise ops (Add/Sub/Mul/Div/Cpr). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

…yout-byvalue

Extends #1095's step-4 treatment to cuCpr, now that #1097 made the GPU Cpr non-contiguous path live. cuCpr_dispatch's non-contiguous kernel no longer cuMalloc_gpu/memcpy/free's five device layout arrays per call; it reuses the shared cuNonContigLayout.cuh (GpuNonContigLayout passed by value + MakeGpuNonContigLayout + ComputeGpuNonContigIndices), matching cuArithmeticDispatch. All GPU elementwise non-contiguous kernels are now free of per-call device allocation. Stacked on #1095 (shared header) and #1097 (live Cpr non-contiguous path, which provides the test coverage). Testing: the CprNonContig tests (from #1097) now exercise the by-value cuCpr kernel -- GPU-vs-CPU across all dtypes + hand-computed identity-pattern literal. On an RTX 4070 Ti (CUDA 13, sm_89): gpu_test_main builds clean, CprNonContig passes, and the full Cpr GPU suite is green. clang-format-14 clean. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

gemini-code-assist

Code Review

This pull request enables non-contiguous tensor comparison on the GPU by passing the layout to the non-contiguous CUDA kernel instead of throwing an error. The reviewer identified a critical bug where the storage size of the left tensor is used as the element count in cuCpr_dispatch, which can cause out-of-bounds writes for sliced non-contiguous tensors; using the output tensor's size is recommended instead. Additionally, the reviewer suggested adding test cases for sliced tensor views to ensure comprehensive coverage.

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 · 2026-07-20T02:45:37Z

+          checkCudaErrors(cudaSetDevice(Rt.device()));
+          cytnx::linalg_internal::cuCpr_dispatch(
+            out._impl->storage()._impl, left._impl->storage()._impl, right._impl->storage()._impl,
+            left._impl->storage()._impl->size(), left._impl->shape(), left._impl->invmapper(),


Using left._impl->storage()._impl->size() as the number of elements (len) to process in cuCpr_dispatch is a critical bug. For non-contiguous tensors (such as sliced views), the underlying storage size can be significantly larger than the logical size of the tensor. Passing the storage size of left will cause the GPU kernel to launch too many threads and perform out-of-bounds writes on the out tensor (which is allocated with the correct logical size of Lt.shape()). This leads to memory corruption or illegal memory access crashes on CUDA. It should be changed to out._impl->storage()._impl->size() (or left.shape().size()), which represents the actual logical size of the tensor being compared.

Suggested change

left._impl->storage()._impl->size(), left._impl->shape(), left._impl->invmapper(),

out._impl->storage()._impl->size(), left._impl->shape(), left._impl->invmapper(),

gemini-code-assist · 2026-07-20T02:45:37Z

+
+          Tensor permuted = arange(0, 9, 1, dtype).reshape({3, 3}).permute({1, 0}).to(Device.cuda);
+          Tensor contig = arange(0, 9, 1, dtype).reshape({3, 3}).to(Device.cuda);
+          ASSERT_FALSE(permuted.is_contiguous());


The current non-contiguous test cases only use permuted tensors where the storage size is equal to the logical size (9 elements). To prevent regressions and ensure that sliced non-contiguous tensors (where storage size is larger than logical size) are handled correctly without triggering out-of-bounds writes, consider adding a test case that uses a sliced tensor view.

in the current design the storage size and logical size are always the same.

chatgpt-codex-connector

💡 Codex Review

Here are some automated review suggestions for this pull request.

Reviewed commit: ce0d11f60e

ℹ️ 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".

chatgpt-codex-connector · 2026-07-20T02:49:04Z

            cuCpr_dispatch_tn_kernel_nonconti<TO>
-              <<<NBlocks, 512, 512 * shape.size() * sizeof(cytnx_uint64)>>>(
-                _out, _Lin, len, _Rin, m_accu_shape, m_old_accu_shapeL, m_old_accu_shapeR,
-                m_invmapper_L, m_invmapper_R, shape.size());
-
-            checkCudaErrors(cudaFree(m_accu_shape));
-            checkCudaErrors(cudaFree(m_old_accu_shapeL));
-            checkCudaErrors(cudaFree(m_old_accu_shapeR));
-            checkCudaErrors(cudaFree(m_invmapper_L));
-            checkCudaErrors(cudaFree(m_invmapper_R));
+              <<<NBlocks, 512, 512 * shape.size() * sizeof(cytnx_uint64)>>>(_out, _Lin, len, _Rin,


Guard high-rank non-contiguous GPU Cpr launches

When a non-contiguous GPU comparison has rank 13-32 and both inputs have more than one element, this launch requests 512 * shape.size() * sizeof(cytnx_uint64) bytes of dynamic shared memory. cuNonContigLayout.cuh notes that the default 48 KiB block budget is exhausted around rank 12, but MakeGpuNonContigLayout only rejects ranks above 32, so this newly reachable Cpr path can fail with a CUDA launch error instead of producing a result. Please guard against the actual shared-memory limit or fall back to a contiguous path for those ranks.

Useful? React with 👍 / 👎.

…Cpr (#1003) Review follow-up (#1097): Gemini flagged (critical) that passing left._impl->storage()._impl->size() as the kernel length risks an out-of-bounds write for a strided view whose storage exceeds its logical size. Use out._impl->storage()._impl->size() -- the logical element count -- matching the CPU non-contiguous path (Add.cpp:120) and the contiguous GPU branch. Verified the specific counterexample is not currently reachable: Cytnx compacts slices (a sliced tensor comes back contiguous), so a non-contiguous tensor's storage always equals its logical size and left/out sizes coincide today. The fix is for correctness/consistency and future-proofing rather than a live crash. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

…or/1003-cucpr-layout-byvalue

…yout-byvalue

yingjerkao · 2026-07-20T03:47:03Z

Addressed, and re-synced with the parent PRs:

Length fix (Gemini, critical): the cuCpr_dispatch length now comes from Cpr.cpp as out._impl->storage()._impl->size() (merged from fix(linalg): support non-contiguous tensor(==)tensor for GPU Cpr (#1003, #988) #1097). Same reasoning as the fix(linalg): support non-contiguous tensor(==)tensor for GPU Cpr (#1003, #988) #1097 thread — the strided-view counterexample isn't currently reachable in Cytnx (slices compact to contiguous), but the fix is correct and matches the CPU path.
Rank-13+ guard (Codex): carried in from refactor(linalg): pass GPU non-contiguous elementwise layout by value (#1003) #1095 — kGpuNonContigMaxRank is now 12 (the true 512 × rank × 8 B ≤ 48 KB limit), so MakeGpuNonContigLayout raises a clear error for a rank the by-value cuCpr kernel can't launch, rather than aborting on the CUDA launch.

Re-verified on an RTX 4070 Ti: the by-value cuCpr non-contiguous kernel passes CprNonContig and the full Cpr GPU suite.

…alue branch

Follows the #1095 review update (strides view, merged in): cuCpr_dispatch's non-contiguous comparison kernel now uses the renamed shared helpers (make_gpu_non_contig_layout / compute_gpu_non_contig_indices) and launches with no dynamic shared memory (the per-thread tmpv multi-index scratch is gone, the recompose folds into the decompose loop). Behaviour unchanged. Testing: CprNonContig (GpuNoncontiguousMatchesCpu + GpuNoncontiguousLiteral), NonContigElementwise, InplacePromote all pass (8/8) on an RTX 4070 Ti SUPER (CUDA 13); GPU build clean. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

ianmccul

Fine, I hope this is a step towards consolidating all of the unary and binary operations through one kernel templated on the operator.

Once that is one there is only a single (or a small number) of functions that need to be replaced with stride arrays.

By the time all this is done, it would have been much easier to rewrite from scratch!

IvanaGyro · 2026-07-20T09:14:57Z

+          // Non-contiguous tensor(==)tensor on the GPU: cuCpr_dispatch's non-contiguous kernel
+          // applies the layout mappers (as Add/Sub/Mul/Div already do), so pass the layout
+          // instead of forcing the caller to contiguous-ize first (#1003, #988).
+          checkCudaErrors(cudaSetDevice(Rt.device()));
+          cytnx::linalg_internal::cuCpr_dispatch(
+            out._impl->storage()._impl, left._impl->storage()._impl, right._impl->storage()._impl,
+            out._impl->storage()._impl->size(), left._impl->shape(), left._impl->invmapper(),
+            right._impl->invmapper());


"Storage + stride + shape" is conceptually a tensor. We should consider to pass tensor directly instead in the future.

yingjerkao and others added 3 commits July 20, 2026 10:32

Merge branch 'fix/1003-gpu-cpr-noncontig' into refactor/1003-cucpr-la…

65a6971

…yout-byvalue

gemini-code-assist Bot reviewed Jul 20, 2026

View reviewed changes

chatgpt-codex-connector Bot reviewed Jul 20, 2026

View reviewed changes

yingjerkao and others added 3 commits July 20, 2026 11:13

Merge branch 'refactor/1003-gpu-noncontig-layout-byvalue' into refact…

2d9d276

…or/1003-cucpr-layout-byvalue

Merge branch 'fix/1003-gpu-cpr-noncontig' into refactor/1003-cucpr-la…

6d1335c

…yout-byvalue

yingjerkao mentioned this pull request Jul 20, 2026

fix(linalg): support non-contiguous tensor(==)tensor for GPU Cpr (#1003, #988) #1097

Merged

yingjerkao and others added 2 commits July 20, 2026 16:42

Merge updated #1095 (strides-view layout, snake_case) into cuCpr by-v…

0aeeaed

…alue branch

yingjerkao requested review from IvanaGyro and ianmccul July 20, 2026 08:58

Base automatically changed from refactor/1003-gpu-noncontig-layout-byvalue to master July 20, 2026 09:02

ianmccul approved these changes Jul 20, 2026

View reviewed changes

IvanaGyro approved these changes Jul 20, 2026

View reviewed changes

yingjerkao merged commit 014a386 into master Jul 20, 2026
8 checks passed

yingjerkao deleted the refactor/1003-cucpr-layout-byvalue branch July 20, 2026 09:48

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

refactor(linalg): pass cuCpr non-contiguous layout by value (#1003)#1098

refactor(linalg): pass cuCpr non-contiguous layout by value (#1003)#1098
yingjerkao merged 8 commits into
masterfrom
refactor/1003-cucpr-layout-byvalue

yingjerkao commented Jul 20, 2026

Uh oh!

gemini-code-assist Bot left a comment

Uh oh!

gemini-code-assist Bot Jul 20, 2026

Uh oh!

gemini-code-assist Bot Jul 20, 2026

Uh oh!

ianmccul Jul 20, 2026

Uh oh!

chatgpt-codex-connector Bot left a comment

Uh oh!

chatgpt-codex-connector Bot Jul 20, 2026

Uh oh!

yingjerkao commented Jul 20, 2026

Uh oh!

ianmccul left a comment

Uh oh!

IvanaGyro Jul 20, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

	left._impl->storage()._impl->size(), left._impl->shape(), left._impl->invmapper(),
	out._impl->storage()._impl->size(), left._impl->shape(), left._impl->invmapper(),

Uh oh!

Conversation

yingjerkao commented Jul 20, 2026

What

Testing

Uh oh!

gemini-code-assist Bot left a comment

Choose a reason for hiding this comment

Code Review

Uh oh!

gemini-code-assist Bot Jul 20, 2026

Choose a reason for hiding this comment

Uh oh!

gemini-code-assist Bot Jul 20, 2026

Choose a reason for hiding this comment

Uh oh!

ianmccul Jul 20, 2026

Choose a reason for hiding this comment

Uh oh!

chatgpt-codex-connector Bot left a comment

Choose a reason for hiding this comment

💡 Codex Review

Uh oh!

chatgpt-codex-connector Bot Jul 20, 2026

Choose a reason for hiding this comment

Uh oh!

yingjerkao commented Jul 20, 2026

Uh oh!

ianmccul left a comment

Choose a reason for hiding this comment

Uh oh!

IvanaGyro Jul 20, 2026

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants