Skip to content

Commit 02d8168

Browse files
badersteffenlarsen
bader
and
steffenlarsen
authored
[SYCL] Run CompileTimePropertiesPass early in the pipeline (#20602)
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. --------- Co-authored-by: Steffen Larsen <[email protected]>
1 parent b765a2e commit 02d8168

File tree

4 files changed

+48
-31
lines changed

4 files changed

+48
-31
lines changed

‎clang/lib/CodeGen/BackendUtil.cpp‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1109,6 +1109,9 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
11091109
MPM.addPass(SYCLPropagateJointMatrixUsagePass());
11101110
// Lowers static/dynamic local memory builtin calls.
11111111
MPM.addPass(SYCLLowerWGLocalMemoryPass());
1112+
// Compile-time properties pass must create standard metadata as early
1113+
// as possible to make them available for other passes.
1114+
MPM.addPass(CompileTimePropertiesPass());
11121115
});
11131116
else if (LangOpts.SYCLIsHost && !LangOpts.SYCLESIMDBuildHostCode)
11141117
PB.registerPipelineStartEPCallback(
@@ -1271,9 +1274,6 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
12711274
MPM.addPass(SPIRITTAnnotationsPass());
12721275
}
12731276

1274-
// Process properties and annotations
1275-
MPM.addPass(CompileTimePropertiesPass());
1276-
12771277
// Record SYCL aspect names (this should come after propagating aspects
12781278
// and before cleaning up metadata)
12791279
MPM.addPass(RecordSYCLAspectNamesPass());

‎clang/test/CodeGenSYCL/kernel-early-optimization-pipeline.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -9,6 +9,7 @@
99
// CHECK: SYCLPropagateAspectsUsagePass
1010
// CHECK: SYCLPropagateJointMatrixUsagePass
1111
// CHECK: SYCLLowerWGLocalMemoryPass
12+
// CHECK: CompileTimePropertiesPass
1213
// CHECK: InferFunctionAttrsPass
1314
// CHECK: AlwaysInlinerPass
1415
// CHECK: ModuleInlinerWrapperPass
@@ -17,7 +18,6 @@
1718
// CHECK: SYCLMutatePrintfAddrspacePass
1819
// CHECK: SYCLPropagateAspectsUsagePass
1920
// CHECK: SYCLAddOptLevelAttributePass
20-
// CHECK: CompileTimePropertiesPass
2121
// CHECK: RecordSYCLAspectNamesPass
2222
// CHECK: CleanupSYCLMetadataPass
2323
//

‎llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp‎

