[OpenCL] Do not guard vload/store_half builtins

[svenvh]

Summary

The vload*_half* and vstore*_half* builtins do not require the
cl_khr_fp16 extension: pointers to half can be declared without the
extension and the _half variants of vload and vstore should be
available without the extension.

This aligns the guards for these builtins for
-fdeclare-opencl-builtins with opencl-c.h.

Fixes https://github.com/llvm/llvm-project/issues/55275

[svenvh]

Add test case.

[Anastasia]

LGTM! Thanks

clang/lib/Headers/opencl-c-base.h

 #endif
 typedef double double2 __attribute__((ext_vector_type(2)));
 typedef double double3 __attribute__((ext_vector_type(3)));
 typedef double double4 __attribute__((ext_vector_type(4)));
 typedef double double8 __attribute__((ext_vector_type(8)));
 typedef double double16 __attribute__((ext_vector_type(16)));
 #endif
 
+// An internal alias for half, for use by OpenCLBuiltins.td.
+#define __half half
+
 #if defined(__OPENCL_CPP_VERSION__)
 #define NULL nullptr
 #elif defined(__OPENCL_C_VERSION__)
 #define NULL ((void*)0)
 #endif
 
 /**
  * Value of maximum non-infinite single-precision floating-point

clang/lib/Sema/OpenCLBuiltins.td

 def Int       : Type<"int",       QualType<"Context.IntTy">>;
 def UInt      : Type<"uint",      QualType<"Context.UnsignedIntTy">>;
 def Long      : Type<"long",      QualType<"Context.LongTy">>;
 def ULong     : Type<"ulong",     QualType<"Context.UnsignedLongTy">>;
 def Float     : Type<"float",     QualType<"Context.FloatTy">>;
 let Extension = Fp64TypeExt in {
   def Double    : Type<"double",    QualType<"Context.DoubleTy">>;
 }
+
+// The half type for builtins that require the cl_khr_fp16 extension.
 let Extension = Fp16TypeExt in {
   def Half      : Type<"half",      QualType<"Context.HalfTy">>;
 }
+
+// Without the cl_khr_fp16 extension, the half type can only be used to declare
+// a pointer.  Define const and non-const pointer types in all address spaces.
+// Use the "__half" alias to allow the TableGen emitter to distinguish the
+// (extensionless) pointee type of these pointer-to-half types from the "half"
+// type defined above that already carries the cl_khr_fp16 extension.
+foreach AS = [PrivateAS, GlobalAS, ConstantAS, LocalAS, GenericAS] in {
+  def "HalfPtr" # AS      : PointerType<Type<"__half", QualType<"Context.HalfTy">>, AS>;
+  def "HalfPtrConst" # AS : PointerType<ConstType<Type<"__half", QualType<"Context.HalfTy">>>, AS>;
+}
+
 def Size      : Type<"size_t",    QualType<"Context.getSizeType()">>;
 def PtrDiff   : Type<"ptrdiff_t", QualType<"Context.getPointerDiffType()">>;
 def IntPtr    : Type<"intptr_t",  QualType<"Context.getIntPtrType()">>;
 def UIntPtr   : Type<"uintptr_t", QualType<"Context.getUIntPtrType()">>;
 def Void      : Type<"void",      QualType<"Context.VoidTy">>;
 
 // OpenCL v1.0/1.2/2.0 s6.1.2: Built-in Vector Data Types.
 // Built-in vector data types are created by TableGen's OpenCLBuiltinEmitter.

 let Extension = FuncExtOpenCLCGenericAddressSpace in {
   defm : VloadVstore<[GenericAS], 1>;
 }
 // vload with constant address space is available regardless of version.
 defm : VloadVstore<[ConstantAS], 0>;
 
 multiclass VloadVstoreHalf<list<AddressSpace> addrspaces, bit defStores> {
   foreach AS = addrspaces in {
-    def : Builtin<"vload_half", [Float, Size, PointerType<ConstType<Half>, AS>], Attr.Pure>;
+    def : Builtin<"vload_half", [Float, Size, !cast<Type>("HalfPtrConst" # AS)], Attr.Pure>;
     foreach VSize = [2, 3, 4, 8, 16] in {
       foreach name = ["vload_half" # VSize, "vloada_half" # VSize] in {
-        def : Builtin<name, [VectorType<Float, VSize>, Size, PointerType<ConstType<Half>, AS>], Attr.Pure>;
+        def : Builtin<name, [VectorType<Float, VSize>, Size, !cast<Type>("HalfPtrConst" # AS)], Attr.Pure>;
       }
     }
     if defStores then {
       foreach rnd = ["", "_rte", "_rtz", "_rtp", "_rtn"] in {
         foreach name = ["vstore_half" # rnd] in {
-          def : Builtin<name, [Void, Float, Size, PointerType<Half, AS>]>;
-          def : Builtin<name, [Void, Double, Size, PointerType<Half, AS>]>;
+          def : Builtin<name, [Void, Float, Size, !cast<Type>("HalfPtr" # AS)]>;
+          def : Builtin<name, [Void, Double, Size, !cast<Type>("HalfPtr" # AS)]>;
         }
         foreach VSize = [2, 3, 4, 8, 16] in {
           foreach name = ["vstore_half" # VSize # rnd, "vstorea_half" # VSize # rnd] in {
-            def : Builtin<name, [Void, VectorType<Float, VSize>, Size, PointerType<Half, AS>]>;
-            def : Builtin<name, [Void, VectorType<Double, VSize>, Size, PointerType<Half, AS>]>;
+            def : Builtin<name, [Void, VectorType<Float, VSize>, Size, !cast<Type>("HalfPtr" # AS)]>;
+            def : Builtin<name, [Void, VectorType<Double, VSize>, Size, !cast<Type>("HalfPtr" # AS)]>;
           }
         }
       }
     }
   }
 }
 
 let Extension = FuncExtOpenCLCNamedAddressSpaceBuiltins in {

clang/test/SemaOpenCL/half.cl

 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -Wno-unused-value -triple spir-unknown-unknown
-// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -Wno-unused-value -triple spir-unknown-unknown -fdeclare-opencl-builtins -finclude-default-header
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -Wno-unused-value -triple spir-unknown-unknown -fdeclare-opencl-builtins -finclude-default-header -DHAVE_BUILTINS
 
 constant float f = 1.0h; // expected-error{{half precision constant requires cl_khr_fp16}}
 
 half half_disabled(half *p, // expected-error{{declaring function return value of type 'half' is not allowed}}
                    half h)  // expected-error{{declaring function parameter of type '__private half' is not allowed}}
 {
   half a[2]; // expected-error{{declaring variable of type '__private half[2]' is not allowed}}
   half b;    // expected-error{{declaring variable of type '__private half' is not allowed}}
   *p; // expected-error{{loading directly from pointer to type '__private half' requires cl_khr_fp16. Use vector data load builtin functions instead}}
   *p = 0; // expected-error{{assigning directly to pointer to type 'half' requires cl_khr_fp16. Use vector data store builtin functions instead}}
   p[1]; // expected-error{{loading directly from pointer to type '__private half' requires cl_khr_fp16. Use vector data load builtin functions instead}}
   p[1] = 0; // expected-error{{assigning directly to pointer to type 'half' requires cl_khr_fp16. Use vector data store builtin functions instead}}
 
   float c = 1.0f;
   b = (half) c;  // expected-error{{casting to type 'half' is not allowed}}
   c = (float) 1.0h;  // expected-error{{half precision constant requires cl_khr_fp16}}
   b = 1.0h; // expected-error{{half precision constant requires cl_khr_fp16}}
 
   half *allowed = &p[1];
   half *allowed2 = &*p;
   half *allowed3 = p + 1;
 
+#ifdef HAVE_BUILTINS
+  (void)ilogb(*p); // expected-error{{loading directly from pointer to type '__private half' requires cl_khr_fp16. Use vector data load builtin functions instead}}
+  vstore_half(42.0f, 0, p);
+#endif
+
   return h;
 }
 
 kernel void half_disabled_kernel(global half *p,
                                  half h);  // expected-error{{declaring function parameter of type '__private half' is not allowed; did you forget * ?}}
 
 // Exactly the same as above but with the cl_khr_fp16 extension enabled.
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
half half_enabled(half *p, half h)
   b = (half) c;
   c = (float) 1.0h;
   b = 1.0h;
 
   half *allowed = &p[1];
   half *allowed2 = &*p;
   half *allowed3 = p + 1;
 
+#ifdef HAVE_BUILTINS
+  (void)ilogb(*p);
+  vstore_half(42.0f, 0, p);
+#endif
+
   return h;
 }
 
 kernel void half_enabled_kernel(global half *p,
                                 half h);