@@ -4617,10 +4617,10 @@ index 05f230088..b4154971e 100644
4617
4617
// This API doesn't have a corresponding API in NCCL and implemented as
4618
4618
diff --git a/xla/service/gpu/spir_compiler.cc b/xla/service/gpu/spir_compiler.cc
4619
4619
new file mode 100644
4620
- index 000000000..1c33d7e50
4620
+ index 000000000..bffa2d412
4621
4621
--- /dev/null
4622
4622
+++ b/xla/service/gpu/spir_compiler.cc
4623
- @@ -0,0 +1,290 @@
4623
+ @@ -0,0 +1,289 @@
4624
4624
+ /* Copyright (c) 2023 Intel Corporation
4625
4625
+
4626
4626
+ Licensed under the Apache License, Version 2.0 (the "License");
@@ -4797,16 +4797,15 @@ index 000000000..1c33d7e50
4797
4797
+
4798
4798
+ // For frontend debugging.
4799
4799
+ FrontendAttributes frontend_attributes;
4800
- + if (IsXeHPC(nullptr)) {
4801
- + frontend_attributes.mutable_map()->emplace("IsXeHPC", "True");
4802
- + }
4803
- + if (HasXMX(nullptr)) {
4804
- + frontend_attributes.mutable_map()->emplace("HasXMX", "True");
4800
+ + bool is_xetla_hardware_support = IsXetlaHardwareSupport();
4801
+ + if (is_xetla_hardware_support) {
4802
+ + frontend_attributes.mutable_map()->emplace("is_xetla_hardware_support",
4803
+ + "True");
4805
4804
+ }
4806
4805
+ hlo_module->add_frontend_attributes(frontend_attributes);
4807
4806
+ bool use_mha = true;
4808
4807
+ TF_CHECK_OK(tsl::ReadBoolFromEnvVar("MHA", true, &use_mha));
4809
- + if (use_mha && IsXetlaHardwareSupport() ) {
4808
+ + if (use_mha && is_xetla_hardware_support ) {
4810
4809
+ HloPassPipeline mha_fusion_pipeline("multi-headed attention fusion");
4811
4810
+ const DebugOptions& debug_options = hlo_module->config().debug_options();
4812
4811
+ // The LayoutAssignment pass may leave behind kCopy instructions which are
0 commit comments