[PATCH] D46871: [AMDGPU] Add interpolation builtins

[Tim Corringham via Phabricator via cfe-commits]

timcorringham added inline comments.


================
Comment at: include/clang/Basic/BuiltinsAMDGPU.def:103-107
+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")
----------------
arsenm wrote:
> You will also need C++ handling for these when the intrinsics are fixed to use name mangling
Can you expand on that please, I don't know what needs to be done here?


================
Comment at: test/CodeGenOpenCL/builtins-amdgcn-interp.cl:2-28
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -S -o - %s | FileCheck %s --check-prefixes=CHECK,GFX9,BANK32
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu fiji -S -o - %s | FileCheck %s --check-prefixes=CHECK,GFX8,BANK32
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx810 -S -o - %s | FileCheck %s --check-prefixes=CHECK,GFX8,BANK16
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+// CHECK-LABEL: test_interp_f16
----------------
arsenm wrote:
> These should be emitting / checking the IR, not the output asm
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.


Repository:
  rC Clang

https://reviews.llvm.org/D46871