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

[bader]

LGTM. Thanks.

[Anastasia]

LGTM! Thanks! Variable name could be fixed before committing.

lib/CodeGen/CodeGenFunction.cpp

624 ↗

(On Diff #106646)

Even though... not consistent here... variable names should start with upper case.

[AlexeySotkin]

Changing case in the variable name

[AlexeySotkin]

Rebasing on tip of trank

cfe/trunk/lib/CodeGen/CodeGenFunction.cpp

if (ty->isPointerType()) {
       if (isPipe)
         typeQuals = "pipe";
     }
 
     argTypeQuals.push_back(llvm::MDString::get(Context, typeQuals));
 
     // 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())
         accessQuals.push_back(llvm::MDString::get(Context, "read_write"));
       else
         accessQuals.push_back(llvm::MDString::get(Context, "read_only"));
     } else
       accessQuals.push_back(llvm::MDString::get(Context, "none"));

cfe/trunk/test/CodeGenOpenCL/kernel-arg-info.cl

 // 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]+]]
 
 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"}
 // CHECK: ![[MD13]] = !{!"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int", !"int", !"int", !"int"}
 // CHECK: ![[MD14]] = !{!"", !"restrict", !"const", !"restrict const", !"const", !"restrict const", !"const volatile", !"restrict const volatile", !"volatile", !"restrict volatile", !"", !"restrict", !"const", !"restrict const", !"const volatile", !"restrict const volatile", !"volatile", !"restrict volatile", !"", !"", !"", !""}
 // ARGINFO: ![[MD15]] = !{!"globalintp", !"globalintrestrictp", !"globalconstintp", !"globalconstintrestrictp", !"constantintp", !"constantintrestrictp", !"globalconstvolatileintp", !"globalconstvolatileintrestrictp", !"globalvolatileintp", !"globalvolatileintrestrictp", !"localintp", !"localintrestrictp", !"localconstintp", !"localconstintrestrictp", !"localconstvolatileintp", !"localconstvolatileintrestrictp", !"localvolatileintp", !"localvolatileintrestrictp", !"X", !"constint", !"constvolatileint", !"volatileint"}
 // CHECK: ![[MD21]] = !{i32 1, i32 1, i32 1, i32 1}
 // CHECK: ![[MD22]] = !{!"read_only", !"read_only", !"write_only", !"read_write"}

 // CHECK: ![[MD44]] = !{!"uint*", !"uint*"}
 // CHECK: ![[MD45]] = !{!"", !""}
 // ARGINFO: ![[MD46]] = !{!"X", !"Y"}
 // CHECK: ![[MD51]] = !{!"read_only", !"write_only"}
 // CHECK: ![[MD52]] = !{!"myImage", !"image1d_t"}
 // 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"}