This revision was automatically updated to reflect the committed changes. Closed by commit rGb250cca11d59: [OpenCL] Do not guard vload/store_half builtins (authored by svenvh).
Repository: rG LLVM Github Monorepo CHANGES SINCE LAST ACTION https://reviews.llvm.org/D125401/new/ https://reviews.llvm.org/D125401 Files: clang/lib/Headers/opencl-c-base.h clang/lib/Sema/OpenCLBuiltins.td clang/test/SemaOpenCL/half.cl
Index: clang/test/SemaOpenCL/half.cl =================================================================== --- clang/test/SemaOpenCL/half.cl +++ clang/test/SemaOpenCL/half.cl @@ -1,5 +1,5 @@ // 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}} @@ -22,6 +22,11 @@ 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; } @@ -49,6 +54,11 @@ half *allowed2 = &*p; half *allowed3 = p + 1; +#ifdef HAVE_BUILTINS + (void)ilogb(*p); + vstore_half(42.0f, 0, p); +#endif + return h; } 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__)
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits