diff options
| author | Alexey Sotkin <alexey.sotkin@intel.com> | 2017-07-26 18:49:54 +0000 |
|---|---|---|
| committer | Alexey Sotkin <alexey.sotkin@intel.com> | 2017-07-26 18:49:54 +0000 |
| commit | 7d7f0dc08b5124d1af9818e2841eeb000eaf9624 (patch) | |
| tree | 47446f40ce004f239028bc4e0d21947b0fc47169 /clang | |
| parent | 833ad37c9035baf52fb44e9300e142a756aaedfb (diff) | |
| download | bcm5719-llvm-7d7f0dc08b5124d1af9818e2841eeb000eaf9624.tar.gz bcm5719-llvm-7d7f0dc08b5124d1af9818e2841eeb000eaf9624.zip | |
[OpenCL] Fix access qualifiers metadata for kernel arguments with typedef
Subscribers: cfe-commits, yaxunl, Anastasia
Differential Revision: https://reviews.llvm.org/D35420
llvm-svn: 309155
Diffstat (limited to 'clang')
| -rw-r--r-- | clang/lib/CodeGen/CodeGenFunction.cpp | 5 | ||||
| -rw-r--r-- | clang/test/CodeGenOpenCL/kernel-arg-info.cl | 22 |
2 files changed, 26 insertions, 1 deletions
diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 93a4a386619..fc7a10d08b2 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -620,7 +620,10 @@ static void GenOpenCLArgMetadata(const FunctionDecl *FD, llvm::Function *Fn, // 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()) diff --git a/clang/test/CodeGenOpenCL/kernel-arg-info.cl b/clang/test/CodeGenOpenCL/kernel-arg-info.cl index 463cc445111..fa48ad28f76 100644 --- a/clang/test/CodeGenOpenCL/kernel-arg-info.cl +++ b/clang/test/CodeGenOpenCL/kernel-arg-info.cl @@ -78,6 +78,21 @@ kernel void foo5(myImage img1, write_only image1d_t img2) { 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 @@ __kernel void foo6(__global char16 arg[]) {} // 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"} |
