Index: cfe/trunk/lib/CodeGen/CodeGenFunction.cpp =================================================================== --- cfe/trunk/lib/CodeGen/CodeGenFunction.cpp +++ cfe/trunk/lib/CodeGen/CodeGenFunction.cpp @@ -435,7 +435,6 @@ // 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 &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 @@ -445,28 +444,21 @@ // MDNode for the kernel argument address space qualifiers. SmallVector addressQuals; - addressQuals.push_back(llvm::MDString::get(Context, "kernel_arg_addr_space")); // MDNode for the kernel argument access qualifiers (images only). SmallVector accessQuals; - accessQuals.push_back(llvm::MDString::get(Context, "kernel_arg_access_qual")); // MDNode for the kernel argument type names. SmallVector argTypeNames; - argTypeNames.push_back(llvm::MDString::get(Context, "kernel_arg_type")); // MDNode for the kernel argument base type names. SmallVector argBaseTypeNames; - argBaseTypeNames.push_back( - llvm::MDString::get(Context, "kernel_arg_base_type")); // MDNode for the kernel argument type qualifiers. SmallVector argTypeQuals; - argTypeQuals.push_back(llvm::MDString::get(Context, "kernel_arg_type_qual")); // MDNode for the kernel argument names. SmallVector 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); @@ -577,13 +569,19 @@ 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, @@ -594,11 +592,7 @@ llvm::LLVMContext &Context = getLLVMContext(); - SmallVector 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()) { QualType hintQTy = A->getTypeHint(); @@ -607,37 +601,29 @@ 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()) { 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()) { 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. Index: cfe/trunk/test/CodeGenOpenCL/kernel-arg-info.cl =================================================================== --- cfe/trunk/test/CodeGenOpenCL/kernel-arg-info.cl +++ cfe/trunk/test/CodeGenOpenCL/kernel-arg-info.cl @@ -1,55 +1,88 @@ -// 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"} + Index: cfe/trunk/test/CodeGenOpenCL/kernel-attributes.cl =================================================================== --- cfe/trunk/test/CodeGenOpenCL/kernel-attributes.cl +++ cfe/trunk/test/CodeGenOpenCL/kernel-attributes.cl @@ -3,14 +3,12 @@ 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} Index: cfe/trunk/test/CodeGenOpenCL/kernel-metadata.cl =================================================================== --- cfe/trunk/test/CodeGenOpenCL/kernel-metadata.cl +++ cfe/trunk/test/CodeGenOpenCL/kernel-metadata.cl @@ -6,10 +6,5 @@ __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]] = !{}