Add half load and store builtins

[jvesely]

Summary

This enables load/stores of half type, without half being a legal type.

[Anastasia]

include/clang/Basic/Builtins.def

1427

I think this should be a language builtin (see above) but perhaps we might need to extend the language version here. Because I believe we only have OpenCL v2.0 currently.

Also this should only be available if cl_khr_fp16 is supported and enabled? I think we are doing similar with some subgroups functions (e.g. get_kernel_sub_group_count_for_ndrange) that are only supported by cl_khr_subgroup but those have custom diagnostic though. May be we could leave this check out since half is not available if cl_khr_fp16 is not enabled anyways.

test/CodeGenOpenCL/no-half.cl

3

It seems strange that cl_khr_fp16 is not enabled too.

[jvesely]

include/clang/Basic/Builtins.def

1427

This is specifically meant to be used when cl_khr_fp16 is not available.
CLC allows using half as storage format and half pointers without the extension,
vstore_half/vload_half are used to load/store half values. (CL1.2 CH 6.1.1.1)

These builtins are not necessary if cl_khr_fp16 is available (we can use regular loads/stores).

I'll take stab at making these CLC only, but similarly to device specific builtins it looked useful beyond that, since these builtins provide access to half type storage.

[jvesely]

restrict builtins to OCLC langauges

[Anastasia]

include/clang/Basic/Builtins.def

1427

Strange. This is not how I would interpret from the extension spec though: https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/cl_khr_fp16.html

But I think for this change is probably fine indeed because this doesn't affect half type itself.

[Anastasia]

LGTM! Thanks!

[Anastasia]

test/CodeGenOpenCL/no-half.cl

20

Would it make sense to add a check for load similarly to store in the test_load_float/test_load_double tests?

28

Minor thing: any reason you are not checking the load fully?

[jvesely]

fully check loads in tests

[jvesely]

include/clang/Basic/Builtins.def

1427

I'm not sure I see the conflict here. cl_khr_fp16 adds support for half scalar and halfn vector types.
without the extension the specs say (OCL 1.2 Ch. 6.1.1.1):

The half data type can only be used to declare a pointer to a buffer that contains half values.

vload_half and vstore_half used to access those buffers without needing half type (or the cl_khr_fp16 extension).

But I think for this change is probably fine indeed because this doesn't affect half type itself.

exactly. this is needed outside of cl_khr_fp16, or the half type

test/CodeGenOpenCL/no-half.cl

20

there is no load. fptrunc double %foo to half uses the function parameter directly

28

just my laziness, I've added full check.

[jvesely]

mark load pointers const

[jvesely]

please let me know if your accept still stands for the modified version.

[Anastasia]

test/CodeGenOpenCL/no-half.cl

28

Could we do the same for the above examples too?

[jvesely]

test/CodeGenOpenCL/no-half.cl

28

I don't understand.
if you mean test_store_*, those functions do not generate any load instructions. the full generated code is:
test_store_double:

entry:
  %0 = fptrunc double %foo to half
  store half %0, half addrspace(1)* %bar, align 2
  ret void

[Anastasia]

Yes, sorry overlooked that. :) LGTM! Thanks!

include/clang/Basic/Builtins.def

 LANGBUILTIN(get_kernel_max_sub_group_size_for_ndrange, "Ui.", "tn", OCLC20_LANG)
 LANGBUILTIN(get_kernel_sub_group_count_for_ndrange, "Ui.", "tn", OCLC20_LANG)
 
 // OpenCL v2.0 s6.13.9 - Address space qualifier functions.
 LANGBUILTIN(to_global, "v*v*", "tn", OCLC20_LANG)
 LANGBUILTIN(to_local, "v*v*", "tn", OCLC20_LANG)
 LANGBUILTIN(to_private, "v*v*", "tn", OCLC20_LANG)
 
+// OpenCL half load/store builtin
+BUILTIN(__builtin_store_half, "vdh*", "n")
+BUILTIN(__builtin_store_halff, "vfh*", "n")
+BUILTIN(__builtin_load_half, "dh*", "nc")
+BUILTIN(__builtin_load_halff, "fh*", "nc")
+
 // Builtins for os_log/os_trace
 BUILTIN(__builtin_os_log_format_buffer_size, "zcC*.", "p:0:nut")
 BUILTIN(__builtin_os_log_format, "v*v*cC*.", "p:0:nt")
 
 // Builtins for XRay
 BUILTIN(__xray_customevent, "vcC*z", "")
 
 // Win64-compatible va_list functions
 BUILTIN(__builtin_ms_va_start, "vc*&.", "nt")
 BUILTIN(__builtin_ms_va_end, "vc*&", "n")
 BUILTIN(__builtin_ms_va_copy, "vc*&c*&", "n")
 
 #undef BUILTIN
 #undef LIBBUILTIN
 #undef LANGBUILTIN

