From: Sameer Sahasrabuddhe Date: Thu, 4 Dec 2014 05:30:58 +0000 (+0000) Subject: Always emit kernel arg info for SPIR. X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=f514ca8aa824579b46ddecbdebd34315c356505d;p=clang Always emit kernel arg info for SPIR. http://llvm.org/bugs/show_bug.cgi?id=21555 Currently, kernel argument metadata is omitted unless the "-cl-kernel-arg-info" option is specified. But the SPIR 1.2 spec requires that all metadata except kernel_arg_name should always be emitted, and kernel_arg_name is only emitted when "-cl-kernel-arg-info" is specified. Patch ported by Ryan Burn from the Khronos SPIR generator. https://github.com/KhronosGroup/SPIR git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@223340 91177308-0d34-0410-b5e6-96231b3b80d8 --- diff --git a/lib/CodeGen/CodeGenFunction.cpp b/lib/CodeGen/CodeGenFunction.cpp index 4291b14893..02cca4bf21 100644 --- a/lib/CodeGen/CodeGenFunction.cpp +++ b/lib/CodeGen/CodeGenFunction.cpp @@ -483,7 +483,8 @@ static void GenOpenCLArgMetadata(const FunctionDecl *FD, llvm::Function *Fn, kernelMDArgs.push_back(llvm::MDNode::get(Context, argTypeNames)); kernelMDArgs.push_back(llvm::MDNode::get(Context, argBaseTypeNames)); kernelMDArgs.push_back(llvm::MDNode::get(Context, argTypeQuals)); - kernelMDArgs.push_back(llvm::MDNode::get(Context, argNames)); + if (CGM.getCodeGenOpts().EmitOpenCLArgMetadata) + kernelMDArgs.push_back(llvm::MDNode::get(Context, argNames)); } void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, @@ -497,9 +498,8 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, SmallVector kernelMDArgs; kernelMDArgs.push_back(Fn); - if (CGM.getCodeGenOpts().EmitOpenCLArgMetadata) - GenOpenCLArgMetadata(FD, Fn, CGM, Context, kernelMDArgs, - Builder, getContext()); + GenOpenCLArgMetadata(FD, Fn, CGM, Context, kernelMDArgs, Builder, + getContext()); if (const VecTypeHintAttr *A = FD->getAttr()) { QualType hintQTy = A->getTypeHint(); diff --git a/test/CodeGenOpenCL/kernel-arg-info.cl b/test/CodeGenOpenCL/kernel-arg-info.cl index 89da3fd7c4..c853e529d6 100644 --- a/test/CodeGenOpenCL/kernel-arg-info.cl +++ b/test/CodeGenOpenCL/kernel-arg-info.cl @@ -1,4 +1,5 @@ -// RUN: %clang_cc1 %s -cl-kernel-arg-info -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s +// 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 kernel void foo(__global int * restrict X, const int Y, volatile int anotherArg, __constant float * restrict Z) { @@ -10,7 +11,8 @@ kernel void foo(__global int * restrict X, const int Y, // CHECK: metadata !{metadata !"kernel_arg_type", metadata !"int*", metadata !"int", metadata !"int", metadata !"float*"} // CHECK: metadata !{metadata !"kernel_arg_base_type", metadata !"int*", metadata !"int", metadata !"int", metadata !"float*"} // CHECK: metadata !{metadata !"kernel_arg_type_qual", metadata !"restrict", metadata !"const", metadata !"volatile", metadata !"restrict const"} -// CHECK: metadata !{metadata !"kernel_arg_name", metadata !"X", metadata !"Y", metadata !"anotherArg", metadata !"Z"} +// ARGINFO: metadata !{metadata !"kernel_arg_name", metadata !"X", metadata !"Y", metadata !"anotherArg", metadata !"Z"} +// NO-ARGINFO-NOT: metadata !{metadata !"kernel_arg_name", metadata !"X", metadata !"Y", metadata !"anotherArg", metadata !"Z"} kernel void foo2(read_only image1d_t img1, image2d_t img2, write_only image2d_array_t img3) { } @@ -19,7 +21,8 @@ kernel void foo2(read_only image1d_t img1, image2d_t img2, write_only image2d_ar // CHECK: metadata !{metadata !"kernel_arg_type", metadata !"image1d_t", metadata !"image2d_t", metadata !"image2d_array_t"} // CHECK: metadata !{metadata !"kernel_arg_base_type", metadata !"image1d_t", metadata !"image2d_t", metadata !"image2d_array_t"} // CHECK: metadata !{metadata !"kernel_arg_type_qual", metadata !"", metadata !"", metadata !""} -// CHECK: metadata !{metadata !"kernel_arg_name", metadata !"img1", metadata !"img2", metadata !"img3"} +// ARGINFO: metadata !{metadata !"kernel_arg_name", metadata !"img1", metadata !"img2", metadata !"img3"} +// NO-ARGINFO-NOT: metadata !{metadata !"kernel_arg_name", metadata !"img1", metadata !"img2", metadata !"img3"} kernel void foo3(__global half * X) { } @@ -28,7 +31,8 @@ kernel void foo3(__global half * X) { // CHECK: metadata !{metadata !"kernel_arg_type", metadata !"half*"} // CHECK: metadata !{metadata !"kernel_arg_base_type", metadata !"half*"} // CHECK: metadata !{metadata !"kernel_arg_type_qual", metadata !""} -// CHECK: metadata !{metadata !"kernel_arg_name", metadata !"X"} +// ARGINFO: metadata !{metadata !"kernel_arg_name", metadata !"X"} +// NO-ARGINFO-NOT: metadata !{metadata !"kernel_arg_name", metadata !"X"} typedef unsigned int myunsignedint; kernel void foo4(__global unsigned int * X, __global myunsignedint * Y) { @@ -38,7 +42,8 @@ kernel void foo4(__global unsigned int * X, __global myunsignedint * Y) { // CHECK: metadata !{metadata !"kernel_arg_type", metadata !"uint*", metadata !"myunsignedint*"} // CHECK: metadata !{metadata !"kernel_arg_base_type", metadata !"uint*", metadata !"uint*"} // CHECK: metadata !{metadata !"kernel_arg_type_qual", metadata !"", metadata !""} -// CHECK: metadata !{metadata !"kernel_arg_name", metadata !"X", metadata !"Y"} +// ARGINFO: metadata !{metadata !"kernel_arg_name", metadata !"X", metadata !"Y"} +// NO-ARGINFO-NOT: metadata !{metadata !"kernel_arg_name", metadata !"X", metadata !"Y"} typedef image1d_t myImage; kernel void foo5(read_only myImage img1, write_only image1d_t img2) { @@ -46,4 +51,5 @@ kernel void foo5(read_only myImage img1, write_only image1d_t img2) { // CHECK: metadata !{metadata !"kernel_arg_access_qual", metadata !"read_only", metadata !"write_only"} // CHECK: metadata !{metadata !"kernel_arg_type", metadata !"myImage", metadata !"image1d_t"} // CHECK: metadata !{metadata !"kernel_arg_base_type", metadata !"image1d_t", metadata !"image1d_t"} -// CHECK: metadata !{metadata !"kernel_arg_name", metadata !"img1", metadata !"img2"} +// ARGINFO: metadata !{metadata !"kernel_arg_name", metadata !"img1", metadata !"img2"} +// NO-ARGINFO-NOT: metadata !{metadata !"kernel_arg_name", metadata !"img1", metadata !"img2"} diff --git a/test/CodeGenOpenCL/kernel-attributes.cl b/test/CodeGenOpenCL/kernel-attributes.cl index 0825ffc033..fc076ed112 100644 --- a/test/CodeGenOpenCL/kernel-attributes.cl +++ b/test/CodeGenOpenCL/kernel-attributes.cl @@ -8,9 +8,9 @@ kernel __attribute__((vec_type_hint(uint4))) __attribute__((work_group_size_hint // CHECK: opencl.kernels = !{[[MDNODE0:![0-9]+]], [[MDNODE3:![0-9]+]]} -// CHECK: [[MDNODE0]] = metadata !{void (i32)* @kernel1, metadata [[MDNODE1:![0-9]+]], metadata [[MDNODE2:![0-9]+]]} +// CHECK: [[MDNODE0]] = metadata !{void (i32)* @kernel1, {{.*}} metadata [[MDNODE1:![0-9]+]], metadata [[MDNODE2:![0-9]+]]} // CHECK: [[MDNODE1]] = metadata !{metadata !"vec_type_hint", i32 undef, i32 1} // CHECK: [[MDNODE2]] = metadata !{metadata !"reqd_work_group_size", i32 1, i32 2, i32 4} -// CHECK: [[MDNODE3]] = metadata !{void (i32)* @kernel2, metadata [[MDNODE4:![0-9]+]], metadata [[MDNODE5:![0-9]+]]} +// CHECK: [[MDNODE3]] = metadata !{void (i32)* @kernel2, {{.*}} metadata [[MDNODE4:![0-9]+]], metadata [[MDNODE5:![0-9]+]]} // CHECK: [[MDNODE4]] = metadata !{metadata !"vec_type_hint", <4 x i32> undef, i32 0} // CHECK: [[MDNODE5]] = metadata !{metadata !"work_group_size_hint", i32 8, i32 16, i32 32} diff --git a/test/CodeGenOpenCL/kernel-metadata.cl b/test/CodeGenOpenCL/kernel-metadata.cl index 3e10a119d0..e3a23d4822 100644 --- a/test/CodeGenOpenCL/kernel-metadata.cl +++ b/test/CodeGenOpenCL/kernel-metadata.cl @@ -7,4 +7,9 @@ __kernel void kernel_function() { } // CHECK: !opencl.kernels = !{!0} -// CHECK: !0 = metadata !{void ()* @kernel_function} +// CHECK: !0 = metadata !{void ()* @kernel_function, metadata !1, metadata !2, metadata !3, metadata !4, metadata !5} +// CHECK: !1 = metadata !{metadata !"kernel_arg_addr_space"} +// CHECK: !2 = metadata !{metadata !"kernel_arg_access_qual"} +// CHECK: !3 = metadata !{metadata !"kernel_arg_type"} +// CHECK: !4 = metadata !{metadata !"kernel_arg_base_type"} +// CHECK: !5 = metadata !{metadata !"kernel_arg_type_qual"}