From 0edaa6e47dcc83e7f09576034cc74e3f85c235ae Mon Sep 17 00:00:00 2001 From: Johnsonms Date: Thu, 22 Jan 2026 20:54:12 -0800 Subject: [PATCH] Fix out-of-bounds TMA access in wgmma_tma_sm90 tutorial (#2945) --- examples/cute/tutorial/hopper/wgmma_tma_sm90.cu | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/examples/cute/tutorial/hopper/wgmma_tma_sm90.cu b/examples/cute/tutorial/hopper/wgmma_tma_sm90.cu index be4d80b8..36a53c46 100644 --- a/examples/cute/tutorial/hopper/wgmma_tma_sm90.cu +++ b/examples/cute/tutorial/hopper/wgmma_tma_sm90.cu @@ -236,7 +236,8 @@ gemm_device(ProblemShape shape_MNK, CtaTiler cta_tiler, ConsumerBarType::arrive(&consumer_mbar[read_pipe]); ++read_state; - if ((warp_idx == 0) && lane_predicate) + // Only issue new TMA copies if there are more tiles to fetch + if ((warp_idx == 0) && lane_predicate && (k_tile_count > 0)) { int pipe = write_state.index(); // Wait for Consumer to complete consumption