[PATCH] D35420: [OpenCL] Fix access qualifiers metadata for kernel arguments with typedef
Alexey Sotkin via Phabricator via cfe-commits
cfe-commits at lists.llvm.org
Wed Jul 26 11:32:48 PDT 2017
AlexeySotkin updated this revision to Diff 108334.
AlexeySotkin added a comment.
Rebasing on tip of trank
https://reviews.llvm.org/D35420
Files:
lib/CodeGen/CodeGenFunction.cpp
test/CodeGenOpenCL/kernel-arg-info.cl
Index: test/CodeGenOpenCL/kernel-arg-info.cl
===================================================================
--- test/CodeGenOpenCL/kernel-arg-info.cl
+++ test/CodeGenOpenCL/kernel-arg-info.cl
@@ -78,6 +78,21 @@
typedef char char16 __attribute__((ext_vector_type(16)));
__kernel void foo6(__global char16 arg[]) {}
// CHECK: !kernel_arg_type ![[MD61:[0-9]+]]
+// ARGINFO: !kernel_arg_name ![[MD62:[0-9]+]]
+
+typedef read_only image1d_t ROImage;
+typedef write_only image1d_t WOImage;
+typedef read_write image1d_t RWImage;
+kernel void foo7(ROImage ro, WOImage wo, RWImage rw) {
+}
+// CHECK: define spir_kernel void @foo7{{[^!]+}}
+// CHECK: !kernel_arg_addr_space ![[MD71:[0-9]+]]
+// CHECK: !kernel_arg_access_qual ![[MD72:[0-9]+]]
+// CHECK: !kernel_arg_type ![[MD73:[0-9]+]]
+// CHECK: !kernel_arg_base_type ![[MD74:[0-9]+]]
+// CHECK: !kernel_arg_type_qual ![[MD75:[0-9]+]]
+// CHECK-NOT: !kernel_arg_name
+// ARGINFO: !kernel_arg_name ![[MD76:[0-9]+]]
// CHECK: ![[MD11]] = !{i32 1, i32 1, i32 1, i32 1, i32 2, i32 2, i32 1, i32 1, i32 1, i32 1, i32 3, i32 3, i32 3, i32 3, i32 3, i32 3, i32 3, i32 3, i32 0, i32 0, i32 0, i32 0}
// CHECK: ![[MD12]] = !{!"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none"}
@@ -105,4 +120,11 @@
// CHECK: ![[MD53]] = !{!"image1d_t", !"image1d_t"}
// ARGINFO: ![[MD54]] = !{!"img1", !"img2"}
// CHECK: ![[MD61]] = !{!"char16*"}
+// ARGINFO: ![[MD62]] = !{!"arg"}
+// CHECK: ![[MD71]] = !{i32 1, i32 1, i32 1}
+// CHECK: ![[MD72]] = !{!"read_only", !"write_only", !"read_write"}
+// CHECK: ![[MD73]] = !{!"ROImage", !"WOImage", !"RWImage"}
+// CHECK: ![[MD74]] = !{!"image1d_t", !"image1d_t", !"image1d_t"}
+// CHECK: ![[MD75]] = !{!"", !"", !""}
+// ARGINFO: ![[MD76]] = !{!"ro", !"wo", !"rw"}
Index: lib/CodeGen/CodeGenFunction.cpp
===================================================================
--- lib/CodeGen/CodeGenFunction.cpp
+++ lib/CodeGen/CodeGenFunction.cpp
@@ -620,7 +620,10 @@
// Get image and pipe access qualifier:
if (ty->isImageType()|| ty->isPipeType()) {
- const OpenCLAccessAttr *A = parm->getAttr<OpenCLAccessAttr>();
+ const Decl *PDecl = parm;
+ if (auto *TD = dyn_cast<TypedefType>(ty))
+ PDecl = TD->getDecl();
+ const OpenCLAccessAttr *A = PDecl->getAttr<OpenCLAccessAttr>();
if (A && A->isWriteOnly())
accessQuals.push_back(llvm::MDString::get(Context, "write_only"));
else if (A && A->isReadWrite())
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D35420.108334.patch
Type: text/x-patch
Size: 2596 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20170726/7c2c80c1/attachment-0001.bin>
More information about the cfe-commits
mailing list