[PATCH] D46871: [AMDGPU] Add interpolation builtins

Matt Arsenault via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Fri May 18 01:33:59 PDT 2018


arsenm 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")
----------------
timcorringham wrote:
> 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?
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


================
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
----------------
timcorringham wrote:
> 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.
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


Repository:
  rC Clang

https://reviews.llvm.org/D46871





More information about the cfe-commits mailing list