Skip to content

Commit e8996c5

Browse files
committed
fix review comments
1 parent c3d1743 commit e8996c5

File tree

2 files changed

+11
-8
lines changed

2 files changed

+11
-8
lines changed

third_party/intel/backend/compiler.py

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -233,8 +233,9 @@ def make_ttgir(mod, metadata, opt, properties):
233233
pm = ir.pass_manager(mod.context)
234234
pm.enable_debug()
235235

236-
if (properties["has_subgroup_2d_block_io"] and properties["has_subgroup_matrix_multiply_accumulate"]
237-
and (os.getenv("TRITON_INTEL_ADVANCED_PATH", "0") == "1" or opt.advanced_path)):
236+
if (os.getenv("TRITON_INTEL_ADVANCED_PATH", "0") == "1" or opt.advanced_path):
237+
assert properties["has_subgroup_2d_block_io"] and properties["has_subgroup_matrix_multiply_accumulate"], \
238+
"Target do not support blocked load/mma"
238239
return XPUBackend.AdvancedPath.make_ttgir(mod, metadata, opt)
239240

240241
passes.ttir.add_convert_to_ttgpuir(pm, "xpu", opt.num_warps, opt.threads_per_warp, opt.num_ctas)

third_party/intel/lib/TritonIntelGPUToLLVM/TritonGPUToLLVM.cpp

Lines changed: 8 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -83,12 +83,14 @@ struct ConvertTritonGPUToLLVM
8383
ModuleOp mod = getOperation();
8484

8585
bool isAdvancedPathEnabled =
86-
mod->hasAttr(triton::gpu::intel::TritonIntelGPUDialect::
87-
getSupportSG2DBlockAttrName()) &&
88-
mod->hasAttr(triton::gpu::intel::TritonIntelGPUDialect::
89-
getSupportDPASAttrName()) &&
90-
(mlir::triton::tools::getBoolEnv("TRITON_INTEL_ADVANCED_PATH") ||
91-
advancedPath);
86+
mlir::triton::tools::getBoolEnv("TRITON_INTEL_ADVANCED_PATH") ||
87+
advancedPath;
88+
if (isAdvancedPathEnabled)
89+
assert(mod->hasAttr(triton::gpu::intel::TritonIntelGPUDialect::
90+
getSupportSG2DBlockAttrName()) &&
91+
mod->hasAttr(triton::gpu::intel::TritonIntelGPUDialect::
92+
getSupportDPASAttrName()) &&
93+
"Target do not support blocked load/mma");
9294
mlir::triton::intel::TritonGPUToLLVMPipelineManager pipelineManager(
9395
mod, context, isAdvancedPathEnabled);
9496
mlir::LowerToLLVMOptions option(context);

0 commit comments

Comments
 (0)