[OpenCL] Added extended tests on metadata generation for half data type and arrays.

[Anastasia]

Btw, is there any reason to add testing specifically for half? Is there anything specific to half in the implementation of this?

test/CodeGenOpenCL/kernel-arg-info.cl

79

could you add space b/w half and *. Also could we format consistently below?

[Anastasia]

In D35000#799705, @Anastasia wrote:

Btw, is there any reason to add testing specifically for half? Is there anything specific to half in the implementation of this?

Trying to understand the reason for this change though...

[echuraev]

In D35000#801132, @Anastasia wrote:

In D35000#799705, @Anastasia wrote:

Btw, is there any reason to add testing specifically for half? Is there anything specific to half in the implementation of this?

Trying to understand the reason for this change though...

Sorry for a delay in response. No it is not any reason to add testing specifically for half. We can also do the same tests for other data types. Here we just check that it is no any qualifiers in metadata.

[Anastasia]

In D35000#807674, @echuraev wrote:

In D35000#801132, @Anastasia wrote:

In D35000#799705, @Anastasia wrote:

Btw, is there any reason to add testing specifically for half? Is there anything specific to half in the implementation of this?

Trying to understand the reason for this change though...

Sorry for a delay in response. No it is not any reason to add testing specifically for half. We can also do the same tests for other data types. Here we just check that it is no any qualifiers in metadata.

Btw, if I read the spec then it feels like we should put the qualifiers of the pointee type instead:

CL_KERNEL_ARG_TYPE_VOLATILE is returned if the argument is a pointer and the referenced type is declared with the volatile qualifier. For example, a kernel argument declared as global int volatile *x returns CL_KERNEL_ARG_TYPE_VOLATILE but a kernel argument declared as global int *volatile x does not. Similarly, CL_KERNEL_ARG_TYPE_RESTRICT or CL_KERNEL_ARG_TYPE_CONST is returned if the argument is a pointer and the referenced type is declared with the restrict or const qualifier. For example, a kernel argument declared as global int const *x returns CL_KERNEL_ARG_TYPE_CONST but a kernel argument declared as global int *const x does not.

It seems that the function foo is attempting to test that too. Perhaps we could unify the testing and create all combinations of pointer and non pointer types which would have qualifier and not.

I don't feel that we should test all possible types though...

[Anastasia]

LGTM! Thanks!

test/CodeGenOpenCL/kernel-arg-info.cl

 // CHECK: !kernel_arg_addr_space ![[MD41:[0-9]+]]
 // CHECK: !kernel_arg_access_qual ![[MD51:[0-9]+]]
 // CHECK: !kernel_arg_type ![[MD52:[0-9]+]]
 // 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]+]]
 
+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*"}
 // CHECK: ![[MD14]] = !{!"restrict", !"", !"", !"restrict const", !"volatile", !"const"}
 // ARGINFO: ![[MD15]] = !{!"X", !"Y", !"anotherArg", !"Z", !"V", !"C"}
 // CHECK: ![[MD21]] = !{i32 1, i32 1, i32 1, i32 1}
 // CHECK: ![[MD22]] = !{!"read_only", !"read_only", !"write_only", !"read_write"}
 // CHECK: ![[MD23]] = !{!"image1d_t", !"image2d_t", !"image2d_array_t", !"image1d_t"}

 // CHECK: ![[MD43]] = !{!"uint*", !"myunsignedint*"}
 // 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]] = !{!"half*", !"half*", !"half*", !"half*", !"half*", !"half*", !"half*", !"half*", !"half*"}
+// CHECK: ![[MD71]] = !{!"char16*"}