Skip to content

Commit c08b0ba

Browse files
Enable block IO store on regular pointer by default (#4749)
Benchmark CI: https://github.com/intel/intel-xpu-backend-for-triton/actions/runs/16389219960 (no geomean regression) Inductor UT CI: https://github.com/intel/intel-xpu-backend-for-triton/actions/runs/16389222996 (passes) Flex Decoding CI: https://github.com/intel/intel-xpu-backend-for-triton/actions/runs/16390730800 (passes) Flex Attention CI: https://github.com/intel/intel-xpu-backend-for-triton/actions/runs/16417773290 (no regression) This PR changes the default behavior. Plan to remove the env var after merging a while and no unexpected regression. Signed-off-by: Whitney Tsang <[email protected]>
1 parent 9a77cdb commit c08b0ba

File tree

1 file changed

+4
-3
lines changed

1 file changed

+4
-3
lines changed

third_party/intel/lib/TritonIntelGPUToLLVM/LoadStoreOpToLLVM.cpp

Lines changed: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -2756,9 +2756,10 @@ struct StoreOpToBlockIOConversion
27562756
offsetBaseX = offsetX;
27572757
offsetBaseY = offsetY;
27582758
} else {
2759-
static const bool enableBlockStore = triton::tools::getBoolEnv(
2760-
"TRITON_INTEL_ENABLE_BLOCK_IO_STORE_ON_REGULAR_PTR");
2761-
if (!enableBlockStore)
2759+
std::optional<bool> enableBlockStore =
2760+
mlir::triton::tools::isEnvValueBool(mlir::triton::tools::getStrEnv(
2761+
"TRITON_INTEL_ENABLE_BLOCK_IO_STORE_ON_REGULAR_PTR"));
2762+
if (enableBlockStore.has_value() && !enableBlockStore.value())
27622763
return failure();
27632764
// Get the LLVM values for pointers
27642765
ptrElems = unpackLLElements(loc, llPtr, rewriter);

0 commit comments

Comments
 (0)