[ExecutionEngine] Modify the way we pass IGC flags. (#912)

Passing multiple finalizer flags together seems to have caused issues in some IGC versions as not being recognized. As a result, passing one finalizer flag at a time.
intel · Oct 3, 2024 · fe1535e · fe1535e
1 parent f31a332
commit fe1535e
Show file tree

Hide file tree

Showing 2 changed files with 7 additions and 8 deletions.
diff --git a/lib/ExecutionEngine/LEVELZERORUNTIME/LevelZeroRuntimeWrappers.cpp b/lib/ExecutionEngine/LEVELZERORUNTIME/LevelZeroRuntimeWrappers.cpp
@@ -402,10 +402,10 @@ static ze_module_handle_t loadModule(GPUL0QUEUE *queue, const void *data,
   }
   // enable large register file if needed
   if (getenv("IMEX_ENABLE_LARGE_REG_FILE")) {
-    build_flags +=
-        " -doubleGRF -Xfinalizer -noLocalSplit -Xfinalizer "
-        "-DPASTokenReduction -Xfinalizer -SWSBDepReduction -Xfinalizer "
-        "'-printregusage -enableBCR' ";
+    build_flags += "-doubleGRF -Xfinalizer -noLocalSplit -Xfinalizer "
+                   "-DPASTokenReduction -Xfinalizer -SWSBDepReduction "
+                   "-Xfinalizer -printregusage -Xfinalizer -enableBCR";
+    ;
   }
 
   desc.format = ZE_MODULE_FORMAT_IL_SPIRV;

diff --git a/lib/ExecutionEngine/SYCLRUNTIME/SyclRuntimeWrappers.cpp b/lib/ExecutionEngine/SYCLRUNTIME/SyclRuntimeWrappers.cpp
@@ -234,10 +234,9 @@ static ze_module_handle_t loadModule(GPUSYCLQUEUE *queue, const void *data,
   }
   // enable large register file if needed
   if (getenv("IMEX_ENABLE_LARGE_REG_FILE")) {
-    build_flags +=
-        " -doubleGRF -Xfinalizer -noLocalSplit -Xfinalizer "
-        "-DPASTokenReduction -Xfinalizer -SWSBDepReduction -Xfinalizer "
-        "'-printregusage -enableBCR' ";
+    build_flags += "-doubleGRF -Xfinalizer -noLocalSplit -Xfinalizer "
+                   "-DPASTokenReduction -Xfinalizer -SWSBDepReduction "
+                   "-Xfinalizer -printregusage -Xfinalizer -enableBCR";
   }
   desc.pBuildFlags = build_flags.c_str();
   auto zeDevice = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(