From 18bc87b3cd59aa78d426d07c62712d6533caad73 Mon Sep 17 00:00:00 2001 From: Premanand M Rao Date: Wed, 19 Aug 2020 19:10:51 -0700 Subject: [PATCH 1/2] [SYCL] Stop emitting kernel arguments metadata 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 --- clang/lib/CodeGen/CodeGenModule.cpp | 28 ++++++++++++---------- clang/test/CodeGenSYCL/kernel-metadata.cpp | 6 +---- 2 files changed, 16 insertions(+), 18 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 3a4626a70d6ed..0744411ba16d8 100644 --- 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) 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 index dd502fa1dd844..80fc954a035f5 100644 --- 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" From a34206193306dea1ca9ebe07f0186f4376bcf393 Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Thu, 20 Aug 2020 12:04:28 +0300 Subject: [PATCH 2/2] [SYCL] Keep OpenCL kernel metadata in ESIMD mode. --- clang/lib/CodeGen/CodeGenModule.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 0744411ba16d8..ad422be30e800 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1555,7 +1555,7 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn, : llvm::ConstantAsMetadata::get(CGF->Builder.getInt32(-1))); } - if (LangOpts.SYCLIsDevice) + if (LangOpts.SYCLIsDevice && !LangOpts.SYCLExplicitSIMD) Fn->setMetadata("kernel_arg_buffer_location", llvm::MDNode::get(VMContext, argSYCLBufferLocationAttr)); else {