Skip to content

Commit b6ef4ce

Browse files
authored
fix stack OP benchmark (flagos-ai#2716)
* fix stack OP benchmard * update * update
1 parent 57739c2 commit b6ef4ce

2 files changed

Lines changed: 10 additions & 5 deletions

File tree

benchmark/test_stack.py

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -33,7 +33,6 @@ def _input_fn(shape, dtype, device):
3333
yield [inp1, inp2, inp3], {"dim": -1},
3434

3535

36-
@pytest.mark.skip(reason="CUDA error - illegal memory access: issue #2675")
3736
@pytest.mark.stack
3837
def test_stack():
3938
bench = StackBenchmark(

src/flag_gems/ops/stack.py

Lines changed: 10 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -47,11 +47,17 @@ def stack_copy_func_kernel_4(
4747
dim_offset = dim_offset_d
4848
total_elements = total_elements_d
4949

50-
block_start = pid_x * BLOCK_X
51-
offsets = tl.arange(0, BLOCK_X)
52-
mask = block_start + offsets < total_elements
53-
50+
block_start = pid_x.to(tl.int64) * BLOCK_X
51+
offsets = tl.arange(0, BLOCK_X).to(tl.int64)
5452
idx = block_start + offsets
53+
scalar_zero = offsets * 0
54+
55+
dim_size_out = dim_size_out + scalar_zero
56+
dim_prod_post = dim_prod_post + scalar_zero
57+
dim_offset = dim_offset + scalar_zero
58+
total_elements = total_elements + scalar_zero
59+
60+
mask = idx < total_elements
5561

5662
pre_idx = idx // dim_prod_post
5763
post_idx = idx % dim_prod_post

0 commit comments

Comments
 (0)