[PATCH] D125401: [OpenCL] Do not guard vload/store_half builtins

[Sven van Haastregt via Phabricator via cfe-commits]

svenvh created this revision.
svenvh added a reviewer: Anastasia.
svenvh added a project: clang.
Herald added subscribers: Naghasan, ldrumm, yaxunl.
Herald added a project: All.
svenvh requested review of this revision.
Herald added a subscriber: cfe-commits.

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


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D125401

Files:
  clang/lib/Headers/opencl-c-base.h
  clang/lib/Sema/OpenCLBuiltins.td


Index: clang/lib/Sema/OpenCLBuiltins.td
===================================================================
--- clang/lib/Sema/OpenCLBuiltins.td
+++ clang/lib/Sema/OpenCLBuiltins.td
@@ -352,9 +352,22 @@
 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()">>;
@@ -877,22 +890,22 @@
 
 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)]>;
           }
         }
       }
Index: clang/lib/Headers/opencl-c-base.h
===================================================================
--- clang/lib/Headers/opencl-c-base.h
+++ clang/lib/Headers/opencl-c-base.h
@@ -202,6 +202,9 @@
 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__)


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D125401.428703.patch
Type: text/x-patch
Size: 3535 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20220511/77b0337f/attachment.bin>