From cf33f60e633c48db01c6d05188efbcf14a3ec239 Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Fri, 7 Nov 2025 15:02:29 -0800 Subject: [PATCH 1/6] [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 // From 2c8d041905ffba4c0d74e09782887475e9024885 Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Thu, 20 Nov 2025 18:01:39 -0800 Subject: [PATCH 2/6] Parametrize CompileTimePropertiesPass with a flag controlling cache controls handling. --- clang/lib/CodeGen/BackendUtil.cpp | 5 ++++- .../test/CodeGenSYCL/kernel-early-optimization-pipeline.cpp | 1 + llvm/include/llvm/SYCLLowerIR/CompileTimePropertiesPass.h | 4 ++++ llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp | 4 +++- 4 files changed, 12 insertions(+), 2 deletions(-) diff --git a/clang/lib/CodeGen/BackendUtil.cpp b/clang/lib/CodeGen/BackendUtil.cpp index 59b6482bd4648..13e469aa2d468 100644 --- a/clang/lib/CodeGen/BackendUtil.cpp +++ b/clang/lib/CodeGen/BackendUtil.cpp @@ -1083,7 +1083,7 @@ void EmitAssemblyHelper::RunOptimizationPipeline( 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()); + MPM.addPass(CompileTimePropertiesPass(false /*ConvertCacheControls*/)); }); else if (LangOpts.SYCLIsHost && !LangOpts.SYCLESIMDBuildHostCode) PB.registerPipelineStartEPCallback( @@ -1245,6 +1245,9 @@ 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 1af3368350bed..4eba192cac322 100644 --- a/clang/test/CodeGenSYCL/kernel-early-optimization-pipeline.cpp +++ b/clang/test/CodeGenSYCL/kernel-early-optimization-pipeline.cpp @@ -18,6 +18,7 @@ // CHECK: SYCLMutatePrintfAddrspacePass // CHECK: SYCLPropagateAspectsUsagePass // CHECK: SYCLAddOptLevelAttributePass +// CHECK: CompileTimePropertiesPass // CHECK: RecordSYCLAspectNamesPass // CHECK: CleanupSYCLMetadataPass // diff --git a/llvm/include/llvm/SYCLLowerIR/CompileTimePropertiesPass.h b/llvm/include/llvm/SYCLLowerIR/CompileTimePropertiesPass.h index 46cbae34debd8..1b4b17c96710e 100644 --- a/llvm/include/llvm/SYCLLowerIR/CompileTimePropertiesPass.h +++ b/llvm/include/llvm/SYCLLowerIR/CompileTimePropertiesPass.h @@ -30,6 +30,8 @@ class IntrinsicInst; class CompileTimePropertiesPass : public PassInfoMixin { public: + CompileTimePropertiesPass(bool ConvertCacheControls = true) + : ConvertCacheControls(ConvertCacheControls) {} // Enriches the module with metadata that describes the found variables for // the SPIRV-LLVM Translator. PreservedAnalyses run(Module &M, ModuleAnalysisManager &MAM); @@ -48,6 +50,8 @@ class CompileTimePropertiesPass // This allows reuse for annotations with the same generated annotation // strings. std::unordered_map ReusableAnnotStrings; + + bool ConvertCacheControls; }; namespace detail { diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp index 13a92f83e0117..776f3df368c2a 100644 --- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp +++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp @@ -852,6 +852,8 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation( // Handle cache control properties if ((*PropName).starts_with("sycl-cache-")) { + if (!ConvertCacheControls) + continue; CacheProp = true; auto DecorValue = PropVal; uint32_t AttrVal; @@ -906,7 +908,7 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation( // If there are no other annotations (except "alignment") then there is no // reason to keep the original intrinsic, so replace it with the first operand // and mark it for removal. - if (!CacheProp && !FPGAProp) { + if (!CacheProp && !FPGAProp && ConvertCacheControls) { IntrInst->replaceAllUsesWith(IntrInst->getOperand(0)); RemovableAnnotations.push_back(IntrInst); return true; From c102d9671edcc6ea8fe9f67c415eb698ab58e89c Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Fri, 21 Nov 2025 15:11:49 -0800 Subject: [PATCH 3/6] Revert "Parametrize CompileTimePropertiesPass with a flag controlling cache controls handling." This reverts commit 2c8d041905ffba4c0d74e09782887475e9024885. --- clang/lib/CodeGen/BackendUtil.cpp | 5 +---- .../test/CodeGenSYCL/kernel-early-optimization-pipeline.cpp | 1 - llvm/include/llvm/SYCLLowerIR/CompileTimePropertiesPass.h | 4 ---- llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp | 4 +--- 4 files changed, 2 insertions(+), 12 deletions(-) diff --git a/clang/lib/CodeGen/BackendUtil.cpp b/clang/lib/CodeGen/BackendUtil.cpp index 13e469aa2d468..59b6482bd4648 100644 --- a/clang/lib/CodeGen/BackendUtil.cpp +++ b/clang/lib/CodeGen/BackendUtil.cpp @@ -1083,7 +1083,7 @@ void EmitAssemblyHelper::RunOptimizationPipeline( 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(false /*ConvertCacheControls*/)); + MPM.addPass(CompileTimePropertiesPass()); }); else if (LangOpts.SYCLIsHost && !LangOpts.SYCLESIMDBuildHostCode) PB.registerPipelineStartEPCallback( @@ -1245,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 4eba192cac322..1af3368350bed 100644 --- a/clang/test/CodeGenSYCL/kernel-early-optimization-pipeline.cpp +++ b/clang/test/CodeGenSYCL/kernel-early-optimization-pipeline.cpp @@ -18,7 +18,6 @@ // CHECK: SYCLMutatePrintfAddrspacePass // CHECK: SYCLPropagateAspectsUsagePass // CHECK: SYCLAddOptLevelAttributePass -// CHECK: CompileTimePropertiesPass // CHECK: RecordSYCLAspectNamesPass // CHECK: CleanupSYCLMetadataPass // diff --git a/llvm/include/llvm/SYCLLowerIR/CompileTimePropertiesPass.h b/llvm/include/llvm/SYCLLowerIR/CompileTimePropertiesPass.h index 1b4b17c96710e..46cbae34debd8 100644 --- a/llvm/include/llvm/SYCLLowerIR/CompileTimePropertiesPass.h +++ b/llvm/include/llvm/SYCLLowerIR/CompileTimePropertiesPass.h @@ -30,8 +30,6 @@ class IntrinsicInst; class CompileTimePropertiesPass : public PassInfoMixin { public: - CompileTimePropertiesPass(bool ConvertCacheControls = true) - : ConvertCacheControls(ConvertCacheControls) {} // Enriches the module with metadata that describes the found variables for // the SPIRV-LLVM Translator. PreservedAnalyses run(Module &M, ModuleAnalysisManager &MAM); @@ -50,8 +48,6 @@ class CompileTimePropertiesPass // This allows reuse for annotations with the same generated annotation // strings. std::unordered_map ReusableAnnotStrings; - - bool ConvertCacheControls; }; namespace detail { diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp index 776f3df368c2a..13a92f83e0117 100644 --- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp +++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp @@ -852,8 +852,6 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation( // Handle cache control properties if ((*PropName).starts_with("sycl-cache-")) { - if (!ConvertCacheControls) - continue; CacheProp = true; auto DecorValue = PropVal; uint32_t AttrVal; @@ -908,7 +906,7 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation( // If there are no other annotations (except "alignment") then there is no // reason to keep the original intrinsic, so replace it with the first operand // and mark it for removal. - if (!CacheProp && !FPGAProp && ConvertCacheControls) { + if (!CacheProp && !FPGAProp) { IntrInst->replaceAllUsesWith(IntrInst->getOperand(0)); RemovableAnnotations.push_back(IntrInst); return true; From 0914222f1bc4783ed04b3a974a54027099b2f071 Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Fri, 21 Nov 2025 15:10:38 -0800 Subject: [PATCH 4/6] Apply metadata to load/store instructions rather than to the instructions producing a pointer argument for load/store instructions. --- .../SYCLLowerIR/CompileTimePropertiesPass.cpp | 29 ++++++++++++++----- 1 file changed, 21 insertions(+), 8 deletions(-) diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp index 13a92f83e0117..a05702e2bbf14 100644 --- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp +++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp @@ -948,16 +948,29 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation( LLVMContext &Ctx = M.getContext(); unsigned MDKindID = Ctx.getMDKindID(SPIRV_DECOR_MD_KIND); if (!FPGAProp && llvm::isa(IntrInst->getArgOperand(0))) { - // If there are no annotations other than cache controls we can apply the - // controls to the pointer and remove the intrinsic. + // Find all load/store instructions using the pointer being annotated and + // apply the cache control metadata to them. + SmallVector, 8> TargetedInstList; + getUserListIgnoringCast(IntrInst, TargetedInstList); + getUserListIgnoringCast(IntrInst, TargetedInstList); + getUserListIgnoringCast(IntrInst, TargetedInstList); + for (const auto &Pair : TargetedInstList) { + auto *Inst = Pair.first; + // Merge with existing metadata if present. + SmallVector MDOps; + if (MDNode *CurrentMD = Inst->getMetadata(MDKindID)) + for (Metadata *Op : CurrentMD->operands()) + MDOps.push_back(Op); + for (Metadata *Op : MDOpsCacheProp) + MDOps.push_back(Op); + MDOps.push_back(ConstantAsMetadata::get(Constant::getIntegerValue( + Type::getInt32Ty(Ctx), APInt(32, Pair.second)))); + Inst->setMetadata(MDKindID, MDTuple::get(Ctx, MDOps)); + } + // Replace all uses of ptr.annotations intrinsic with first operand and + // delete the original intrinsic. Instruction *PtrInstr = cast(IntrInst->getArgOperand(0)); - if (MDNode *CurrentMD = PtrInstr->getMetadata(MDKindID)) - for (Metadata *Op : CurrentMD->operands()) - MDOpsCacheProp.push_back(Op); - PtrInstr->setMetadata(MDKindID, MDTuple::get(Ctx, MDOpsCacheProp)); - // Replace all uses of IntrInst with first operand IntrInst->replaceAllUsesWith(PtrInstr); - // Delete the original IntrInst RemovableAnnotations.push_back(IntrInst); } else { // If there were FPGA annotations then we retain the original intrinsic From 1026c0efa14da0e22bb4d7ab6a3d489fc42bb2b8 Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Fri, 21 Nov 2025 16:23:06 -0800 Subject: [PATCH 5/6] Update properties_cache_control test checks. --- .../properties/properties_cache_control.cpp | 43 +++++++++++-------- 1 file changed, 24 insertions(+), 19 deletions(-) diff --git a/sycl/test/check_device_code/extensions/properties/properties_cache_control.cpp b/sycl/test/check_device_code/extensions/properties/properties_cache_control.cpp index 6497dd0e70730..6d1b21b41fe37 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_cache_control.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_cache_control.cpp @@ -171,57 +171,62 @@ SYCL_EXTERNAL void annotated_ptr_func_param_test(float *p) { } // CHECK: spir_func{{.*}}annotated_ptr_func_param_test -// CHECK: {{.*}}call ptr addrspace(4) @llvm.ptr.annotation.p4.p1{{.*}}!spirv.Decorations [[WHINT:.*]] +// CHECK: store float 4.200000e+01, ptr addrspace(4) %{{.*}}, !spirv.Decorations ![[WHINT:[0-9]+]] // CHECK: ret void // CHECK: spir_kernel{{.*}}cache_control_read_hint_func -// CHECK: {{.*}}addrspacecast ptr addrspace(1){{.*}}!spirv.Decorations [[RHINT:.*]] +// CHECK: store float 5.500000e+01, ptr addrspace(1) %{{.*}}, !spirv.Decorations ![[RHINT:[0-9]+]] // CHECK: ret void // CHECK: spir_kernel{{.*}}cache_control_read_assertion_func -// CHECK: {{.*}}addrspacecast ptr addrspace(1){{.*}}!spirv.Decorations [[RASSERT:.*]] +// CHECK: store i32 66, ptr addrspace(1) %{{.*}}, !spirv.Decorations ![[RASSERT:[0-9]+]] // CHECK: ret void // CHECK: spir_kernel{{.*}}cache_control_write_hint_func -// CHECK: {{.*}}addrspacecast ptr addrspace(1){{.*}}!spirv.Decorations [[WHINT]] +// CHECK: store float 7.700000e+01, ptr addrspace(1) %{{.*}}, !spirv.Decorations ![[WHINT]] // CHECK: ret void // CHECK: spir_kernel{{.*}}cache_control_read_write_func -// CHECK: {{.*}}addrspacecast ptr addrspace(1){{.*}}!spirv.Decorations [[RWHINT:.*]] +// CHECK: store float 7.700000e+01, ptr addrspace(1) %{{.*}}, !spirv.Decorations ![[RWHINT:[0-9]+]] // CHECK: ret void // CHECK: spir_kernel{{.*}}cache_control_load_store_func -// CHECK: {{.*}}getelementptr{{.*}}addrspace(4){{.*}}!spirv.Decorations [[LDSTHINT_A:.*]] -// CHECK: {{.*}}getelementptr{{.*}}addrspace(4){{.*}}!spirv.Decorations [[LDSTHINT_B:.*]] +// CHECK: store double 1.000000e+00, ptr addrspace(1) %[[PTR_A:.*]], align 8{{.*}}, !spirv.Decorations ![[STHINT_A:[0-9]+]] +// CHECK: store double 1.000000e+00, ptr addrspace(1) %[[PTR_B:.*]], align 8{{.*}}, !spirv.Decorations ![[STHINT_B:[0-9]+]] +// CHECK: load double, ptr addrspace(1) %[[PTR_A]], align 8{{.*}}, !spirv.Decorations ![[LDHINT_A:[0-9]+]] +// CHECK: load double, ptr addrspace(1) %[[PTR_B]], align 8{{.*}}, !spirv.Decorations ![[LDHINT_B:[0-9]+]] // CHECK: ret void -// CHECK: [[WHINT]] = !{[[WHINT1:.*]], [[WHINT2:.*]], [[WHINT3:.*]], [[WHINT4:.*]]} +// CHECK: [[WHINT]] = !{[[WHINT1:.*]], [[WHINT2:.*]], [[WHINT3:.*]], [[WHINT4:.*]], i32 1} // CHECK: [[WHINT1]] = !{i32 6443, i32 3, i32 3} // CHECK: [[WHINT2]] = !{i32 6443, i32 0, i32 1} // CHECK: [[WHINT3]] = !{i32 6443, i32 1, i32 2} // CHECK: [[WHINT4]] = !{i32 6443, i32 2, i32 2} -// CHECK: [[RHINT]] = !{[[RHINT1:.*]], [[RHINT2:.*]], [[RHINT3:.*]]} +// CHECK: [[RHINT]] = !{[[RHINT1:.*]], [[RHINT2:.*]], [[RHINT3:.*]], i32 1} // CHECK: [[RHINT1]] = !{i32 6442, i32 1, i32 0} // CHECK: [[RHINT2]] = !{i32 6442, i32 2, i32 0} // CHECK: [[RHINT3]] = !{i32 6442, i32 0, i32 1} -// CHECK: [[RASSERT]] = !{[[RASSERT1:.*]], [[RASSERT2:.*]], [[RASSERT3:.*]]} +// CHECK: [[RASSERT]] = !{[[RASSERT1:.*]], [[RASSERT2:.*]], [[RASSERT3:.*]], i32 1} // CHECK: [[RASSERT1]] = !{i32 6442, i32 1, i32 3} // CHECK: [[RASSERT2]] = !{i32 6442, i32 2, i32 3} // CHECK: [[RASSERT3]] = !{i32 6442, i32 0, i32 4} -// CHECK: [[RWHINT]] = !{[[RWHINT1:.*]], [[RWHINT2:.*]], [[RWHINT3:.*]]} +// CHECK: [[RWHINT]] = !{[[RWHINT1:.*]], [[RWHINT2:.*]], [[RWHINT3:.*]], i32 1} // CHECK: [[RWHINT1]] = !{i32 6442, i32 2, i32 1} // CHECK: [[RWHINT2]] = !{i32 6442, i32 3, i32 4} // CHECK: [[RWHINT3]] = !{i32 6443, i32 3, i32 1} -// CHECK: [[LDSTHINT_A]] = !{[[RHINT1]], [[RHINT2]], [[RHINT3]], [[LDSTHINT_A1:.*]], [[LDSTHINT_A2:.*]], [[LDSTHINT_A3:.*]]} -// CHECK: [[LDSTHINT_A1]] = !{i32 6443, i32 0, i32 0} -// CHECK: [[LDSTHINT_A2]] = !{i32 6443, i32 1, i32 0} -// CHECK: [[LDSTHINT_A3]] = !{i32 6443, i32 2, i32 0} +// CHECK: [[STHINT_A]] = !{[[STHINT_A1:.*]], [[STHINT_A2:.*]], [[STHINT_A3:.*]], i32 1} +// CHECK: [[STHINT_A1]] = !{i32 6443, i32 0, i32 0} +// CHECK: [[STHINT_A2]] = !{i32 6443, i32 1, i32 0} +// CHECK: [[STHINT_A3]] = !{i32 6443, i32 2, i32 0} -// CHECK: [[LDSTHINT_B]] = !{[[LDSTHINT_B1:.*]], [[RWHINT1]], [[LDSTHINT_B2:.*]], [[LDSTHINT_A2]], [[LDSTHINT_A3]], [[LDSTHINT_B3:.*]]} -// CHECK: [[LDSTHINT_B1]] = !{i32 6442, i32 1, i32 1} -// CHECK: [[LDSTHINT_B2]] = !{i32 6442, i32 0, i32 2} -// CHECK: [[LDSTHINT_B3]] = !{i32 6443, i32 0, i32 2} +// CHECK: [[STHINT_B]] = !{[[STHINT_A2]], [[STHINT_A3]], [[STHINT_B1:.*]], i32 1} +// CHECK: [[STHINT_B1]] = !{i32 6443, i32 0, i32 2} + +// CHECK: [[LDHINT_A]] = !{[[RHINT1]], [[RHINT2]], [[RHINT3]], i32 0} +// CHECK: [[LDHINT_B]] = !{[[LDHINT_B1:.*]], [[RWHINT1]], [[LDHINT_B2:.*]], i32 0} +// CHECK: [[LDHINT_B1]] = !{i32 6442, i32 1, i32 1} +// CHECK: [[LDHINT_B2]] = !{i32 6442, i32 0, i32 2} From 178a4452db31482d67617e7f086bc67a626759c1 Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Mon, 24 Nov 2025 06:48:52 -0800 Subject: [PATCH 6/6] Update llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp Co-authored-by: Steffen Larsen --- llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp index 8e3626d34672a..ca79ecc028915 100644 --- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp +++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp @@ -949,8 +949,7 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation( getUserListIgnoringCast(IntrInst, TargetedInstList); getUserListIgnoringCast(IntrInst, TargetedInstList); getUserListIgnoringCast(IntrInst, TargetedInstList); - for (const auto &Pair : TargetedInstList) { - auto *Inst = Pair.first; + for (const auto &[Inst, MDVal] : TargetedInstList) { // Merge with existing metadata if present. SmallVector MDOps; if (MDNode *CurrentMD = Inst->getMetadata(MDKindID)) @@ -959,7 +958,7 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation( for (Metadata *Op : MDOpsCacheProp) MDOps.push_back(Op); MDOps.push_back(ConstantAsMetadata::get(Constant::getIntegerValue( - Type::getInt32Ty(Ctx), APInt(32, Pair.second)))); + Type::getInt32Ty(Ctx), APInt(32, MDVal)))); Inst->setMetadata(MDKindID, MDTuple::get(Ctx, MDOps)); } // Replace all uses of ptr.annotations intrinsic with first operand and