From: Tim Corringham <tcorring@amd.com>
Date: Mon, 28 Jan 2019 13:50:37 +0000 (+0000)
Subject: [AMDGPU] Add interpolation builtins
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=62765d829f254f8888a8fa18a2ea091193f0e40e;p=clang

[AMDGPU] Add interpolation builtins

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

Reviewers: arsenm, tpr, dstuttard, #amdgpu

Reviewed By: arsenm, #amdgpu

Subscribers: kzhuravl, wdng, nhaehnle, yaxunl, dstuttard, tpr, t-tye, cfe-commits

Differential Revision: https://reviews.llvm.org/D46871

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@352358 91177308-0d34-0410-b5e6-96231b3b80d8
---

diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def
index 068b61ec86..547870447d 100644
--- a/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/include/clang/Basic/BuiltinsAMDGPU.def
@@ -105,6 +105,15 @@ BUILTIN(__builtin_amdgcn_ds_fmaxf, "ff*3fIiIiIb", "n")
 TARGET_BUILTIN(__builtin_amdgcn_s_dcache_inv_vol, "v", "n", "ci-insts")
 TARGET_BUILTIN(__builtin_amdgcn_buffer_wbinvl1_vol, "v", "n", "ci-insts")
 
+//===----------------------------------------------------------------------===//
+// 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.
 //===----------------------------------------------------------------------===//
diff --git a/test/CodeGenOpenCL/builtins-amdgcn-interp.cl b/test/CodeGenOpenCL/builtins-amdgcn-interp.cl
new file mode 100644
index 0000000000..39d913e902
--- /dev/null
+++ b/test/CodeGenOpenCL/builtins-amdgcn-interp.cl
@@ -0,0 +1,34 @@
+// 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);
+}