From cf33f60e633c48db01c6d05188efbcf14a3ec239 Mon Sep 17 00:00:00 2001
From: Alexey Bader <alexey.bader@intel.com>
Date: Fri, 7 Nov 2025 15:02:29 -0800
Subject: [PATCH] [SYCL] Run CompileTimePropertiesPass early in the pipeline

Some compile time properties work as a replacement for kernel
attributes. For example, work_group_size semantics must be identical to
sycl::reqd_work_group_size kernel attribute. The problem is kernel
attributes are lowered to LLVM metadata by Clang, but work_group_size
represented as an LLVM attribute.
CompileTimePropertiesPass converts attribute to canonical metadata
representation, but does it late in the opimization pipeline.
This patch moves CompileTimePropertiesPass to the beginning of the
optimization pipeline to keep canonical representation for SYCL kernel
attributes information passes via compile-time properties.
---
 clang/lib/CodeGen/BackendUtil.cpp                           | 6 +++---
 .../test/CodeGenSYCL/kernel-early-optimization-pipeline.cpp | 2 +-
 2 files changed, 4 insertions(+), 4 deletions(-)

diff --git a/clang/lib/CodeGen/BackendUtil.cpp b/clang/lib/CodeGen/BackendUtil.cpp
index e510e10575026..59b6482bd4648 100644
--- a/clang/lib/CodeGen/BackendUtil.cpp
+++ b/clang/lib/CodeGen/BackendUtil.cpp
@@ -1081,6 +1081,9 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
         MPM.addPass(SYCLPropagateJointMatrixUsagePass());
         // Lowers static/dynamic local memory builtin calls.
         MPM.addPass(SYCLLowerWGLocalMemoryPass());
+        // Compile-time properties pass must create standard metadata as early
+        // as possible to make them available for other passes.
+        MPM.addPass(CompileTimePropertiesPass());
       });
     else if (LangOpts.SYCLIsHost && !LangOpts.SYCLESIMDBuildHostCode)
       PB.registerPipelineStartEPCallback(
@@ -1242,9 +1245,6 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
         MPM.addPass(SPIRITTAnnotationsPass());
       }
 
-      // Process properties and annotations
-      MPM.addPass(CompileTimePropertiesPass());
-
       // Record SYCL aspect names (this should come after propagating aspects
       // and before cleaning up metadata)
       MPM.addPass(RecordSYCLAspectNamesPass());
diff --git a/clang/test/CodeGenSYCL/kernel-early-optimization-pipeline.cpp b/clang/test/CodeGenSYCL/kernel-early-optimization-pipeline.cpp
index d352f1bcca39a..1af3368350bed 100644
--- a/clang/test/CodeGenSYCL/kernel-early-optimization-pipeline.cpp
+++ b/clang/test/CodeGenSYCL/kernel-early-optimization-pipeline.cpp
@@ -9,6 +9,7 @@
 // CHECK: SYCLPropagateAspectsUsagePass
 // CHECK: SYCLPropagateJointMatrixUsagePass
 // CHECK: SYCLLowerWGLocalMemoryPass
+// CHECK: CompileTimePropertiesPass
 // CHECK: InferFunctionAttrsPass
 // CHECK: AlwaysInlinerPass
 // CHECK: ModuleInlinerWrapperPass
@@ -17,7 +18,6 @@
 // CHECK: SYCLMutatePrintfAddrspacePass
 // CHECK: SYCLPropagateAspectsUsagePass
 // CHECK: SYCLAddOptLevelAttributePass
-// CHECK: CompileTimePropertiesPass
 // CHECK: RecordSYCLAspectNamesPass
 // CHECK: CleanupSYCLMetadataPass
 //