r223340 - Always emit kernel arg info for SPIR.
Sameer Sahasrabuddhe
sameer.sahasrabuddhe at amd.com
Wed Dec 3 21:30:58 PST 2014
Author: sameerds
Date: Wed Dec 3 23:30:58 2014
New Revision: 223340
URL: http://llvm.org/viewvc/llvm-project?rev=223340&view=rev
Log:
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
Modified:
cfe/trunk/lib/CodeGen/CodeGenFunction.cpp
cfe/trunk/test/CodeGenOpenCL/kernel-arg-info.cl
cfe/trunk/test/CodeGenOpenCL/kernel-attributes.cl
cfe/trunk/test/CodeGenOpenCL/kernel-metadata.cl
Modified: cfe/trunk/lib/CodeGen/CodeGenFunction.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenFunction.cpp?rev=223340&r1=223339&r2=223340&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CodeGenFunction.cpp (original)
+++ cfe/trunk/lib/CodeGen/CodeGenFunction.cpp Wed Dec 3 23:30:58 2014
@@ -483,7 +483,8 @@ static void GenOpenCLArgMetadata(const F
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::EmitOpenCLKernelMe
SmallVector <llvm::Value*, 5> 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<VecTypeHintAttr>()) {
QualType hintQTy = A->getTypeHint();
Modified: cfe/trunk/test/CodeGenOpenCL/kernel-arg-info.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/kernel-arg-info.cl?rev=223340&r1=223339&r2=223340&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/kernel-arg-info.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/kernel-arg-info.cl Wed Dec 3 23:30:58 2014
@@ -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
// 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 img
// 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 *
// 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,
// 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"}
Modified: cfe/trunk/test/CodeGenOpenCL/kernel-attributes.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/kernel-attributes.cl?rev=223340&r1=223339&r2=223340&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/kernel-attributes.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/kernel-attributes.cl Wed Dec 3 23:30:58 2014
@@ -8,9 +8,9 @@ kernel __attribute__((vec_type_hint(uint
// 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}
Modified: cfe/trunk/test/CodeGenOpenCL/kernel-metadata.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/kernel-metadata.cl?rev=223340&r1=223339&r2=223340&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/kernel-metadata.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/kernel-metadata.cl Wed Dec 3 23:30:58 2014
@@ -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"}
More information about the cfe-commits
mailing list