[AMDGPU] Add interpolation builtins

[timcorringham]

Summary

Added builtins for the interpolation intrinsics, and related LIT
test.

[arsenm]

include/clang/Basic/BuiltinsAMDGPU.def

103–107

You will also need C++ handling for these when the intrinsics are fixed to use name mangling

test/CodeGenOpenCL/builtins-amdgcn-interp.cl

3–29

These should be emitting / checking the IR, not the output asm

[timcorringham]

include/clang/Basic/BuiltinsAMDGPU.def

103–107

Can you expand on that please, I don't know what needs to be done here?

test/CodeGenOpenCL/builtins-amdgcn-interp.cl

3–29

Checking the IR doesn't allow you to check the behavior very well, as there are differences depending on the target-cpu which are not apparent in the IR.

[arsenm]

include/clang/Basic/BuiltinsAMDGPU.def

103–107

In the other review I commented you should change the existing intrinsics to use name mangling instead of adding a second set. Once you do that, GCCBuiltin no longer works. You have to add some code like the other type mangled intrinsics in lib/CodeGen/CGBuiltin.cpp

test/CodeGenOpenCL/builtins-amdgcn-interp.cl

3–29

That is a codegen test which we have (or should have) in the backend. I know we have some for the f32 variants.

Clang tests should only be checking that the produced IR is correct

[timcorringham]

[AMDGPU] Add interpolation builtins

Updated LIT test so that it only tests for the existence of the builtins
by checking the IR.

[arsenm]

LGTM. Checking the full operands wouldn't hurt though.

include/clang/Basic/BuiltinsAMDGPU.def

 BUILTIN(__builtin_amdgcn_readfirstlane, "ii", "nc")
 BUILTIN(__builtin_amdgcn_readlane, "iii", "nc")
 BUILTIN(__builtin_amdgcn_fmed3f, "ffff", "nc")
 BUILTIN(__builtin_amdgcn_ds_fadd, "ff*3fiib", "n")
 BUILTIN(__builtin_amdgcn_ds_fmin, "ff*3fiib", "n")
 BUILTIN(__builtin_amdgcn_ds_fmax, "ff*3fiib", "n")
 
 //===----------------------------------------------------------------------===//
+// Interpolation builtins.
+//===----------------------------------------------------------------------===//
+BUILTIN(__builtin_amdgcn_interp_p1_f16, "ffUiUibUi", "nc")
+BUILTIN(__builtin_amdgcn_interp_p2_f16, "hffUiUibUi", "nc")
+BUILTIN(__builtin_amdgcn_interp_p1, "ffUiUiUi", "nc")
+BUILTIN(__builtin_amdgcn_interp_p2, "fffUiUiUi", "nc")
+BUILTIN(__builtin_amdgcn_interp_mov, "fUiUiUiUi", "nc")
+
+//===----------------------------------------------------------------------===//
 // VI+ only builtins.
 //===----------------------------------------------------------------------===//
 
 TARGET_BUILTIN(__builtin_amdgcn_div_fixuph, "hhhh", "nc", "16-bit-insts")
 TARGET_BUILTIN(__builtin_amdgcn_rcph, "hh", "nc", "16-bit-insts")
 TARGET_BUILTIN(__builtin_amdgcn_rsqh, "hh", "nc", "16-bit-insts")
 TARGET_BUILTIN(__builtin_amdgcn_sinh, "hh", "nc", "16-bit-insts")
 TARGET_BUILTIN(__builtin_amdgcn_cosh, "hh", "nc", "16-bit-insts")

test/CodeGenOpenCL/builtins-amdgcn-interp.cl

+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -S -emit-llvm -o - %s | FileCheck %s
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+// CHECK-LABEL: test_interp_f16
+// CHECK: call float @llvm.amdgcn.interp.p1.f16
+// CHECK: call half @llvm.amdgcn.interp.p2.f16
+// CHECK: call float @llvm.amdgcn.interp.p1.f16
+// CHECK: call half @llvm.amdgcn.interp.p2.f16
+void test_interp_f16(global half* out, float i, float j, int m0)
+{
+  float p1_0 = __builtin_amdgcn_interp_p1_f16(i, 2, 3, false, m0);
+  half p2_0 = __builtin_amdgcn_interp_p2_f16(p1_0, j, 2, 3, false, m0);
+  float p1_1 = __builtin_amdgcn_interp_p1_f16(i, 2, 3, true, m0);
+  half p2_1 = __builtin_amdgcn_interp_p2_f16(p1_1, j, 2, 3, true, m0);
+  *out = p2_0 + p2_1;
+}
+
+// CHECK-LABEL: test_interp_f32
+// CHECK: call float @llvm.amdgcn.interp.p1
+// CHECK: call float @llvm.amdgcn.interp.p2
+void test_interp_f32(global float* out, float i, float j, int m0)
+{
+  float p1 = __builtin_amdgcn_interp_p1(i, 1, 4, m0);
+  *out = __builtin_amdgcn_interp_p2(p1, j, 1, 4, m0);
+}
+
+// CHECK-LABEL: test_interp_mov
+// CHECK: call float @llvm.amdgcn.interp.mov
+void test_interp_mov(global float* out, float i, float j, int m0)
+{
+  *out = __builtin_amdgcn_interp_mov(2, 3, 4, m0);
+}