r297436 - [AMDGPU] Add builtin functions readlane ds_permute mov_dpp
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Thu Mar 9 17:30:47 PST 2017
Author: yaxunl
Date: Thu Mar 9 19:30:46 2017
New Revision: 297436
URL: http://llvm.org/viewvc/llvm-project?rev=297436&view=rev
Log:
[AMDGPU] Add builtin functions readlane ds_permute mov_dpp
Differential Revision: https://reviews.llvm.org/D30551
Modified:
cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
cfe/trunk/lib/Basic/Targets.cpp
cfe/trunk/lib/CodeGen/CGBuiltin.cpp
cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl
Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def?rev=297436&r1=297435&r2=297436&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original)
+++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Thu Mar 9 19:30:46 2017
@@ -86,6 +86,10 @@ BUILTIN(__builtin_amdgcn_sicmpl, "LUiLiL
BUILTIN(__builtin_amdgcn_fcmp, "LUiddIi", "nc")
BUILTIN(__builtin_amdgcn_fcmpf, "LUiffIi", "nc")
BUILTIN(__builtin_amdgcn_ds_swizzle, "iiIi", "nc")
+BUILTIN(__builtin_amdgcn_ds_permute, "iii", "nc")
+BUILTIN(__builtin_amdgcn_ds_bpermute, "iii", "nc")
+BUILTIN(__builtin_amdgcn_readfirstlane, "ii", "nc")
+BUILTIN(__builtin_amdgcn_readlane, "iii", "nc")
BUILTIN(__builtin_amdgcn_fmed3f, "ffff", "nc")
//===----------------------------------------------------------------------===//
@@ -103,6 +107,7 @@ TARGET_BUILTIN(__builtin_amdgcn_frexp_ex
TARGET_BUILTIN(__builtin_amdgcn_fracth, "hh", "nc", "16-bit-insts")
TARGET_BUILTIN(__builtin_amdgcn_classh, "bhi", "nc", "16-bit-insts")
TARGET_BUILTIN(__builtin_amdgcn_s_memrealtime, "LUi", "n", "s-memrealtime")
+TARGET_BUILTIN(__builtin_amdgcn_mov_dpp, "iiIiIiIiIb", "nc", "dpp")
//===----------------------------------------------------------------------===//
// GFX9+ only builtins.
Modified: cfe/trunk/lib/Basic/Targets.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets.cpp?rev=297436&r1=297435&r2=297436&view=diff
==============================================================================
--- cfe/trunk/lib/Basic/Targets.cpp (original)
+++ cfe/trunk/lib/Basic/Targets.cpp Thu Mar 9 19:30:46 2017
@@ -2387,6 +2387,7 @@ bool AMDGPUTargetInfo::initFeatureMap(
case GK_GFX8:
Features["s-memrealtime"] = true;
Features["16-bit-insts"] = true;
+ Features["dpp"] = true;
break;
case GK_NONE:
Modified: cfe/trunk/lib/CodeGen/CGBuiltin.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGBuiltin.cpp?rev=297436&r1=297435&r2=297436&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGBuiltin.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGBuiltin.cpp Thu Mar 9 19:30:46 2017
@@ -8401,6 +8401,14 @@ Value *CodeGenFunction::EmitAMDGPUBuilti
case AMDGPU::BI__builtin_amdgcn_ds_swizzle:
return emitBinaryBuiltin(*this, E, Intrinsic::amdgcn_ds_swizzle);
+ case AMDGPU::BI__builtin_amdgcn_mov_dpp: {
+ llvm::SmallVector<llvm::Value *, 5> Args;
+ for (unsigned I = 0; I != 5; ++I)
+ Args.push_back(EmitScalarExpr(E->getArg(I)));
+ Value *F = CGM.getIntrinsic(Intrinsic::amdgcn_mov_dpp,
+ Args[0]->getType());
+ return Builder.CreateCall(F, Args);
+ }
case AMDGPU::BI__builtin_amdgcn_div_fixup:
case AMDGPU::BI__builtin_amdgcn_div_fixupf:
case AMDGPU::BI__builtin_amdgcn_div_fixuph:
Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl?rev=297436&r1=297435&r2=297436&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl Thu Mar 9 19:30:46 2017
@@ -81,3 +81,11 @@ void test_s_memrealtime(global ulong* ou
{
*out = __builtin_amdgcn_s_memrealtime();
}
+
+// CHECK-LABEL: @test_mov_dpp
+// CHECK: call i32 @llvm.amdgcn.mov.dpp.i32(i32 %src, i32 0, i32 0, i32 0, i1 false)
+void test_mov_dpp(global int* out, int src)
+{
+ *out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0, false);
+}
+
Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl?rev=297436&r1=297435&r2=297436&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Thu Mar 9 19:30:46 2017
@@ -235,6 +235,34 @@ void test_ds_swizzle(global int* out, in
*out = __builtin_amdgcn_ds_swizzle(a, 32);
}
+// CHECK-LABEL: @test_ds_permute
+// CHECK: call i32 @llvm.amdgcn.ds.permute(i32 %a, i32 %b)
+void test_ds_permute(global int* out, int a, int b)
+{
+ out[0] = __builtin_amdgcn_ds_permute(a, b);
+}
+
+// CHECK-LABEL: @test_ds_bpermute
+// CHECK: call i32 @llvm.amdgcn.ds.bpermute(i32 %a, i32 %b)
+void test_ds_bpermute(global int* out, int a, int b)
+{
+ *out = __builtin_amdgcn_ds_bpermute(a, b);
+}
+
+// CHECK-LABEL: @test_readfirstlane
+// CHECK: call i32 @llvm.amdgcn.readfirstlane(i32 %a)
+void test_readfirstlane(global int* out, int a)
+{
+ *out = __builtin_amdgcn_readfirstlane(a);
+}
+
+// CHECK-LABEL: @test_readlane
+// CHECK: call i32 @llvm.amdgcn.readlane(i32 %a, i32 %b)
+void test_readlane(global int* out, int a, int b)
+{
+ *out = __builtin_amdgcn_readlane(a, b);
+}
+
// CHECK-LABEL: @test_fcmp_f32
// CHECK: call i64 @llvm.amdgcn.fcmp.f32(float %a, float %b, i32 5)
void test_fcmp_f32(global ulong* out, float a, float b)
Modified: cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl?rev=297436&r1=297435&r2=297436&view=diff
==============================================================================
--- cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl (original)
+++ cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl Thu Mar 9 19:30:46 2017
@@ -1,16 +1,17 @@
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s
-// FIXME: We only get one error if the functions are the other order in the
-// file.
-
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
typedef unsigned long ulong;
typedef unsigned int uint;
-ulong test_s_memrealtime()
+// To get all errors for feature checking we need to put them in one function
+// since Clang will stop codegen for the next function if it finds error during
+// codegen of the previous function.
+void test_target_builtin(global int* out, int a)
{
- return __builtin_amdgcn_s_memrealtime(); // expected-error {{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}}
+ __builtin_amdgcn_s_memrealtime(); // expected-error {{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}}
+ *out = __builtin_amdgcn_mov_dpp(a, 0, 0, 0, false); // expected-error {{'__builtin_amdgcn_mov_dpp' needs target feature dpp}}
}
void test_s_sleep(int x)
@@ -92,3 +93,12 @@ void test_s_getreg(global int* out, int
{
*out = __builtin_amdgcn_s_getreg(a); // expected-error {{argument to '__builtin_amdgcn_s_getreg' must be a constant integer}}
}
+
+void test_mov_dpp2(global int* out, int a, int b, int c, int d, bool e)
+{
+ *out = __builtin_amdgcn_mov_dpp(a, b, 0, 0, false); // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_mov_dpp(a, 0, c, 0, false); // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_mov_dpp(a, 0, 0, d, false); // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_mov_dpp(a, 0, 0, 0, e); // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
+}
+
More information about the cfe-commits
mailing list