[PATCH] D35420: [OpenCL] Fix access qualifiers metadata for kernel arguments with typedef

[Alexey Sotkin via Phabricator via cfe-commits]

This revision was automatically updated to reflect the committed changes.
Closed by commit rL309155: [OpenCL] Fix access qualifiers metadata for kernel arguments with typedef (authored by AlexeySotkin).

Repository:
  rL LLVM

https://reviews.llvm.org/D35420

Files:
  cfe/trunk/lib/CodeGen/CodeGenFunction.cpp
  cfe/trunk/test/CodeGenOpenCL/kernel-arg-info.cl


Index: cfe/trunk/lib/CodeGen/CodeGenFunction.cpp
===================================================================
--- cfe/trunk/lib/CodeGen/CodeGenFunction.cpp
+++ cfe/trunk/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())
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
@@ -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"}
 


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D35420.108338.patch
Type: text/x-patch
Size: 2656 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20170726/66fb9ccf/attachment.bin>