Lines changed: 20 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -943,16 +943,28 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation(
943943
LLVMContext &Ctx = M.getContext();
944944
unsigned MDKindID = Ctx.getMDKindID(SpirvDecorMdKind);
945945
if (!FPGAProp && llvm::isa<llvm::Instruction>(IntrInst->getArgOperand(0))) {
946-
// If there are no annotations other than cache controls we can apply the
947-
// controls to the pointer and remove the intrinsic.
946+
// Find all load/store instructions using the pointer being annotated and
947+
// apply the cache control metadata to them.
948+
SmallVector<std::pair<Instruction *, int>, 8> TargetedInstList;
949+
getUserListIgnoringCast<LoadInst>(IntrInst, TargetedInstList);
950+
getUserListIgnoringCast<StoreInst>(IntrInst, TargetedInstList);
951+
getUserListIgnoringCast<MemTransferInst>(IntrInst, TargetedInstList);
952+
for (const auto &[Inst, MDVal] : TargetedInstList) {
953+
// Merge with existing metadata if present.
954+
SmallVector<Metadata *, 8> MDOps;
955+
if (MDNode *CurrentMD = Inst->getMetadata(MDKindID))
956+
for (Metadata *Op : CurrentMD->operands())
957+
MDOps.push_back(Op);
958+
for (Metadata *Op : MDOpsCacheProp)
959+
MDOps.push_back(Op);
960+
MDOps.push_back(ConstantAsMetadata::get(Constant::getIntegerValue(
961+
Type::getInt32Ty(Ctx), APInt(32, MDVal))));
962+
Inst->setMetadata(MDKindID, MDTuple::get(Ctx, MDOps));
963+
}
964+
// Replace all uses of ptr.annotations intrinsic with first operand and
965+
// delete the original intrinsic.
948966
Instruction *PtrInstr = cast<Instruction>(IntrInst->getArgOperand(0));
949-
if (MDNode *CurrentMD = PtrInstr->getMetadata(MDKindID))
950-
for (Metadata *Op : CurrentMD->operands())
951-
MDOpsCacheProp.push_back(Op);
952-
PtrInstr->setMetadata(MDKindID, MDTuple::get(Ctx, MDOpsCacheProp));
953-
// Replace all uses of IntrInst with first operand
954967
IntrInst->replaceAllUsesWith(PtrInstr);
955-
// Delete the original IntrInst
956968
RemovableAnnotations.push_back(IntrInst);
957969
} else {
958970
// If there were FPGA annotations then we retain the original intrinsic

‎sycl/test/check_device_code/extensions/properties/properties_cache_control.cpp‎

Lines changed: 24 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -171,57 +171,62 @@ SYCL_EXTERNAL void annotated_ptr_func_param_test(float *p) {
171171
}
172172

173173
// CHECK: spir_func{{.*}}annotated_ptr_func_param_test
174-
// CHECK: {{.*}}call ptr addrspace(4) @llvm.ptr.annotation.p4.p1{{.*}}!spirv.Decorations [[WHINT:.*]]
174+
// CHECK: store float 4.200000e+01, ptr addrspace(4) %{{.*}}, !spirv.Decorations ![[WHINT:[0-9]+]]
175175
// CHECK: ret void
176176

177177
// CHECK: spir_kernel{{.*}}cache_control_read_hint_func
178-
// CHECK: {{.*}}addrspacecast ptr addrspace(1){{.*}}!spirv.Decorations [[RHINT:.*]]
178+
// CHECK: store float 5.500000e+01, ptr addrspace(1) %{{.*}}, !spirv.Decorations ![[RHINT:[0-9]+]]
179179
// CHECK: ret void
180180

181181
// CHECK: spir_kernel{{.*}}cache_control_read_assertion_func
182-
// CHECK: {{.*}}addrspacecast ptr addrspace(1){{.*}}!spirv.Decorations [[RASSERT:.*]]
182+
// CHECK: store i32 66, ptr addrspace(1) %{{.*}}, !spirv.Decorations ![[RASSERT:[0-9]+]]
183183
// CHECK: ret void
184184

185185
// CHECK: spir_kernel{{.*}}cache_control_write_hint_func
186-
// CHECK: {{.*}}addrspacecast ptr addrspace(1){{.*}}!spirv.Decorations [[WHINT]]
186+
// CHECK: store float 7.700000e+01, ptr addrspace(1) %{{.*}}, !spirv.Decorations ![[WHINT]]
187187
// CHECK: ret void
188188

189189
// CHECK: spir_kernel{{.*}}cache_control_read_write_func
190-
// CHECK: {{.*}}addrspacecast ptr addrspace(1){{.*}}!spirv.Decorations [[RWHINT:.*]]
190+
// CHECK: store float 7.700000e+01, ptr addrspace(1) %{{.*}}, !spirv.Decorations ![[RWHINT:[0-9]+]]
191191
// CHECK: ret void
192192

193193
// CHECK: spir_kernel{{.*}}cache_control_load_store_func
194-
// CHECK: {{.*}}getelementptr{{.*}}addrspace(4){{.*}}!spirv.Decorations [[LDSTHINT_A:.*]]
195-
// CHECK: {{.*}}getelementptr{{.*}}addrspace(4){{.*}}!spirv.Decorations [[LDSTHINT_B:.*]]
194+
// CHECK: store double 1.000000e+00, ptr addrspace(1) %[[PTR_A:.*]], align 8{{.*}}, !spirv.Decorations ![[STHINT_A:[0-9]+]]
195+
// CHECK: store double 1.000000e+00, ptr addrspace(1) %[[PTR_B:.*]], align 8{{.*}}, !spirv.Decorations ![[STHINT_B:[0-9]+]]
196+
// CHECK: load double, ptr addrspace(1) %[[PTR_A]], align 8{{.*}}, !spirv.Decorations ![[LDHINT_A:[0-9]+]]
197+
// CHECK: load double, ptr addrspace(1) %[[PTR_B]], align 8{{.*}}, !spirv.Decorations ![[LDHINT_B:[0-9]+]]
196198
// CHECK: ret void
197199

198-
// CHECK: [[WHINT]] = !{[[WHINT1:.*]], [[WHINT2:.*]], [[WHINT3:.*]], [[WHINT4:.*]]}
200+
// CHECK: [[WHINT]] = !{[[WHINT1:.*]], [[WHINT2:.*]], [[WHINT3:.*]], [[WHINT4:.*]], i32 1}
199201
// CHECK: [[WHINT1]] = !{i32 6443, i32 3, i32 3}
200202
// CHECK: [[WHINT2]] = !{i32 6443, i32 0, i32 1}
201203
// CHECK: [[WHINT3]] = !{i32 6443, i32 1, i32 2}
202204
// CHECK: [[WHINT4]] = !{i32 6443, i32 2, i32 2}
203205

204-
// CHECK: [[RHINT]] = !{[[RHINT1:.*]], [[RHINT2:.*]], [[RHINT3:.*]]}
206+
// CHECK: [[RHINT]] = !{[[RHINT1:.*]], [[RHINT2:.*]], [[RHINT3:.*]], i32 1}
205207
// CHECK: [[RHINT1]] = !{i32 6442, i32 1, i32 0}
206208
// CHECK: [[RHINT2]] = !{i32 6442, i32 2, i32 0}
207209
// CHECK: [[RHINT3]] = !{i32 6442, i32 0, i32 1}
208210

209-
// CHECK: [[RASSERT]] = !{[[RASSERT1:.*]], [[RASSERT2:.*]], [[RASSERT3:.*]]}
211+
// CHECK: [[RASSERT]] = !{[[RASSERT1:.*]], [[RASSERT2:.*]], [[RASSERT3:.*]], i32 1}
210212
// CHECK: [[RASSERT1]] = !{i32 6442, i32 1, i32 3}
211213
// CHECK: [[RASSERT2]] = !{i32 6442, i32 2, i32 3}
212214
// CHECK: [[RASSERT3]] = !{i32 6442, i32 0, i32 4}
213215

214-
// CHECK: [[RWHINT]] = !{[[RWHINT1:.*]], [[RWHINT2:.*]], [[RWHINT3:.*]]}
216+
// CHECK: [[RWHINT]] = !{[[RWHINT1:.*]], [[RWHINT2:.*]], [[RWHINT3:.*]], i32 1}
215217
// CHECK: [[RWHINT1]] = !{i32 6442, i32 2, i32 1}
216218
// CHECK: [[RWHINT2]] = !{i32 6442, i32 3, i32 4}
217219
// CHECK: [[RWHINT3]] = !{i32 6443, i32 3, i32 1}
218220

219-
// CHECK: [[LDSTHINT_A]] = !{[[RHINT1]], [[RHINT2]], [[RHINT3]], [[LDSTHINT_A1:.*]], [[LDSTHINT_A2:.*]], [[LDSTHINT_A3:.*]]}
220-
// CHECK: [[LDSTHINT_A1]] = !{i32 6443, i32 0, i32 0}
221-
// CHECK: [[LDSTHINT_A2]] = !{i32 6443, i32 1, i32 0}
222-
// CHECK: [[LDSTHINT_A3]] = !{i32 6443, i32 2, i32 0}
221+
// CHECK: [[STHINT_A]] = !{[[STHINT_A1:.*]], [[STHINT_A2:.*]], [[STHINT_A3:.*]], i32 1}
222+
// CHECK: [[STHINT_A1]] = !{i32 6443, i32 0, i32 0}
223+
// CHECK: [[STHINT_A2]] = !{i32 6443, i32 1, i32 0}
224+
// CHECK: [[STHINT_A3]] = !{i32 6443, i32 2, i32 0}
223225

224-
// CHECK: [[LDSTHINT_B]] = !{[[LDSTHINT_B1:.*]], [[RWHINT1]], [[LDSTHINT_B2:.*]], [[LDSTHINT_A2]], [[LDSTHINT_A3]], [[LDSTHINT_B3:.*]]}
225-
// CHECK: [[LDSTHINT_B1]] = !{i32 6442, i32 1, i32 1}
226-
// CHECK: [[LDSTHINT_B2]] = !{i32 6442, i32 0, i32 2}
227-
// CHECK: [[LDSTHINT_B3]] = !{i32 6443, i32 0, i32 2}
226+
// CHECK: [[STHINT_B]] = !{[[STHINT_A2]], [[STHINT_A3]], [[STHINT_B1:.*]], i32 1}
227+
// CHECK: [[STHINT_B1]] = !{i32 6443, i32 0, i32 2}
228+
229+
// CHECK: [[LDHINT_A]] = !{[[RHINT1]], [[RHINT2]], [[RHINT3]], i32 0}
230+
// CHECK: [[LDHINT_B]] = !{[[LDHINT_B1:.*]], [[RWHINT1]], [[LDHINT_B2:.*]], i32 0}
231+
// CHECK: [[LDHINT_B1]] = !{i32 6442, i32 1, i32 1}
232+
// CHECK: [[LDHINT_B2]] = !{i32 6442, i32 0, i32 2}

0 commit comments

Comments
 (0)