Skip to content

Commit eb5633d

Browse files
committed
test(cfo): use i32 offsets in dynamic block pointers
Keep descriptor storage wide while satisfying Triton make_block_ptr's 32-bit offset contract in both if branches. This fixes the Group A pytest input without changing product lowering or marker semantics.
1 parent a51e902 commit eb5633d

1 file changed

Lines changed: 5 additions & 2 deletions

File tree

third_party/ascend/unittest/pytest_ut/test_control_flow_pointer_boundary.py

Lines changed: 5 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -35,8 +35,11 @@ def block_ptr_if_descriptor_kernel(a, b, out, choose_ptr, descriptor_ptr, BLOCK:
3535
size_b = tl.load(descriptor_ptr + 1)
3636
stride_a = tl.load(descriptor_ptr + 2)
3737
stride_b = tl.load(descriptor_ptr + 3)
38-
offset_a = tl.load(descriptor_ptr + 4)
39-
offset_b = tl.load(descriptor_ptr + 5)
38+
# Block-pointer offsets are required to use Triton's 32-bit index type.
39+
# Keep the descriptor storage wide, but narrow the values at the frontend
40+
# boundary instead of changing the product's make_block_ptr contract.
41+
offset_a = tl.load(descriptor_ptr + 4).to(tl.int32)
42+
offset_b = tl.load(descriptor_ptr + 5).to(tl.int32)
4043
if choose != 0:
4144
pointer = tl.make_block_ptr(
4245
base=a + 3,

0 commit comments

Comments
 (0)