Skip to content

[Fix][CUDA] Allocate complete words for shared sub-byte buffers - #20310

Open
LittlehamsterXu wants to merge 1 commit into
apache:mainfrom
LittlehamsterXu:fix/tvm-20273-ceil-shared-subbyte
Open

LittlehamsterXu wants to merge 1 commit into
apache:mainfrom
LittlehamsterXu:fix/tvm-20273-ceil-shared-subbyte

Conversation

@LittlehamsterXu

@LittlehamsterXu LittlehamsterXu commented Sep 10, 2026

Copy link
Copy Markdown

Description

CUDA codegen packs int1, int4, and uint4 values in shared memory using 32-bit words. The static shared-memory allocation size was previously computed with floor division.

For shapes that are smaller than one packing word, this could generate a zero-sized __shared__ array. For non-divisible shapes, it could also under-allocate the final partial word.

This patch uses ceil division so that every partial final word is allocated. It also adds regression coverage for int1, int4, and uint4 shapes.

Fixes #20273

Testing

  • pytest -q tests/python/tirx/codegen/test_codegen_cuda.py -k subbyte_shared_alloc_uses_ceil_div
  • pytest -q tests/python/tirx/codegen/test_codegen_cuda.py
  • ruff check tests/python/tirx/codegen/test_codegen_cuda.py
  • git diff --check

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.

[Bug][CUDA] Static shared sub-byte allocation uses floor-sized storage

1 participant