[PR41008][OpenCL] Add OpenCL C compatibility mode to C++ for OpenCL

[kpet]

Summary

... and support the restrict keyword in that mode.

In this mode, compatibility with OpenCL C is favoured over
compatibility with C++. Intended uses include migrating
existing OpenCL C code and mixing OpenCL C and C++ for OpenCL
code.

Signed-off-by: Kevin Petit <[e-mail]>

[lebedev.ri]

Is there a test that shows that that keyword is still not provided for C++ after this?

[kpet]

There doesn't seem to be a test for this (enabling restrict in C++ doesn't make any existing test fail). I couldn't find any test checking that OpenCL or C++/OpenCL keywords generally aren't enabled in C or C++. Doing this generically would essentially equate to writing unit tests for getKeywordStatus.

EDIT: I have verified that restrict is indeed rejected in C++.

[Anastasia]

Can you please explain why we are deviating from C++ on this? My concern is that this can potentially uncover bugs in C++ parsing due to un-handled restrict cases that would have to be fixed in some unclear way... I would appreciate more detailed analysis before we go ahead.

[kpet]

Something I should probably have explained in the commit message is that the functionality provided by restrict is already present in C++, where it is already exposed using different keywords. If you look at clang/include/clang/Basic/TokenKinds.def, there are __restrict__ and __restrict aliases that end up producing tok::kw_restrict. The only thing this patch does is enabling the restrict keyword in C++ for OpenCL for compatibility with OpenCL C (see more on that on the bug). I don't expect any special handling will be required since the functionality is already supported both in C++ and OpenCL C.

[kpet]

As discussed offline, let's only enable restrict under a compatibility flag. PTAL.

[Anastasia]

I think we should use -fpermissive rather than adding similar flag to Clang. At the end we might end up with other cases where we need similar mechanism it doesn't make sense to add a flag for each case.

Also I am not sure that what you are doing here is really enabling OpenCL C compatibility mode that would imply to me that restrict is expected to behave in C++ mode as it was in C but is it really the case?

[kpet]

Thanks for the feedback.

I think we should use -fpermissive rather than adding similar flag to Clang.

Could you point me at a use of -fpermissive for a similar case? Also it seems -fpermissive is not mentioned in the Clang manual or the help text. Am I missing something?

At the end we might end up with other cases where we need similar mechanism it doesn't make sense to add a flag for each case.

Completely agree and that was one of the reasons I decided to add a generic flag like this. I felt that compatibility with existing OpenCL C source was an important enough use-case to warrant adding a flag for it especially since, as you say, we expect such a flag could be useful to control several bits of behaviour.

Also I am not sure that what you are doing here is really enabling OpenCL C compatibility mode that would imply to me that restrict is expected to behave in C++ mode as it was in C but is it really the case?

Isn't restrict mainly an optimisation feature? That's certainly my reading of 6.7.3-7 in the C99 spec:

The intended use of the restrict qualifier (like the register storage class) is to promote
optimization, and deleting all instances of the qualifier from all preprocessing translation
units composing a conforming program does not change its meaning (i.e., observable
behavior).

The main goal here is to be able to consume existing OpenCL C source without modifications. I thought that enabling existing behaviour for restrict/__restrict that may result in additional optimisations was more useful than just discarding restrict, which would be a valid way of achieving compatibility. Is there something specific that worries you?

[Anastasia]

In D68388#1756701, @kpet wrote:

Thanks for the feedback.

I think we should use -fpermissive rather than adding similar flag to Clang.

Could you point me at a use of -fpermissive for a similar case? Also it seems -fpermissive is not mentioned in the Clang manual or the help text. Am I missing something?

bin/clang -fpermissive test.cl

Clang already has the flag added. So you can just start using it I presume. Feel free to document it.

At the end we might end up with other cases where we need similar mechanism it doesn't make sense to add a flag for each case.

Completely agree and that was one of the reasons I decided to add a generic flag like this. I felt that compatibility with existing OpenCL C source was an important enough use-case to warrant adding a flag for it especially since, as you say, we expect such a flag could be useful to control several bits of behaviour.