lib/CodeGen/CGBuiltin.cpp

const char *Name =
             : "__get_kernel_sub_group_count_for_ndrange_impl";
     return RValue::get(Builder.CreateCall(
         CGM.CreateRuntimeFunction(
             llvm::FunctionType::get(
                 IntTy, {NDRange->getType(), GenericVoidPtrTy}, false),
             Name),
         {NDRange, Block}));
   }
+
+  case Builtin::BI__builtin_store_half:
+  case Builtin::BI__builtin_store_halff: {
+    Value *Val = EmitScalarExpr(E->getArg(0));
+    Address Address = EmitPointerWithAlignment(E->getArg(1));
+    Value *HalfVal = Builder.CreateFPTrunc(Val, Builder.getHalfTy());
+    return RValue::get(Builder.CreateStore(HalfVal, Address));
+  }
+  case Builtin::BI__builtin_load_half: {
+    Address Address = EmitPointerWithAlignment(E->getArg(0));
+    Value *HalfVal = Builder.CreateLoad(Address);
+    return RValue::get(Builder.CreateFPExt(HalfVal, Builder.getDoubleTy()));
+  }
+  case Builtin::BI__builtin_load_halff: {
+    Address Address = EmitPointerWithAlignment(E->getArg(0));
+    Value *HalfVal = Builder.CreateLoad(Address);
+    return RValue::get(Builder.CreateFPExt(HalfVal, Builder.getFloatTy()));
+  }
   case Builtin::BIprintf:
     if (getTarget().getTriple().isNVPTX())
       return EmitNVPTXDevicePrintfCallExpr(E, ReturnValue);
     break;
   case Builtin::BI__builtin_canonicalize:
   case Builtin::BI__builtin_canonicalizef:
   case Builtin::BI__builtin_canonicalizel:
     return RValue::get(emitUnaryBuiltin(*this, E, Intrinsic::canonicalize));

test/CodeGenOpenCL/no-half.cl

+// RUN: %clang_cc1 %s -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s
+
+#pragma OPENCL EXTENSION cl_khr_fp64:enable
+
+// CHECK-LABEL: @test_store_float(float %foo, half addrspace({{.}}){{.*}} %bar)
+__kernel void test_store_float(float foo, __global half* bar)
+{
+  __builtin_store_halff(foo, bar);
+// CHECK: [[HALF_VAL:%.*]] = fptrunc float %foo to half
+// CHECK: store half [[HALF_VAL]], half addrspace({{.}})* %bar, align 2
+}
+
+// CHECK-LABEL: @test_store_double(double %foo, half addrspace({{.}}){{.*}} %bar)
+__kernel void test_store_double(double foo, __global half* bar)
+{
+  __builtin_store_half(foo, bar);
+// CHECK: [[HALF_VAL:%.*]] = fptrunc double %foo to half
+// CHECK: store half [[HALF_VAL]], half addrspace({{.}})* %bar, align 2
+}
+
+// CHECK-LABEL: @test_load_float(float addrspace({{.}}){{.*}} %foo, half addrspace({{.}}){{.*}} %bar)
+__kernel void test_load_float(__global float* foo, __global half* bar)
+{
+  foo[0] = __builtin_load_halff(bar);
+// CHECK: [[HALF_VAL:%.*]] = load
+// CHECK: [[FULL_VAL:%.*]] = fpext half [[HALF_VAL]] to float
+// CHECK: store float [[FULL_VAL]], float addrspace({{.}})* %foo
+}
+
+// CHECK-LABEL: @test_load_double(double addrspace({{.}}){{.*}} %foo, half addrspace({{.}}){{.*}} %bar)
+__kernel void test_load_double(__global double* foo, __global half* bar)
+{
+  foo[0] = __builtin_load_half(bar);
+// CHECK: [[HALF_VAL:%.*]] = load
+// CHECK: [[FULL_VAL:%.*]] = fpext half [[HALF_VAL]] to double
+// CHECK: store double [[FULL_VAL]], double addrspace({{.}})* %foo
+}