Author: timcorringham Date: Mon Jan 28 05:50:37 2019 New Revision: 352358 URL: http://llvm.org/viewvc/llvm-project?rev=352358&view=rev Log: [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 Added: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-interp.cl Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def?rev=352358&r1=352357&r2=352358&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original) +++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Mon Jan 28 05:50:37 2019 @@ -106,6 +106,15 @@ TARGET_BUILTIN(__builtin_amdgcn_s_dcache 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. //===----------------------------------------------------------------------===// Added: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-interp.cl URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-interp.cl?rev=352358&view=auto ============================================================================== --- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-interp.cl (added) +++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-interp.cl Mon Jan 28 05:50:37 2019 @@ -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); +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