Also I am not sure that what you are doing here is really enabling OpenCL C compatibility mode that would imply to me that restrict is expected to behave in C++ mode as it was in C but is it really the case?

Isn't restrict mainly an optimisation feature? That's certainly my reading of 6.7.3-7 in the C99 spec:

The intended use of the restrict qualifier (like the register storage class) is to promote
optimization, and deleting all instances of the qualifier from all preprocessing translation
units composing a conforming program does not change its meaning (i.e., observable
behavior).

The main goal here is to be able to consume existing OpenCL C source without modifications. I thought that enabling existing behaviour for restrict/__restrict that may result in additional optimisations was more useful than just discarding restrict, which would be a valid way of achieving compatibility. Is there something specific that worries you?

C++ deviates from C so restrict in C++ might not be restrict in C. Can we confidently say we will support restrict exactly like OpenCL C does i.e. C99?

In other words I believe it is much safe to say that we support restrict in adhoc way that deviates C++ for OpenCL documentation without implying developers should refer to OpenCL C and C99 spec to get how it works. We can also use this mechanism for any behavior that deviates the documentation rather than adding a flag for each (i.e. we might want to deviate in something that doesn't come from OpenCL C at all).

clang/docs/ClangCommandLineReference.rst

 .. option:: -cl-uniform-work-group-size
 
 OpenCL only. Defines that the global work-size be a multiple of the work-group size specified to clEnqueueNDRangeKernel
 
 .. option:: -cl-unsafe-math-optimizations
 
 OpenCL only. Allow unsafe floating-point optimizations.  Also implies -cl-no-signed-zeros and -cl-mad-enable.
 
+C++ for OpenCL flags
+--------------------
+
+.. option:: -clcxx-clc-compatibility
+
+C++ for OpenCL only. Accept OpenCL C constructs that would otherwise be undefined in C++ for OpenCL. Favours compatibility with OpenCL C rather than with C++.
+
 Target-dependent compilation options
 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
 .. option:: -G<size>, -G=<arg>, -msmall-data-threshold=<arg>
 
 Put objects of at most <size> bytes into small data section (MIPS / Hexagon)
 
 .. option:: -m16
 

clang/include/clang/Basic/LangOptions.def

 ENUM_LANGOPT(MSPointerToMemberRepresentationMethod, PragmaMSPointersToMembersKind, 2, PPTMK_BestCase, "member-pointer representation method")
 ENUM_LANGOPT(DefaultCallingConv, DefaultCallingConvention, 3, DCC_None, "default calling convention")
 
 LANGOPT(ShortEnums        , 1, 0, "short enum types")
 
 LANGOPT(OpenCL            , 1, 0, "OpenCL")
 LANGOPT(OpenCLVersion     , 32, 0, "OpenCL C version")
 LANGOPT(OpenCLCPlusPlus   , 1, 0, "C++ for OpenCL")
+LANGOPT(OpenCLCPlusPlusCLCCompat , 1, 0, "C++ for OpenCL - OpenCL C compatibility mode")
 LANGOPT(OpenCLCPlusPlusVersion     , 32, 0, "C++ for OpenCL version")
 LANGOPT(NativeHalfType    , 1, 0, "Native half type support")
 LANGOPT(NativeHalfArgsAndReturns, 1, 0, "Native half args and returns")
 LANGOPT(HalfArgsAndReturns, 1, 0, "half args and returns")
 LANGOPT(CUDA              , 1, 0, "CUDA")
 LANGOPT(HIP               , 1, 0, "HIP")
 LANGOPT(OpenMP            , 32, 0, "OpenMP support and version of OpenMP (31, 40 or 45)")
 LANGOPT(OpenMPSimd        , 1, 0, "Use SIMD only OpenMP support.")

clang/include/clang/Basic/TokenKinds.def

 KEYWORD(float                       , KEYALL)
 KEYWORD(for                         , KEYALL)
 KEYWORD(goto                        , KEYALL)
 KEYWORD(if                          , KEYALL)
 KEYWORD(inline                      , KEYC99|KEYCXX|KEYGNU)
 KEYWORD(int                         , KEYALL)
 KEYWORD(long                        , KEYALL)
 KEYWORD(register                    , KEYALL)
-KEYWORD(restrict                    , KEYC99)
+KEYWORD(restrict                    , KEYC99|KEYOPENCLCXXCOMPAT)
 KEYWORD(return                      , KEYALL)
 KEYWORD(short                       , KEYALL)
 KEYWORD(signed                      , KEYALL)
 KEYWORD(sizeof                      , KEYALL)
 KEYWORD(static                      , KEYALL)
 KEYWORD(struct                      , KEYALL)
 KEYWORD(switch                      , KEYALL)
 KEYWORD(typedef                     , KEYALL)

clang/include/clang/Driver/Options.td

 
 def f_clang_Group : OptionGroup<"<f (clang-only) group>">,
                     Group<CompileOnly_Group>, DocFlatten;
 def pedantic_Group : OptionGroup<"<pedantic group>">, Group<f_Group>,
                      DocFlatten;
 def opencl_Group : OptionGroup<"<opencl group>">, Group<f_Group>,
                    DocName<"OpenCL flags">;
 
+def openclcxx_Group : OptionGroup<"<c++ opencl group>">, Group<f_Group>,
+                   DocName<"C++ for OpenCL flags">;
+
 def m_Group : OptionGroup<"<m group>">, Group<CompileOnly_Group>,
               DocName<"Target-dependent compilation options">;
 
 // Feature groups - these take command line options that correspond directly to
 // target specific features and can be translated directly from command line
 // options.
 def m_aarch64_Features_Group : OptionGroup<"<aarch64 features group>">,
                                Group<m_Group>, DocName<"AARCH64">;

 def cl_std_EQ : Joined<["-"], "cl-std=">, Group<opencl_Group>, Flags<[CC1Option]>,
   HelpText<"OpenCL language standard to compile for.">, Values<"cl,CL,cl1.1,CL1.1,cl1.2,CL1.2,cl2.0,CL2.0,clc++,CLC++">;
 def cl_denorms_are_zero : Flag<["-"], "cl-denorms-are-zero">, Group<opencl_Group>, Flags<[CC1Option]>,
   HelpText<"OpenCL only. Allow denormals to be flushed to zero.">;
 def cl_fp32_correctly_rounded_divide_sqrt : Flag<["-"], "cl-fp32-correctly-rounded-divide-sqrt">, Group<opencl_Group>, Flags<[CC1Option]>,
   HelpText<"OpenCL only. Specify that single precision floating-point divide and sqrt used in the program source are correctly rounded.">;
 def cl_uniform_work_group_size : Flag<["-"], "cl-uniform-work-group-size">, Group<opencl_Group>, Flags<[CC1Option]>,
   HelpText<"OpenCL only. Defines that the global work-size be a multiple of the work-group size specified to clEnqueueNDRangeKernel">;
+def clcxx_clc_compatibility : Flag<["-"], "clcxx-clc-compatibility">, Group<openclcxx_Group>, Flags<[CC1Option]>,
+  HelpText<"C++ for OpenCL only. Accept OpenCL C constructs that would otherwise be undefined in C++ for OpenCL. Favours compatibility with OpenCL C rather than with C++.">;
 def client__name : JoinedOrSeparate<["-"], "client_name">;
 def combine : Flag<["-", "--"], "combine">, Flags<[DriverOption, Unsupported]>;
 def compatibility__version : JoinedOrSeparate<["-"], "compatibility_version">;
 def config : Separate<["--"], "config">, Flags<[DriverOption]>,
   HelpText<"Specifies configuration file">;
 def config_system_dir_EQ : Joined<["--"], "config-system-dir=">, Flags<[DriverOption, HelpHidden]>,
   HelpText<"System directory for configuration files">;
 def config_user_dir_EQ : Joined<["--"], "config-user-dir=">, Flags<[DriverOption, HelpHidden]>,

clang/lib/Basic/IdentifierTable.cpp

 //===----------------------------------------------------------------------===//
 // Language Keyword Implementation
 //===----------------------------------------------------------------------===//
 
 // Constants for TokenKinds.def
 namespace {
 
   enum {
-    KEYC99        = 0x1,
-    KEYCXX        = 0x2,
-    KEYCXX11      = 0x4,
-    KEYGNU        = 0x8,
-    KEYMS         = 0x10,
-    BOOLSUPPORT   = 0x20,
-    KEYALTIVEC    = 0x40,
-    KEYNOCXX      = 0x80,
-    KEYBORLAND    = 0x100,
-    KEYOPENCLC    = 0x200,
-    KEYC11        = 0x400,
-    KEYNOMS18     = 0x800,
-    KEYNOOPENCL   = 0x1000,
-    WCHARSUPPORT  = 0x2000,
-    HALFSUPPORT   = 0x4000,
-    CHAR8SUPPORT  = 0x8000,
-    KEYCONCEPTS   = 0x10000,
-    KEYOBJC       = 0x20000,
-    KEYZVECTOR    = 0x40000,
-    KEYCOROUTINES = 0x80000,
-    KEYMODULES    = 0x100000,
-    KEYCXX2A      = 0x200000,
-    KEYOPENCLCXX  = 0x400000,
-    KEYMSCOMPAT   = 0x800000,
+    KEYC99             = 0x1,
+    KEYCXX             = 0x2,
+    KEYCXX11           = 0x4,
+    KEYGNU             = 0x8,
+    KEYMS              = 0x10,
+    BOOLSUPPORT        = 0x20,
+    KEYALTIVEC         = 0x40,
+    KEYNOCXX           = 0x80,
+    KEYBORLAND         = 0x100,
+    KEYOPENCLC         = 0x200,
+    KEYC11             = 0x400,
+    KEYNOMS18          = 0x800,
+    KEYNOOPENCL        = 0x1000,
+    WCHARSUPPORT       = 0x2000,
+    HALFSUPPORT        = 0x4000,
+    CHAR8SUPPORT       = 0x8000,
+    KEYCONCEPTS        = 0x10000,
+    KEYOBJC            = 0x20000,
+    KEYZVECTOR         = 0x40000,
+    KEYCOROUTINES      = 0x80000,
+    KEYMODULES         = 0x100000,
+    KEYCXX2A           = 0x200000,
+    KEYOPENCLCXX       = 0x400000,
+    KEYMSCOMPAT        = 0x800000,
+    KEYOPENCLCXXCOMPAT = 0x1000000,
     KEYALLCXX = KEYCXX | KEYCXX11 | KEYCXX2A,
     KEYALL = (0xffffff & ~KEYNOMS18 &
               ~KEYNOOPENCL) // KEYNOMS18 and KEYNOOPENCL are used to exclude.
   };
 
   /// How a keyword is treated in the selected standard.
   enum KeywordStatus {
     KS_Disabled,    // Disabled
static KeywordStatus getKeywordStatus(const LangOptions &LangOpts,
   if (LangOpts.Half && (Flags & HALFSUPPORT)) return KS_Enabled;
   if (LangOpts.WChar && (Flags & WCHARSUPPORT)) return KS_Enabled;
   if (LangOpts.Char8 && (Flags & CHAR8SUPPORT)) return KS_Enabled;
   if (LangOpts.AltiVec && (Flags & KEYALTIVEC)) return KS_Enabled;
   if (LangOpts.ZVector && (Flags & KEYZVECTOR)) return KS_Enabled;
   if (LangOpts.OpenCL && !LangOpts.OpenCLCPlusPlus && (Flags & KEYOPENCLC))
     return KS_Enabled;
   if (LangOpts.OpenCLCPlusPlus && (Flags & KEYOPENCLCXX)) return KS_Enabled;
+  if (LangOpts.OpenCLCPlusPlusCLCCompat && (Flags & KEYOPENCLCXXCOMPAT))
+    return KS_Extension;
   if (!LangOpts.CPlusPlus && (Flags & KEYNOCXX)) return KS_Enabled;
   if (LangOpts.C11 && (Flags & KEYC11)) return KS_Enabled;
   // We treat bridge casts as objective-C keywords so we can warn on them
   // in non-arc mode.
   if (LangOpts.ObjC && (Flags & KEYOBJC)) return KS_Enabled;
   if (LangOpts.ConceptsTS && (Flags & KEYCONCEPTS)) return KS_Enabled;
   if (LangOpts.Coroutines && (Flags & KEYCOROUTINES)) return KS_Enabled;
   if (LangOpts.ModulesTS && (Flags & KEYMODULES)) return KS_Enabled;

clang/lib/Driver/ToolChains/Clang.cpp

const unsigned ForwardedArguments[] = {
       options::OPT_cl_finite_math_only,
       options::OPT_cl_kernel_arg_info,
       options::OPT_cl_unsafe_math_optimizations,
       options::OPT_cl_fast_relaxed_math,
       options::OPT_cl_mad_enable,
       options::OPT_cl_no_signed_zeros,
       options::OPT_cl_denorms_are_zero,
       options::OPT_cl_fp32_correctly_rounded_divide_sqrt,
-      options::OPT_cl_uniform_work_group_size
+      options::OPT_cl_uniform_work_group_size,
+      options::OPT_clcxx_clc_compatibility
   };
 
   if (Arg *A = Args.getLastArg(options::OPT_cl_std_EQ)) {
     std::string CLStdStr = std::string("-cl-std=") + A->getValue();
     CmdArgs.push_back(Args.MakeArgString(CLStdStr));
   }
 
   for (const auto &Arg : ForwardedArguments)

clang/lib/Frontend/CompilerInvocation.cpp

#include "clang/Basic/LangStandards.def"
   // this option was added for compatibility with OpenCL 1.0.
   if (Args.getLastArg(OPT_cl_strict_aliasing)
        && Opts.OpenCLVersion > 100) {
     Diags.Report(diag::warn_option_invalid_ocl_version)
         << Opts.getOpenCLVersionTuple().getAsString()
         << Args.getLastArg(OPT_cl_strict_aliasing)->getAsString(Args);
   }
 
+  if (Args.getLastArg(OPT_clcxx_clc_compatibility)) {
+    Opts.OpenCLCPlusPlusCLCCompat = true;
+  }
+
   // We abuse '-f[no-]gnu-keywords' to force overriding all GNU-extension
   // keywords. This behavior is provided by GCC's poorly named '-fasm' flag,
   // while a subset (the non-C++ GNU keywords) is provided by GCC's
   // '-fgnu-keywords'. Clang conflates the two for simplicity under the single
   // name, as it doesn't seem a useful distinction.
   Opts.GNUKeywords = Args.hasFlag(OPT_fgnu_keywords, OPT_fno_gnu_keywords,
                                   Opts.GNUKeywords);
 

clang/test/CodeGenOpenCLCXX/compat-clc-kernel-arg-info.cl

+// RUN: %clang_cc1 %s -cl-std=clc++ -clcxx-clc-compatibility -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=clc++ -clcxx-clc-compatibility -emit-llvm -o - -triple spir-unknown-unknown -cl-kernel-arg-info | FileCheck %s -check-prefix ARGINFO
+
+kernel void foo(global int * globalintp, global int * restrict globalintrestrictp,
+                global const int * globalconstintp,
+                global const int * restrict globalconstintrestrictp,
+                constant int * constantintp, constant int * restrict constantintrestrictp,
+                global const volatile int * globalconstvolatileintp,
+                global const volatile int * restrict globalconstvolatileintrestrictp,
+                global volatile int * globalvolatileintp,
+                global volatile int * restrict globalvolatileintrestrictp,
+                local int * localintp, local int * restrict localintrestrictp,
+                local const int * localconstintp,
+                local const int * restrict localconstintrestrictp,
+                local const volatile int * localconstvolatileintp,
+                local const volatile int * restrict localconstvolatileintrestrictp,
+                local volatile int * localvolatileintp,
+                local volatile int * restrict localvolatileintrestrictp,
+                int X, const int constint, const volatile int constvolatileint,
+                volatile int volatileint) {
+  *globalintrestrictp = constint + volatileint;
+}
+// CHECK: define spir_kernel void @foo{{[^!]+}}
+// CHECK: !kernel_arg_addr_space ![[MD11:[0-9]+]]
+// CHECK: !kernel_arg_access_qual ![[MD12:[0-9]+]]
+// CHECK: !kernel_arg_type ![[MD13:[0-9]+]]
+// CHECK: !kernel_arg_base_type ![[MD13]]
+// CHECK: !kernel_arg_type_qual ![[MD14:[0-9]+]]
+// CHECK-NOT: !kernel_arg_name
+// ARGINFO: !kernel_arg_name ![[MD15:[0-9]+]]
+
+kernel void foo2(read_only image1d_t img1, image2d_t img2, write_only image2d_array_t img3, read_write image1d_t img4) {
+}
+// CHECK: define spir_kernel void @foo2{{[^!]+}}
+// CHECK: !kernel_arg_addr_space ![[MD21:[0-9]+]]
+// CHECK: !kernel_arg_access_qual ![[MD22:[0-9]+]]
+// CHECK: !kernel_arg_type ![[MD23:[0-9]+]]
+// CHECK: !kernel_arg_base_type ![[MD23]]
+// CHECK: !kernel_arg_type_qual ![[MD24:[0-9]+]]
+// CHECK-NOT: !kernel_arg_name
+// ARGINFO: !kernel_arg_name ![[MD25:[0-9]+]]
+
+kernel void foo3(__global half * X) {
+}
+// CHECK: define spir_kernel void @foo3{{[^!]+}}
+// CHECK: !kernel_arg_addr_space ![[MD31:[0-9]+]]
+// CHECK: !kernel_arg_access_qual ![[MD32:[0-9]+]]
+// CHECK: !kernel_arg_type ![[MD33:[0-9]+]]
+// CHECK: !kernel_arg_base_type ![[MD33]]
+// CHECK: !kernel_arg_type_qual ![[MD34:[0-9]+]]
+// CHECK-NOT: !kernel_arg_name
+// ARGINFO: !kernel_arg_name ![[MD35:[0-9]+]]
+
+typedef unsigned int myunsignedint;
+kernel void foo4(__global unsigned int * X, __global myunsignedint * Y) {
+}
+// CHECK: define spir_kernel void @foo4{{[^!]+}}
+// CHECK: !kernel_arg_addr_space ![[MD41:[0-9]+]]
+// CHECK: !kernel_arg_access_qual ![[MD42:[0-9]+]]
+// CHECK: !kernel_arg_type ![[MD43:[0-9]+]]
+// CHECK: !kernel_arg_base_type ![[MD44:[0-9]+]]
+// CHECK: !kernel_arg_type_qual ![[MD45:[0-9]+]]
+// CHECK-NOT: !kernel_arg_name
+// ARGINFO: !kernel_arg_name ![[MD46:[0-9]+]]
+
+typedef image1d_t myImage;
+kernel void foo5(myImage img1, write_only image1d_t img2) {
+}
+// CHECK: define spir_kernel void @foo5{{[^!]+}}
+// 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]+]]
+
+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: ![[MD23]] = !{!"image1d_t", !"image2d_t", !"image2d_array_t", !"image1d_t"}
+// CHECK: ![[MD24]] = !{!"", !"", !"", !""}
+// ARGINFO: ![[MD25]] = !{!"img1", !"img2", !"img3", !"img4"}
+// CHECK: ![[MD31]] = !{i32 1}
+// CHECK: ![[MD32]] = !{!"none"}
+// CHECK: ![[MD33]] = !{!"half*"}
+// CHECK: ![[MD34]] = !{!""}
+// ARGINFO: ![[MD35]] = !{!"X"}
+// CHECK: ![[MD41]] = !{i32 1, i32 1}
+// CHECK: ![[MD42]] = !{!"none", !"none"}
+// 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]] = !{!"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"}
+