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.
*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);
-}