echuraev created this revision. Herald added a subscriber: yaxunl. https://reviews.llvm.org/D35000
Files: 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 @@ -61,6 +61,37 @@ // CHECK-NOT: !kernel_arg_name // ARGINFO: !kernel_arg_name ![[MD54:[0-9]+]] +kernel void foo6(constant half*constanthalfp, + constant half *restrict constanthalfrestrictp, + global half*globalhalfp, + global half *restrict globalhalfrestrictp, + global const half* globalconsthalfp, + global const half * restrict globalconsthalfrestrictp, + + global volatile half*globalvolatilehalfp, + global volatile half *restrict globalvolatilehalfrestrictp, + global const volatile half* globalconstvolatilehalfp) +{} +// CHECK: !kernel_arg_type ![[MD61:[0-9]+]] +// CHECK: !kernel_arg_base_type ![[MD61]] + +kernel void foo6_2(global const volatile half * restrict globalconstvolatilehalfrestrictp, + local half*localhalfp, + local half *restrict localhalfrestrictp, + local const half* localconsthalfp, + local const half * restrict localconsthalfrestrictp, + local volatile half*localvolatilehalfp, + local volatile half *restrict localvolatilehalfrestrictp, + local const volatile half* localconstvolatilehalfp, + local const volatile half * restrict localconstvolatilehalfrestrictp) +{} +// CHECK: !kernel_arg_type ![[MD61]] +// CHECK: !kernel_arg_base_type ![[MD61]] + +typedef char char16 __attribute__((ext_vector_type(16))); +__kernel void foo7(__global char16 arg[]) {} +// CHECK: !kernel_arg_type ![[MD71:[0-9]+]] + // CHECK: ![[MD11]] = !{i32 1, i32 0, i32 0, i32 2, i32 1, i32 1} // CHECK: ![[MD12]] = !{!"none", !"none", !"none", !"none", !"none", !"none"} // CHECK: ![[MD13]] = !{!"int*", !"int", !"int", !"float*", !"int*", !"int*"} @@ -86,4 +117,6 @@ // CHECK: ![[MD52]] = !{!"myImage", !"image1d_t"} // CHECK: ![[MD53]] = !{!"image1d_t", !"image1d_t"} // ARGINFO: ![[MD54]] = !{!"img1", !"img2"} +// CHECK: ![[MD61]] = !{!"half*", !"half*", !"half*", !"half*", !"half*", !"half*", !"half*", !"half*", !"half*"} +// CHECK: ![[MD71]] = !{!"char16*"}
Index: test/CodeGenOpenCL/kernel-arg-info.cl =================================================================== --- test/CodeGenOpenCL/kernel-arg-info.cl +++ test/CodeGenOpenCL/kernel-arg-info.cl @@ -61,6 +61,37 @@ // CHECK-NOT: !kernel_arg_name // ARGINFO: !kernel_arg_name ![[MD54:[0-9]+]] +kernel void foo6(constant half*constanthalfp, + constant half *restrict constanthalfrestrictp, + global half*globalhalfp, + global half *restrict globalhalfrestrictp, + global const half* globalconsthalfp, + global const half * restrict globalconsthalfrestrictp, + + global volatile half*globalvolatilehalfp, + global volatile half *restrict globalvolatilehalfrestrictp, + global const volatile half* globalconstvolatilehalfp) +{} +// CHECK: !kernel_arg_type ![[MD61:[0-9]+]] +// CHECK: !kernel_arg_base_type ![[MD61]] + +kernel void foo6_2(global const volatile half * restrict globalconstvolatilehalfrestrictp, + local half*localhalfp, + local half *restrict localhalfrestrictp, + local const half* localconsthalfp, + local const half * restrict localconsthalfrestrictp, + local volatile half*localvolatilehalfp, + local volatile half *restrict localvolatilehalfrestrictp, + local const volatile half* localconstvolatilehalfp, + local const volatile half * restrict localconstvolatilehalfrestrictp) +{} +// CHECK: !kernel_arg_type ![[MD61]] +// CHECK: !kernel_arg_base_type ![[MD61]] + +typedef char char16 __attribute__((ext_vector_type(16))); +__kernel void foo7(__global char16 arg[]) {} +// CHECK: !kernel_arg_type ![[MD71:[0-9]+]] + // CHECK: ![[MD11]] = !{i32 1, i32 0, i32 0, i32 2, i32 1, i32 1} // CHECK: ![[MD12]] = !{!"none", !"none", !"none", !"none", !"none", !"none"} // CHECK: ![[MD13]] = !{!"int*", !"int", !"int", !"float*", !"int*", !"int*"} @@ -86,4 +117,6 @@ // CHECK: ![[MD52]] = !{!"myImage", !"image1d_t"} // CHECK: ![[MD53]] = !{!"image1d_t", !"image1d_t"} // ARGINFO: ![[MD54]] = !{!"img1", !"img2"} +// CHECK: ![[MD61]] = !{!"half*", !"half*", !"half*", !"half*", !"half*", !"half*", !"half*", !"half*", !"half*"} +// CHECK: ![[MD71]] = !{!"char16*"}
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits