r323890 - [AMDGPU] Add ds_fadd, ds_fmin, ds_fmax builtins functions

[Daniil Fukalov via cfe-commits]

Author: dfukalov
Date: Wed Jan 31 08:55:09 2018
New Revision: 323890

URL: http://llvm.org/viewvc/llvm-project?rev=323890&view=rev
Log:
[AMDGPU] Add ds_fadd, ds_fmin, ds_fmax builtins functions

Reviewed by arsenm

Differential Revision: https://reviews.llvm.org/D42578

Modified:
    cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
    cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl

Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def?rev=323890&r1=323889&r2=323890&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original)
+++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Wed Jan 31 08:55:09 2018
@@ -93,6 +93,9 @@ BUILTIN(__builtin_amdgcn_ds_bpermute, "i
 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.

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=323890&r1=323889&r2=323890&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl Wed Jan 31 08:55:09 2018
@@ -89,3 +89,23 @@ void test_mov_dpp(global int* out, int s
   *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(local 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(local 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(local float *out, float src)
+{
+  *out = __builtin_amdgcn_ds_fmax(out, src, 0, 0, false);
+}