[OpenCL] Use function metadata to represent kernel attributes

[yaxunl]

Summary

This patch uses function metadata to represent reqd_work_group_size, work_group_size_hint and vector_type_hint kernel attributes and kernel argument info.

[Anastasia]

Looking good generally, I am just not sure about mixing two different representations.

[yaxunl]

In D20979#452463, @Anastasia wrote:

Looking good generally, I am just not sure about mixing two different representations.

If we choose only one form of representation, would you suggest to use function metadata or function attribute?

[Anastasia]

In D20979#452616, @yaxunl wrote:

In D20979#452463, @Anastasia wrote:

Looking good generally, I am just not sure about mixing two different representations.

If we choose only one form of representation, would you suggest to use function metadata or function attribute?

I am still not sure if this is the intended use of target-dependent attributes to be honest. So I would prefer metadata representation.

Also if we use metadata could we avoid parsing values from strings potentially in contrast to attributes that represent all values as strings?

Related to your earlier comments about inflexibility of metadata, would it be possible to extend MDNode to be able to insert new operands?

[yaxunl]

In D20979#453632, @Anastasia wrote:

In D20979#452616, @yaxunl wrote:

In D20979#452463, @Anastasia wrote:

Looking good generally, I am just not sure about mixing two different representations.

If we choose only one form of representation, would you suggest to use function metadata or function attribute?

I am still not sure if this is the intended use of target-dependent attributes to be honest. So I would prefer metadata representation.

Also if we use metadata could we avoid parsing values from strings potentially in contrast to attributes that represent all values as strings?

OK due to the controversy of target-dependent function attributes, let's concur to use function metadata for all the kernel attributes, and do not convert everything to string in the metadata.

Related to your earlier comments about inflexibility of metadata, would it be possible to extend MDNode to be able to insert new operands?

I will see how difficult is that.

[yaxunl]

I updated the patch to use function metadata only. Please review. Thanks.

[Anastasia]

LGTM!

[yaxunl]

Ping!

Alexey/Xiuli, are you OK with this change? Thanks.

[bader]

LGTM.

test/CodeGenOpenCL/kernel-attributes.cl

9 ↗

