[SYCL] Stop emitting kernel arguments metadata (#2348)

OpenCL kernel arguments metadata - kernel_arg_addr_space, kernel_arg_access_qual, kernel_arg_type, kernel_arg_base_type, kernel_arg_type_qual and kernel_arg_name - are not needed for SYCL. But this causes optimizations to unnecessarily handle them, so stop emitting them. Signed-off-by: Premanand M Rao <premanand.m.rao@intel.com> Keep OpenCL kernel metadata in ESIMD mode. Co-authored-by: Alexey Bader <alexey.bader@intel.com>
intel · Aug 20, 2020 · f658918 · f658918
1 parent e1e3658
commit f658918
Show file tree

Hide file tree

Showing 2 changed files with 17 additions and 19 deletions.
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -1555,22 +1555,24 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn,
               : llvm::ConstantAsMetadata::get(CGF->Builder.getInt32(-1)));
     }
 
-  Fn->setMetadata("kernel_arg_addr_space",
-                  llvm::MDNode::get(VMContext, addressQuals));
-  Fn->setMetadata("kernel_arg_access_qual",
-                  llvm::MDNode::get(VMContext, accessQuals));
-  Fn->setMetadata("kernel_arg_type",
-                  llvm::MDNode::get(VMContext, argTypeNames));
-  Fn->setMetadata("kernel_arg_base_type",
-                  llvm::MDNode::get(VMContext, argBaseTypeNames));
-  Fn->setMetadata("kernel_arg_type_qual",
-                  llvm::MDNode::get(VMContext, argTypeQuals));
-  if (getCodeGenOpts().EmitOpenCLArgMetadata)
-    Fn->setMetadata("kernel_arg_name",
-                    llvm::MDNode::get(VMContext, argNames));
-  if (LangOpts.SYCLIsDevice)
+  if (LangOpts.SYCLIsDevice && !LangOpts.SYCLExplicitSIMD)
     Fn->setMetadata("kernel_arg_buffer_location",
                     llvm::MDNode::get(VMContext, argSYCLBufferLocationAttr));
+  else {
+    Fn->setMetadata("kernel_arg_addr_space",
+                    llvm::MDNode::get(VMContext, addressQuals));
+    Fn->setMetadata("kernel_arg_access_qual",
+                    llvm::MDNode::get(VMContext, accessQuals));
+    Fn->setMetadata("kernel_arg_type",
+                    llvm::MDNode::get(VMContext, argTypeNames));
+    Fn->setMetadata("kernel_arg_base_type",
+                    llvm::MDNode::get(VMContext, argBaseTypeNames));
+    Fn->setMetadata("kernel_arg_type_qual",
+                    llvm::MDNode::get(VMContext, argTypeQuals));
+    if (getCodeGenOpts().EmitOpenCLArgMetadata)
+      Fn->setMetadata("kernel_arg_name",
+                      llvm::MDNode::get(VMContext, argNames));
+  }
 }
 
 /// Determines whether the language options require us to model

diff --git a/clang/test/CodeGenSYCL/kernel-metadata.cpp b/clang/test/CodeGenSYCL/kernel-metadata.cpp
@@ -1,10 +1,6 @@
 // RUN: %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s
 
-// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE15kernel_function{{.*}} !kernel_arg_addr_space ![[MDAS:[0-9]+]] !kernel_arg_access_qual ![[MDAC:[0-9]+]] !kernel_arg_type ![[MDAT:[0-9]+]] !kernel_arg_base_type ![[MDAT:[0-9]+]] !kernel_arg_type_qual ![[MDATQ:[0-9]+]]
-// CHECK: ![[MDAS]] = !{i32 1, i32 0, i32 0, i32 0}
-// CHECK: ![[MDAC]] = !{!"none", !"none", !"none", !"none"}
-// CHECK: ![[MDAT]] = !{!"int*", !"cl::sycl::range<1>", !"cl::sycl::range<1>", !"cl::sycl::id<1>"}
-// CHECK: ![[MDATQ]] = !{!"", !"", !"", !""}
+// CHECK-NOT: define {{.*}}spir_kernel void @{{.*}}kernel_function{{.*}} !kernel_arg_addr_space {{.*}} !kernel_arg_access_qual {{.*}} !kernel_arg_type {{.*}} !kernel_arg_base_type {{.*}} !kernel_arg_type_qual {{.*}}
 
 #include "sycl.hpp"