Remove env var TRITON_INTEL_ENABLE_BLOCK_IO_STORE_ON_REGULAR_PTR (#4851)

There was no regression discovered after enabling
`TRITON_INTEL_ENABLE_BLOCK_IO_STORE_ON_REGULAR_PTR` by default. Since we
always want to enable this behavior, the env var is not needed.

Signed-off-by: Whitney Tsang <[email protected]>

Author: whitneywhtsang

include/triton/Tools/Sys/GetEnv.hpp

@@ -48,7 +48,6 @@ inline const std::set<std::string> CACHE_INVALIDATING_ENV_VARS = {
     "TRITON_INTEL_AGGRESSIVE_DPAS_REUSE",
     "TRITON_INTEL_DO_NOT_SINK_INSTR_ACROSS_RGN",
     "TRITON_INTEL_ENABLE_BLOCK_IO_ALL_LAYOUTS",
-    "TRITON_INTEL_ENABLE_BLOCK_IO_STORE_ON_REGULAR_PTR",
     "TRITON_INTEL_ENABLE_FIRST_LOAD_TO_SLM",
     "TRITON_INTEL_ENABLE_INSTR_SCHED",
     "TRITON_INTEL_FAST_MATH",

python/test/unit/intel/test_block_store.py

@@ -9,7 +9,6 @@
 import triton
 from triton._internal_testing import is_xpu
 
-os.environ["TRITON_INTEL_ENABLE_BLOCK_IO_STORE_ON_REGULAR_PTR"] = "1"
 os.environ["TRITON_INTEL_ENABLE_BLOCK_IO_ALL_LAYOUTS"] = "1"
 
 

test/TritonIntelGPU/tensor-pointer-store-block-2d.mlir

@@ -1,5 +1,5 @@
-// RUN: env TRITON_INTEL_ENABLE_BLOCK_IO_STORE_ON_REGULAR_PTR=1 triton-opt %s -split-input-file --intel-allocate-shared-memory --convert-triton-intel-gpu-to-llvm | FileCheck %s --implicit-check-not=llvm.inline_asm
-// RUN: env TRITON_INTEL_ENABLE_BLOCK_IO_STORE_ON_REGULAR_PTR=1 TRITON_INTEL_ENABLE_BLOCK_IO_ALL_LAYOUTS=1 triton-opt %s -split-input-file --intel-allocate-shared-memory --convert-triton-intel-gpu-to-llvm | FileCheck %s --implicit-check-not=llvm.inline_asm  --check-prefixes=ALL-LAYOUT
+// RUN: triton-opt %s -split-input-file --intel-allocate-shared-memory --convert-triton-intel-gpu-to-llvm | FileCheck %s --implicit-check-not=llvm.inline_asm
+// RUN: env TRITON_INTEL_ENABLE_BLOCK_IO_ALL_LAYOUTS=1 triton-opt %s -split-input-file --intel-allocate-shared-memory --convert-triton-intel-gpu-to-llvm | FileCheck %s --implicit-check-not=llvm.inline_asm  --check-prefixes=ALL-LAYOUT
 
 #blocked = #ttg.blocked<{sizePerThread = [1, 4, 2], threadsPerWarp = [1, 1, 32], warpsPerCTA = [1, 8, 2], order = [2, 1, 0]}>
 #slice = #ttg.slice<{dim = 1, parent = #blocked}>

third_party/intel/lib/TritonIntelGPUToLLVM/LoadStoreOpToLLVM.cpp

@@ -2772,11 +2772,6 @@ struct StoreOpToBlockIOConversion
       offsetBaseX = offsetX;
       offsetBaseY = offsetY;
     } else {
-      std::optional<bool> enableBlockStore =
-          mlir::triton::tools::isEnvValueBool(mlir::triton::tools::getStrEnv(
-              "TRITON_INTEL_ENABLE_BLOCK_IO_STORE_ON_REGULAR_PTR"));
-      if (enableBlockStore.has_value() && !enableBlockStore.value())
-        return failure();
       // Get the LLVM values for pointers
       ptrElems = unpackLLElements(loc, llPtr, rewriter);
       assert(ptrElems.size() == numElems &&