diff options
Diffstat (limited to 'clang')
| -rw-r--r-- | clang/include/clang/Basic/BuiltinsAMDGPU.def | 5 | ||||
| -rw-r--r-- | clang/lib/Basic/Targets.cpp | 1 | ||||
| -rw-r--r-- | clang/lib/CodeGen/CGBuiltin.cpp | 8 | ||||
| -rw-r--r-- | clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl | 8 | ||||
| -rw-r--r-- | clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 28 | ||||
| -rw-r--r-- | clang/test/SemaOpenCL/builtins-amdgcn-error.cl | 20 |
6 files changed, 65 insertions, 5 deletions
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index 15482775488..a8ab657c379 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -86,6 +86,10 @@ BUILTIN(__builtin_amdgcn_sicmpl, "LUiLiLiIi", "nc") 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_exph, "sh", "nc", "16-bit-insts") 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. diff --git a/clang/lib/Basic/Targets.cpp b/clang/lib/Basic/Targets.cpp index 6b0b9100b3e..e23ff3cc380 100644 --- a/clang/lib/Basic/Targets.cpp +++ b/clang/lib/Basic/Targets.cpp @@ -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: diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index 1c7ee2e6d49..85ed2b86777 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -8401,6 +8401,14 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, 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: diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl index 0c88fdfb724..1dad6749183 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl @@ -81,3 +81,11 @@ void test_s_memrealtime(global ulong* out) { *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); +} + diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl index e33c3ca5266..a19ce2f1ed5 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -235,6 +235,34 @@ void test_ds_swizzle(global int* out, int a) *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) diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error.cl index 6cb2f388225..2639bf27752 100644 --- a/clang/test/SemaOpenCL/builtins-amdgcn-error.cl +++ b/clang/test/SemaOpenCL/builtins-amdgcn-error.cl @@ -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 a) { *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}} +} + |
