diff --git a/src/op/parallel.cc b/src/op/parallel.cc index 8a4790fd4..96d0dfda7 100644 --- a/src/op/parallel.cc +++ b/src/op/parallel.cc @@ -422,6 +422,7 @@ LayoutMap ParallelOpNode::InferLayout(const LayoutInferArgs &T, // Step 1: try to infer loop's partition from a source fragment Buffer source_buffer, read_source_buffer; + bool source_buffer_is_write = false; Buffer replicated_write_buffer; // Backup: fully replicated write buffer for (const auto &buffer : access_order_) { @@ -437,6 +438,7 @@ LayoutMap ParallelOpNode::InferLayout(const LayoutInferArgs &T, if (access.is_write) { source_buffer = buffer; + source_buffer_is_write = true; } else { // Keep the buffer with largest number of indices // (which means the inference based on that buffer is more accurate) @@ -455,6 +457,7 @@ LayoutMap ParallelOpNode::InferLayout(const LayoutInferArgs &T, if (frag.has_value() && is_one(frag.value()->ReplicateExtent()) && !source_buffer.defined()) { source_buffer = buffer; + source_buffer_is_write = false; } } } @@ -476,7 +479,7 @@ LayoutMap ParallelOpNode::InferLayout(const LayoutInferArgs &T, predicate_ = annotated_predicate_.value(); } } else if (!loop_layout_.defined() && source_buffer.defined() && - allow_layout_propgate) { + (allow_layout_propgate || source_buffer_is_write)) { loop_layout_ = ComputeLoopLayoutFromBuffer(source_buffer, T); } else if (!loop_layout_.defined() && level == InferLevel::kFree) { // For free layout inference diff --git a/testing/python/layout/test_tilelang_layout_inference.py b/testing/python/layout/test_tilelang_layout_inference.py index 13b58cf23..ad8f03ce0 100644 --- a/testing/python/layout/test_tilelang_layout_inference.py +++ b/testing/python/layout/test_tilelang_layout_inference.py @@ -60,5 +60,28 @@ def test_reject_fragment_write_from_non_owner_threads(): _invalid_fragment_write_owner_layout() +@tilelang.jit +def _scalar_reduce_accumulator_owner_layout(): + @T.prim_func + def main(out: T.Tensor((1,), T.float32)): + with T.Kernel(1, threads=128): + acc = T.alloc_fragment((128,), T.float32) + total = T.alloc_fragment((1,), T.float32) + total[0] = T.float32(0.0) + for k in T.Parallel(128): + acc[k] = T.float32(1.0) + T.reduce_sum(acc, total, clear=False) + if T.get_thread_binding() == 0: + out[0] = total[0] + + return main + + +def test_scalar_reduce_accumulator_owner_layout(): + kernel = _scalar_reduce_accumulator_owner_layout() + source = kernel.get_kernel_source() + assert "AllReduce" in source + + if __name__ == "__main__": tilelang.testing.main()