(On Diff #60545)

AFAIK, this named metadata node were also useful for kernel functions traversal.
What is the new BKM to differentiate kernel/non-kernel functions?
Check for !kernel_arg_* function attribute?

[yaxunl]

test/CodeGenOpenCL/kernel-attributes.cl

9 ↗

(On Diff #60545)

I think it could be used for that purpose, since kernel functions and only kernel functions have this function attribute.

cfe/trunk/lib/CodeGen/CodeGenFunction.cpp

void CodeGenFunction::EmitMCountInstrumentation() {
   EmitNounwindRuntimeCall(MCountFn);
 }
 
 // OpenCL v1.2 s5.6.4.6 allows the compiler to store kernel argument
 // information in the program executable. The argument information stored
 // includes the argument name, its type, the address and access qualifiers used.
 static void GenOpenCLArgMetadata(const FunctionDecl *FD, llvm::Function *Fn,
                                  CodeGenModule &CGM, llvm::LLVMContext &Context,
-                                 SmallVector<llvm::Metadata *, 5> &kernelMDArgs,
                                  CGBuilderTy &Builder, ASTContext &ASTCtx) {
   // Create MDNodes that represent the kernel arg metadata.
   // Each MDNode is a list in the form of "key", N number of values which is
   // the same number of values as their are kernel arguments.
 
   const PrintingPolicy &Policy = ASTCtx.getPrintingPolicy();
 
   // MDNode for the kernel argument address space qualifiers.
   SmallVector<llvm::Metadata *, 8> addressQuals;
-  addressQuals.push_back(llvm::MDString::get(Context, "kernel_arg_addr_space"));
 
   // MDNode for the kernel argument access qualifiers (images only).
   SmallVector<llvm::Metadata *, 8> accessQuals;
-  accessQuals.push_back(llvm::MDString::get(Context, "kernel_arg_access_qual"));
 
   // MDNode for the kernel argument type names.
   SmallVector<llvm::Metadata *, 8> argTypeNames;
-  argTypeNames.push_back(llvm::MDString::get(Context, "kernel_arg_type"));
 
   // MDNode for the kernel argument base type names.
   SmallVector<llvm::Metadata *, 8> argBaseTypeNames;
-  argBaseTypeNames.push_back(
-      llvm::MDString::get(Context, "kernel_arg_base_type"));
 
   // MDNode for the kernel argument type qualifiers.
   SmallVector<llvm::Metadata *, 8> argTypeQuals;
-  argTypeQuals.push_back(llvm::MDString::get(Context, "kernel_arg_type_qual"));
 
   // MDNode for the kernel argument names.
   SmallVector<llvm::Metadata *, 8> argNames;
-  argNames.push_back(llvm::MDString::get(Context, "kernel_arg_name"));
 
   for (unsigned i = 0, e = FD->getNumParams(); i != e; ++i) {
     const ParmVarDecl *parm = FD->getParamDecl(i);
     QualType ty = parm->getType();
     std::string typeQuals;
 
     if (ty->isPointerType()) {
       QualType pointeeTy = ty->getPointeeType();
if (ty->isImageType()|| ty->isPipeType()) {
         accessQuals.push_back(llvm::MDString::get(Context, "read_only"));
     } else
       accessQuals.push_back(llvm::MDString::get(Context, "none"));
 
     // Get argument name.
     argNames.push_back(llvm::MDString::get(Context, parm->getName()));
   }
 
-  kernelMDArgs.push_back(llvm::MDNode::get(Context, addressQuals));
-  kernelMDArgs.push_back(llvm::MDNode::get(Context, accessQuals));
-  kernelMDArgs.push_back(llvm::MDNode::get(Context, argTypeNames));
-  kernelMDArgs.push_back(llvm::MDNode::get(Context, argBaseTypeNames));
-  kernelMDArgs.push_back(llvm::MDNode::get(Context, argTypeQuals));
+  Fn->setMetadata("kernel_arg_addr_space",
+                  llvm::MDNode::get(Context, addressQuals));
+  Fn->setMetadata("kernel_arg_access_qual",
+                  llvm::MDNode::get(Context, accessQuals));
+  Fn->setMetadata("kernel_arg_type",
+                  llvm::MDNode::get(Context, argTypeNames));
+  Fn->setMetadata("kernel_arg_base_type",
+                  llvm::MDNode::get(Context, argBaseTypeNames));
+  Fn->setMetadata("kernel_arg_type_qual",
+                  llvm::MDNode::get(Context, argTypeQuals));
   if (CGM.getCodeGenOpts().EmitOpenCLArgMetadata)
-    kernelMDArgs.push_back(llvm::MDNode::get(Context, argNames));
+    Fn->setMetadata("kernel_arg_name",
+                    llvm::MDNode::get(Context, argNames));
 }
 
 void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
                                                llvm::Function *Fn)
 {
   if (!FD->hasAttr<OpenCLKernelAttr>())
     return;
 
   llvm::LLVMContext &Context = getLLVMContext();
 
-  SmallVector<llvm::Metadata *, 5> kernelMDArgs;
-  kernelMDArgs.push_back(llvm::ConstantAsMetadata::get(Fn));
-
-  GenOpenCLArgMetadata(FD, Fn, CGM, Context, kernelMDArgs, Builder,
-                       getContext());
+  GenOpenCLArgMetadata(FD, Fn, CGM, Context, Builder, getContext());
 
   if (const VecTypeHintAttr *A = FD->getAttr<VecTypeHintAttr>()) {
     QualType hintQTy = A->getTypeHint();
     const ExtVectorType *hintEltQTy = hintQTy->getAs<ExtVectorType>();
     bool isSignedInteger =
         hintQTy->isSignedIntegerType() ||
         (hintEltQTy && hintEltQTy->getElementType()->isSignedIntegerType());
     llvm::Metadata *attrMDArgs[] = {
-        llvm::MDString::get(Context, "vec_type_hint"),
         llvm::ConstantAsMetadata::get(llvm::UndefValue::get(
             CGM.getTypes().ConvertType(A->getTypeHint()))),
         llvm::ConstantAsMetadata::get(llvm::ConstantInt::get(
             llvm::IntegerType::get(Context, 32),
             llvm::APInt(32, (uint64_t)(isSignedInteger ? 1 : 0))))};
-    kernelMDArgs.push_back(llvm::MDNode::get(Context, attrMDArgs));
+    Fn->setMetadata("vec_type_hint", llvm::MDNode::get(Context, attrMDArgs));
   }
 
   if (const WorkGroupSizeHintAttr *A = FD->getAttr<WorkGroupSizeHintAttr>()) {
     llvm::Metadata *attrMDArgs[] = {
-        llvm::MDString::get(Context, "work_group_size_hint"),
         llvm::ConstantAsMetadata::get(Builder.getInt32(A->getXDim())),
         llvm::ConstantAsMetadata::get(Builder.getInt32(A->getYDim())),
         llvm::ConstantAsMetadata::get(Builder.getInt32(A->getZDim()))};
-    kernelMDArgs.push_back(llvm::MDNode::get(Context, attrMDArgs));
+    Fn->setMetadata("work_group_size_hint", llvm::MDNode::get(Context, attrMDArgs));
   }
 
   if (const ReqdWorkGroupSizeAttr *A = FD->getAttr<ReqdWorkGroupSizeAttr>()) {
     llvm::Metadata *attrMDArgs[] = {
-        llvm::MDString::get(Context, "reqd_work_group_size"),
         llvm::ConstantAsMetadata::get(Builder.getInt32(A->getXDim())),
         llvm::ConstantAsMetadata::get(Builder.getInt32(A->getYDim())),
         llvm::ConstantAsMetadata::get(Builder.getInt32(A->getZDim()))};
-    kernelMDArgs.push_back(llvm::MDNode::get(Context, attrMDArgs));
+    Fn->setMetadata("reqd_work_group_size", llvm::MDNode::get(Context, attrMDArgs));
   }
-
-  llvm::MDNode *kernelMDNode = llvm::MDNode::get(Context, kernelMDArgs);
-  llvm::NamedMDNode *OpenCLKernelMetadata =
-    CGM.getModule().getOrInsertNamedMetadata("opencl.kernels");
-  OpenCLKernelMetadata->addOperand(kernelMDNode);
 }
 
 /// Determine whether the function F ends with a return stmt.
 static bool endsWithReturn(const Decl* F) {
   const Stmt *Body = nullptr;
   if (auto *FD = dyn_cast_or_null<FunctionDecl>(F))
     Body = FD->getBody();
   else if (auto *OMD = dyn_cast_or_null<ObjCMethodDecl>(F))

cfe/trunk/test/CodeGenOpenCL/kernel-arg-info.cl

-// RUN: %clang_cc1 %s -cl-kernel-arg-info -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s -check-prefix ARGINFO
-// RUN: %clang_cc1 %s -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s -check-prefix NO-ARGINFO
+// RUN: %clang_cc1 %s -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s
+// RUN: %clang_cc1 %s -emit-llvm -o - -triple spir-unknown-unknown -cl-kernel-arg-info | FileCheck %s -check-prefix ARGINFO
 
 kernel void foo(__global int * restrict X, const int Y, 
                 volatile int anotherArg, __constant float * restrict Z) {
   *X = Y + anotherArg;
 }
-
-// CHECK:  !{!"kernel_arg_addr_space", i32 1, i32 0, i32 0, i32 2}
-// CHECK:  !{!"kernel_arg_access_qual", !"none", !"none", !"none", !"none"}
-// CHECK:  !{!"kernel_arg_type", !"int*", !"int", !"int", !"float*"}
-// CHECK:  !{!"kernel_arg_base_type", !"int*", !"int", !"int", !"float*"}
-// CHECK:  !{!"kernel_arg_type_qual", !"restrict", !"const", !"volatile", !"restrict const"}
-// ARGINFO: !{!"kernel_arg_name", !"X", !"Y", !"anotherArg", !"Z"}
-// NO-ARGINFO-NOT: !{!"kernel_arg_name", !"X", !"Y", !"anotherArg", !"Z"}
+// CHECK: define spir_kernel void @foo{{[^!]+}}
+// CHECK: !kernel_arg_addr_space ![[MD11:[0-9]+]]
+// CHECK: !kernel_arg_access_qual ![[MD12:[0-9]+]]
+// CHECK: !kernel_arg_type ![[MD13:[0-9]+]]
+// CHECK: !kernel_arg_base_type ![[MD13]]
+// CHECK: !kernel_arg_type_qual ![[MD14:[0-9]+]]
+// CHECK-NOT: !kernel_arg_name
+// ARGINFO: !kernel_arg_name ![[MD15:[0-9]+]]
 
 kernel void foo2(read_only image1d_t img1, image2d_t img2, write_only image2d_array_t img3) {
 }
-// CHECK:  !{!"kernel_arg_addr_space", i32 1, i32 1, i32 1}
-// CHECK:  !{!"kernel_arg_access_qual", !"read_only", !"read_only", !"write_only"}
-// CHECK:  !{!"kernel_arg_type", !"image1d_t", !"image2d_t", !"image2d_array_t"}
-// CHECK:  !{!"kernel_arg_base_type", !"image1d_t", !"image2d_t", !"image2d_array_t"}
-// CHECK:  !{!"kernel_arg_type_qual", !"", !"", !""}
-// ARGINFO: !{!"kernel_arg_name", !"img1", !"img2", !"img3"}
-// NO-ARGINFO-NOT: !{!"kernel_arg_name", !"img1", !"img2", !"img3"}
+// CHECK: define spir_kernel void @foo2{{[^!]+}}
+// CHECK: !kernel_arg_addr_space ![[MD21:[0-9]+]]
+// CHECK: !kernel_arg_access_qual ![[MD22:[0-9]+]]
+// CHECK: !kernel_arg_type ![[MD23:[0-9]+]]
+// CHECK: !kernel_arg_base_type ![[MD23]]
+// CHECK: !kernel_arg_type_qual ![[MD24:[0-9]+]]
+// CHECK-NOT: !kernel_arg_name
+// ARGINFO: !kernel_arg_name ![[MD25:[0-9]+]]
 
 kernel void foo3(__global half * X) {
 }
-// CHECK:  !{!"kernel_arg_addr_space", i32 1}
-// CHECK:  !{!"kernel_arg_access_qual", !"none"}
-// CHECK:  !{!"kernel_arg_type", !"half*"}
-// CHECK:  !{!"kernel_arg_base_type", !"half*"}
-// CHECK:  !{!"kernel_arg_type_qual", !""}
-// ARGINFO: !{!"kernel_arg_name", !"X"}
-// NO-ARGINFO-NOT: !{!"kernel_arg_name", !"X"}
+// CHECK: define spir_kernel void @foo3{{[^!]+}}
+// CHECK: !kernel_arg_addr_space ![[MD31:[0-9]+]]
+// CHECK: !kernel_arg_access_qual ![[MD32:[0-9]+]]
+// CHECK: !kernel_arg_type ![[MD33:[0-9]+]]
+// CHECK: !kernel_arg_base_type ![[MD33]]
+// CHECK: !kernel_arg_type_qual ![[MD34:[0-9]+]]
+// CHECK-NOT: !kernel_arg_name
+// ARGINFO: !kernel_arg_name ![[MD35:[0-9]+]]
 
 typedef unsigned int myunsignedint;
 kernel void foo4(__global unsigned int * X, __global myunsignedint * Y) {
 }
-// CHECK:  !{!"kernel_arg_addr_space", i32 1, i32 1}
-// CHECK:  !{!"kernel_arg_access_qual", !"none", !"none"}
-// CHECK:  !{!"kernel_arg_type", !"uint*", !"myunsignedint*"}
-// CHECK:  !{!"kernel_arg_base_type", !"uint*", !"uint*"}
-// CHECK:  !{!"kernel_arg_type_qual", !"", !""}
-// ARGINFO: !{!"kernel_arg_name", !"X", !"Y"}
-// NO-ARGINFO-NOT: !{!"kernel_arg_name", !"X", !"Y"}
+// CHECK: define spir_kernel void @foo4{{[^!]+}}
+// CHECK: !kernel_arg_addr_space ![[MD41:[0-9]+]]
+// CHECK: !kernel_arg_access_qual ![[MD42:[0-9]+]]
+// CHECK: !kernel_arg_type ![[MD43:[0-9]+]]
+// CHECK: !kernel_arg_base_type ![[MD44:[0-9]+]]
+// CHECK: !kernel_arg_type_qual ![[MD45:[0-9]+]]
+// CHECK-NOT: !kernel_arg_name
+// ARGINFO: !kernel_arg_name ![[MD46:[0-9]+]]
 
 typedef image1d_t myImage;
 kernel void foo5(read_only myImage img1, write_only image1d_t img2) {
 }
-// CHECK:  !{!"kernel_arg_access_qual", !"read_only", !"write_only"}
-// CHECK:  !{!"kernel_arg_type", !"myImage", !"image1d_t"}
-// CHECK:  !{!"kernel_arg_base_type", !"image1d_t", !"image1d_t"}
-// ARGINFO: !{!"kernel_arg_name", !"img1", !"img2"}
-// NO-ARGINFO-NOT: !{!"kernel_arg_name", !"img1", !"img2"}
+// CHECK: define spir_kernel void @foo5{{[^!]+}}
+// CHECK: !kernel_arg_addr_space ![[MD41:[0-9]+]]
+// CHECK: !kernel_arg_access_qual ![[MD51:[0-9]+]]
+// CHECK: !kernel_arg_type ![[MD52:[0-9]+]]
+// CHECK: !kernel_arg_base_type ![[MD53:[0-9]+]]
+// CHECK: !kernel_arg_type_qual ![[MD45]]
+// CHECK-NOT: !kernel_arg_name
+// ARGINFO: !kernel_arg_name ![[MD54:[0-9]+]]
+
+// CHECK: ![[MD11]] = !{i32 1, i32 0, i32 0, i32 2}
+// CHECK: ![[MD12]] = !{!"none", !"none", !"none", !"none"}
+// CHECK: ![[MD13]] = !{!"int*", !"int", !"int", !"float*"}
+// CHECK: ![[MD14]] = !{!"restrict", !"const", !"volatile", !"restrict const"}
+// ARGINFO: ![[MD15]] = !{!"X", !"Y", !"anotherArg", !"Z"}
+// CHECK: ![[MD21]] = !{i32 1, i32 1, i32 1}
+// CHECK: ![[MD22]] = !{!"read_only", !"read_only", !"write_only"}
+// CHECK: ![[MD23]] = !{!"__read_only image1d_t", !"__read_only image2d_t", !"__write_only image2d_array_t"}
+// CHECK: ![[MD24]] = !{!"", !"", !""}
+// ARGINFO: ![[MD25]] = !{!"img1", !"img2", !"img3"}
+// CHECK: ![[MD31]] = !{i32 1}
+// CHECK: ![[MD32]] = !{!"none"}
+// CHECK: ![[MD33]] = !{!"half*"}
+// CHECK: ![[MD34]] = !{!""}
+// ARGINFO: ![[MD35]] = !{!"X"}
+// CHECK: ![[MD41]] = !{i32 1, i32 1}
+// CHECK: ![[MD42]] = !{!"none", !"none"}
+// CHECK: ![[MD43]] = !{!"uint*", !"myunsignedint*"}
+// CHECK: ![[MD44]] = !{!"uint*", !"uint*"}
+// CHECK: ![[MD45]] = !{!"", !""}
+// ARGINFO: ![[MD46]] = !{!"X", !"Y"}
+// CHECK: ![[MD51]] = !{!"read_only", !"write_only"}
+// CHECK: ![[MD52]] = !{!"myImage", !"__write_only image1d_t"}
+// CHECK: ![[MD53]] = !{!"__read_only image1d_t", !"__write_only image1d_t"}
+// ARGINFO: ![[MD54]] = !{!"img1", !"img2"}
+

cfe/trunk/test/CodeGenOpenCL/kernel-attributes.cl

 // RUN: %clang_cc1 -emit-llvm -o - %s | FileCheck %s
 
 typedef unsigned int uint4 __attribute__((ext_vector_type(4)));
 
 kernel  __attribute__((vec_type_hint(int))) __attribute__((reqd_work_group_size(1,2,4))) void kernel1(int a) {}
+// CHECK: define void @kernel1(i32 %a) {{[^{]+}} !vec_type_hint ![[MD1:[0-9]+]] !reqd_work_group_size ![[MD2:[0-9]+]]
 
 kernel __attribute__((vec_type_hint(uint4))) __attribute__((work_group_size_hint(8,16,32))) void kernel2(int a) {}
+// CHECK: define void @kernel2(i32 %a) {{[^{]+}} !vec_type_hint ![[MD3:[0-9]+]] !work_group_size_hint ![[MD4:[0-9]+]]
 
-// CHECK: opencl.kernels = !{[[MDNODE0:![0-9]+]], [[MDNODE3:![0-9]+]]}
-
-// CHECK: [[MDNODE0]] = !{void (i32)* @kernel1, {{.*}} [[MDNODE1:![0-9]+]], [[MDNODE2:![0-9]+]]}
-// CHECK: [[MDNODE1]] = !{!"vec_type_hint", i32 undef, i32 1}
-// CHECK: [[MDNODE2]] = !{!"reqd_work_group_size", i32 1, i32 2, i32 4}
-// CHECK: [[MDNODE3]] = !{void (i32)* @kernel2, {{.*}} [[MDNODE4:![0-9]+]], [[MDNODE5:![0-9]+]]}
-// CHECK: [[MDNODE4]] = !{!"vec_type_hint", <4 x i32> undef, i32 0}
-// CHECK: [[MDNODE5]] = !{!"work_group_size_hint", i32 8, i32 16, i32 32}
+// CHECK: [[MD1]] = !{i32 undef, i32 1}
+// CHECK: [[MD2]] = !{i32 1, i32 2, i32 4}
+// CHECK: [[MD3]] = !{<4 x i32> undef, i32 0}
+// CHECK: [[MD4]] = !{i32 8, i32 16, i32 32}

cfe/trunk/test/CodeGenOpenCL/kernel-metadata.cl

 // RUN: %clang_cc1 %s -emit-llvm -o - | FileCheck %s
 
 void normal_function() {
 }
 
 __kernel void kernel_function() {
 }
 
-// CHECK: !opencl.kernels = !{!0}
-// CHECK: !0 = !{void ()* @kernel_function, !1, !2, !3, !4, !5}
-// CHECK: !1 = !{!"kernel_arg_addr_space"}
-// CHECK: !2 = !{!"kernel_arg_access_qual"}
-// CHECK: !3 = !{!"kernel_arg_type"}
-// CHECK: !4 = !{!"kernel_arg_base_type"}
-// CHECK: !5 = !{!"kernel_arg_type_qual"}
+// CHECK: define void @kernel_function() {{[^{]+}} !kernel_arg_addr_space ![[MD:[0-9]+]] !kernel_arg_access_qual ![[MD]] !kernel_arg_type ![[MD]] !kernel_arg_base_type ![[MD]] !kernel_arg_type_qual ![[MD]] {
+// CHECK: ![[MD]] = !{}