This is an archive of the discontinued LLVM Phabricator instance.

Differential D42578

[AMDGPU] Add ds_fadd, ds_fmin, ds_fmax builtins functions
ClosedPublic

Authored by dfukalov on Jan 26 2018, 7:25 AM.

Details

Reviewers
arsenm
b-sumner
Commits
rGe72cde57d297: [AMDGPU] Add ds_fadd, ds_fmin, ds_fmax builtins functions
rC323890: [AMDGPU] Add ds_fadd, ds_fmin, ds_fmax builtins functions
rL323890: [AMDGPU] Add ds_fadd, ds_fmin, ds_fmax builtins functions

Diff Detail

Repository
rL LLVM

Event Timeline

dfukalov created this revision.Jan 26 2018, 7:25 AM
Herald added subscribers: cfe-commits, t-tye, tpr and 5 others. · View Herald TranscriptJan 26 2018, 7:25 AM
b-sumner added a comment.Jan 26 2018, 7:48 AM

Were you going to add min and max separately?

dfukalov updated this revision to Diff 131593.Jan 26 2018, 8:29 AM
dfukalov retitled this revision from [AMDGPU] Add ds_fadd builtin function to [AMDGPU] Add ds_fadd, ds_fmin, ds_fmax builtins functions.

Sorry, missed them

arsenm added inline comments.Jan 29 2018, 8:13 AM
include/clang/Basic/BuiltinsAMDGPU.def
96–98 ↗(On Diff #131593)

These are definitely not const intrinsics and need the c removed

dfukalov updated this revision to Diff 131820.Jan 29 2018, 9:47 AM

fixed builtins descriptions

dfukalov marked an inline comment as done.Jan 29 2018, 9:47 AM
b-sumner added a comment.Jan 29 2018, 2:29 PM

Should we expect that the last 3 arguments have any effect? Do we want to test to ensure they have the expected effects?

arsenm accepted this revision.Jan 29 2018, 10:21 PM

LGTM

This revision is now accepted and ready to land.Jan 29 2018, 10:21 PM
Closed by commit rL323890: [AMDGPU] Add ds_fadd, ds_fmin, ds_fmax builtins functions (authored by dfukalov). · Explain WhyJan 31 2018, 8:59 AM
Closed by commit rC323890: [AMDGPU] Add ds_fadd, ds_fmin, ds_fmax builtins functions (authored by dfukalov). · Explain Why
This revision was automatically updated to reflect the committed changes.
This revision was automatically updated to reflect the committed changes.
Herald added a subscriber: llvm-commits. · View Herald TranscriptJan 31 2018, 8:59 AM

Revision Contents

PathSize
cfe/
trunk/
include/
clang/
Basic/
3 lines
test/
CodeGenOpenCL/
20 lines

Diff 132191

cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def

Show First 20 Lines • Show All 87 Lines • ▼ Show 20 Lines
BUILTIN(__builtin_amdgcn_fcmp, "LUiddIi", "nc") BUILTIN(__builtin_amdgcn_fcmp, "LUiddIi", "nc")
BUILTIN(__builtin_amdgcn_fcmpf, "LUiffIi", "nc") BUILTIN(__builtin_amdgcn_fcmpf, "LUiffIi", "nc")
BUILTIN(__builtin_amdgcn_ds_swizzle, "iiIi", "nc") BUILTIN(__builtin_amdgcn_ds_swizzle, "iiIi", "nc")
BUILTIN(__builtin_amdgcn_ds_permute, "iii", "nc") BUILTIN(__builtin_amdgcn_ds_permute, "iii", "nc")
BUILTIN(__builtin_amdgcn_ds_bpermute, "iii", "nc") BUILTIN(__builtin_amdgcn_ds_bpermute, "iii", "nc")
BUILTIN(__builtin_amdgcn_readfirstlane, "ii", "nc") BUILTIN(__builtin_amdgcn_readfirstlane, "ii", "nc")
BUILTIN(__builtin_amdgcn_readlane, "iii", "nc") BUILTIN(__builtin_amdgcn_readlane, "iii", "nc")
BUILTIN(__builtin_amdgcn_fmed3f, "ffff", "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. // VI+ only builtins.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
TARGET_BUILTIN(__builtin_amdgcn_div_fixuph, "hhhh", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_div_fixuph, "hhhh", "nc", "16-bit-insts")
TARGET_BUILTIN(__builtin_amdgcn_rcph, "hh", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_rcph, "hh", "nc", "16-bit-insts")
TARGET_BUILTIN(__builtin_amdgcn_rsqh, "hh", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_rsqh, "hh", "nc", "16-bit-insts")
▲ Show 20 Lines • Show All 42 Lines • Show Last 20 Lines

cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl

Show First 20 Lines • Show All 83 Lines • ▼ Show 20 Lines
// CHECK-LABEL: @test_mov_dpp // CHECK-LABEL: @test_mov_dpp
// CHECK: call i32 @llvm.amdgcn.mov.dpp.i32(i32 %src, i32 0, i32 0, i32 0, i1 false) // 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) void test_mov_dpp(global int* out, int src)
{ {
*out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0, false); *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);
}