Repository navigation
Execute Widen arrays on CUDA - #11303
Open
connortsui20 wants to merge 1 commit into
Open
connortsui20 wants to merge 1 commit into
connortsui20 wants to merge 1 commit into
Conversation
connortsui20
commented
Nov 6, 2026
Comment on lines
47
to
64
| @@ -63,7 +64,7 @@ __device__ __forceinline__ Output decimal_cast_value(Input value) { | |||
| } | |||
Member
Author
There was a problem hiding this comment.
I am unfamiliar with the GPU code but from my LLM conversations, it seems like we do not know if the optimizer can remove intermediate conversions? i.e. i8->i64->i16?
connortsui20
marked this pull request as ready for review
November 6, 2026 17:06
Merging this PR will improve performance by 14.57%1 warning
Performance Changes
Tip Curious why performance improved? Comment Comparing 
|
Device-resident Widen arrays cannot use the CPU fallback. Restore their logical width on the device so narrow integer and decimal children can execute through CUDA. Progress towards #11284. Signed-off-by: Connor Tsui <connor.tsui20@gmail.com>
connortsui20
force-pushed
the
ct/cuda-widen
branch
from
November 6, 2026 19:06
1741725 to
37897fe
Compare
Contributor
|
Will have a look next week first thing! @connortsui20 |
This branch has not been deployed
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Progress towards #11284 (part 1).
Stacked on #11302, which this change doesn't depend on. The stack only exists so that CD5, which needs both, can sit on top.
CUDA has no kernel for
vortex.widen, and the CPU fallback only runs when every buffer is on the host, so executing a device-residentWidenfails with "CPU fallback with device-resident buffers is not supported". After CD5, every DecimalByteParts array with a narrow MSP reads from a file with an integerWidenas its MSP, so CD5 can't merge until GPU reads can execute it.Changes
Widenis not fused into the dynamic dispatch path. Dynamic dispatch zero-extends signed narrow BitPacked and FoR children, so fusing a signedWidenthere needs care to sign-extend instead. TheTODOinplan_builder.rstracks that. Until then,WidenExecutorexecutes the child on the device and widens it in a separate kernel.Decimal storage that is wider than the
Widen's logical type is an error instead of a narrowing copy on the device. Narrowing needs either an overflow check or trust in the stored precision, and that choice is the open question in part 2 of #11284.The decimal path reuses the kernels in
decimal_cast.cu, and the Arrow export now shareswiden_decimal_valueswith it, so there is one copy of the widening code. The tests compare every integer and decimal widening pair against the CPUWiden(with nulls, negative values, a FoR child, and the kernel's block boundaries). They only run in the GPU CI.