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*"}