Magic-Mirror/llvm-capstone
1
0
Fork 0
mirror of https://github.com/capstone-engine/llvm-capstone.git synced 2024-12-31 20:44:18 +00:00
Code Issues Actions 2 Packages Projects Releases Wiki Activity

[AMDGPU] Add ds_fadd, ds_fmin, ds_fmax builtins functions

Browse Source 
Reviewed by arsenm

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

llvm-svn: 323890
This commit is contained in:
Daniil Fukalov 2018-01-31 16:55:09 +00:00
parent 82a83391d3
commit e72cde57d2
2 changed files with 23 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
clang/include/clang/Basic/BuiltinsAMDGPU.def
View File

@ -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.

20
clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
View File

@ -89,3 +89,23 @@ 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(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);
}