Review fixes

google · May 28, 2024 · 918bbe6 · 918bbe6
1 parent 48a3e1c
commit 918bbe6
Show file tree

Hide file tree

Showing 3 changed files with 9 additions and 8 deletions.
diff --git a/lib/Compiler.cpp b/lib/Compiler.cpp
@@ -525,6 +525,10 @@ int RunPassPipeline(llvm::Module &M, llvm::raw_svector_ostream *binaryStream) {
     pm.addPass(clspv::AutoPodArgsPass());
     pm.addPass(clspv::DeclarePushConstantsPass());
     pm.addPass(clspv::DefineOpenCLWorkItemBuiltinsPass());
+    // Replace the LLVM intrinsics. This will give them a chance to be better
+    // optimized through the pipeline. It also helps with generic address space
+    // lowering.
+    pm.addPass(clspv::ReplaceLLVMIntrinsicsPass());
 
     // RewritePackedStructsPass will rewrite packed struct types, and
     // ReplacePointerBitcastPass will lower the new packed struct type. So,
@@ -571,9 +575,6 @@ int RunPassPipeline(llvm::Module &M, llvm::raw_svector_ostream *binaryStream) {
     // This pass needs to be after every inlining to make sure we are capable of
     // removing every addrspacecast. It only needs to run if generic addrspace
     // is used.
-
-    pm.addPass(clspv::ReplaceLLVMIntrinsicsPass());
-
     if (clspv::Option::LanguageUsesGenericAddressSpace()) {
       pm.addPass(clspv::ReplaceOpenCLBuiltinPass());
       pm.addPass(clspv::LowerAddrSpaceCastPass());

diff --git a/test/LLVMIntrinsics/descend_into_array.cl b/test/LLVMIntrinsics/descend_into_array.cl
@@ -18,7 +18,7 @@ __kernel void foo(__global outer* out, global outer* in) {
 // CHECK: [[zero:%[0-9a-zA-Z_]+]] = OpConstant {{.*}} 0
 // CHECK: [[dst:%[0-9a-zA-Z_]+]] = OpVariable {{.*}} StorageBuffer
 // CHECK: [[src:%[0-9a-zA-Z_]+]] = OpVariable {{.*}} StorageBuffer
-// CHECK-DAG: [[src_gep:%[0-9a-zA-Z_]+]] = OpAccessChain {{.*}} [[src]] [[zero]] [[zero]]
-// CHECK-DAG: OpLoad {{.*}} [[src_gep]]
-// CHECK-DAG: [[dst_gep:%[0-9a-zA-Z_]+]] = OpAccessChain {{.*}} [[dst]] [[zero]] [[zero]]
-// CHECK-DAG: OpStore [[dst_gep]]
+// CHECK: [[dst_gep:%[0-9a-zA-Z_]+]] = OpAccessChain {{.*}} [[dst]] [[zero]] [[zero]]
+// CHECK: [[src_gep:%[0-9a-zA-Z_]+]] = OpAccessChain {{.*}} [[src]] [[zero]] [[zero]]
+// CHECK: OpLoad {{.*}} [[src_gep]]
+// CHECK: OpStore [[dst_gep]]
diff --git a/test/LLVMIntrinsics/memcpy_from_constant.cl b/test/LLVMIntrinsics/memcpy_from_constant.cl
@@ -51,4 +51,4 @@ void kernel memcpy_from_constant(global float* result) {
 // CHECK: OpStore [[ssbo_gep3]] [[gep3]]
 
 // CHECK: [[ssbo_gep4:%[a-zA-Z0-9_]+]] = OpAccessChain [[ptr_ssbo_float]] [[ssbo]] [[uint_0]] [[uint_4]]
-// CHECK: OpStore [[ssbo_gep4]] [[gep4]]
+// CHECK: OpStore [[ssbo_gep4]] [[gep4]]