From: Daniil Fukalov <daniil.fukalov@amd.com>
Date: Sun, 4 Feb 2018 22:32:07 +0000 (+0000)
Subject: Recommit rL323890: [AMDGPU] Add ds_fadd, ds_fmin, ds_fmax builtins functions
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=b9bdf3d4e882e565e4af6a67d419648ef572cd4f;p=clang

Recommit rL323890: [AMDGPU] Add ds_fadd, ds_fmin, ds_fmax builtins functions

Fixed asserts in tests.

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@324201 91177308-0d34-0410-b5e6-96231b3b80d8
---

diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def
index ec6a0fb917..e63f6e4099 100644
--- a/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/include/clang/Basic/BuiltinsAMDGPU.def
@@ -93,6 +93,9 @@ 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")
+BUILTIN(__builtin_amdgcn_ds_fadd, "ff*3fiib", "n")
+BUILTIN(__builtin_amdgcn_ds_fmin, "ff*3fiib", "n")
+BUILTIN(__builtin_amdgcn_ds_fmax, "ff*3fiib", "n")
 
 //===----------------------------------------------------------------------===//
 // VI+ only builtins.
diff --git a/test/CodeGenOpenCL/builtins-amdgcn-vi.cl b/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
index 1dad674918..caa76e2374 100644
--- a/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
+++ b/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
@@ -89,3 +89,20 @@ void test_mov_dpp(global int* out, int src)
   *out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0, false);
 }
 
+// CHECK-LABEL: @test_ds_fadd
+// CHECK: call float @llvm.amdgcn.ds.fadd(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
+void test_ds_fadd(__attribute__((address_space(3))) float *out, float src) {
+  *out = __builtin_amdgcn_ds_fadd(out, src, 0, 0, false);
+}
+
+// CHECK-LABEL: @test_ds_fmin
+// CHECK: call float @llvm.amdgcn.ds.fmin(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
+void test_ds_fmin(__attribute__((address_space(3))) float *out, float src) {
+  *out = __builtin_amdgcn_ds_fmin(out, src, 0, 0, false);
+}
+
+// CHECK-LABEL: @test_ds_fmax
+// CHECK: call float @llvm.amdgcn.ds.fmax(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
+void test_ds_fmax(__attribute__((address_space(3))) float *out, float src) {
+  *out = __builtin_amdgcn_ds_fmax(out, src, 0, 0, false);
+}