This is an archive of the discontinued LLVM Phabricator instance.

Differential D30360

AMDGPU: expose l1 cache invalidation intrinsics
ClosedPublic

Authored by jvesely on Feb 24 2017, 3:47 PM.

Details

Reviewers
arsenm
Commits
rGc25509751713: AMDGPU: export l1 cache invalidation intrinsics
rC296240: AMDGPU: export l1 cache invalidation intrinsics
rL296240: AMDGPU: export l1 cache invalidation intrinsics

Diff Detail

Repository
rL LLVM

Event Timeline

jvesely created this revision.Feb 24 2017, 3:47 PM
Herald added subscribers: tpr, dstuttard, tony-tye and 4 others. · View Herald TranscriptFeb 24 2017, 3:47 PM
arsenm accepted this revision.Feb 24 2017, 4:37 PM

LGTM

This revision is now accepted and ready to land.Feb 24 2017, 4:37 PM
Closed by commit rL296240: AMDGPU: export l1 cache invalidation intrinsics (authored by jvesely). · Explain WhyFeb 24 2017, 8:32 PM
This revision was automatically updated to reflect the committed changes.

Revision Contents

PathSize
cfe/
trunk/
include/
clang/
Basic/
2 lines
test/
CodeGenOpenCL/
14 lines

Diff 89771

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

Show All 33 Lines
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// Instruction builtins. // Instruction builtins.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
BUILTIN(__builtin_amdgcn_s_getreg, "UiIi", "n") BUILTIN(__builtin_amdgcn_s_getreg, "UiIi", "n")
BUILTIN(__builtin_amdgcn_s_waitcnt, "vIi", "n") BUILTIN(__builtin_amdgcn_s_waitcnt, "vIi", "n")
BUILTIN(__builtin_amdgcn_s_barrier, "v", "n") BUILTIN(__builtin_amdgcn_s_barrier, "v", "n")
BUILTIN(__builtin_amdgcn_wave_barrier, "v", "n") BUILTIN(__builtin_amdgcn_wave_barrier, "v", "n")
BUILTIN(__builtin_amdgcn_s_dcache_inv, "v", "n")
BUILTIN(__builtin_amdgcn_buffer_wbinvl1, "v", "n")
BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n") BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n")
BUILTIN(__builtin_amdgcn_div_scalef, "fffbb*", "n") BUILTIN(__builtin_amdgcn_div_scalef, "fffbb*", "n")
BUILTIN(__builtin_amdgcn_div_fmas, "ddddb", "nc") BUILTIN(__builtin_amdgcn_div_fmas, "ddddb", "nc")
BUILTIN(__builtin_amdgcn_div_fmasf, "ffffb", "nc") BUILTIN(__builtin_amdgcn_div_fmasf, "ffffb", "nc")
BUILTIN(__builtin_amdgcn_div_fixup, "dddd", "nc") BUILTIN(__builtin_amdgcn_div_fixup, "dddd", "nc")
BUILTIN(__builtin_amdgcn_div_fixupf, "ffff", "nc") BUILTIN(__builtin_amdgcn_div_fixupf, "ffff", "nc")
BUILTIN(__builtin_amdgcn_trig_preop, "ddi", "nc") BUILTIN(__builtin_amdgcn_trig_preop, "ddi", "nc")
BUILTIN(__builtin_amdgcn_trig_preopf, "ffi", "nc") BUILTIN(__builtin_amdgcn_trig_preopf, "ffi", "nc")
▲ Show 20 Lines • Show All 83 Lines • Show Last 20 Lines

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

Show First 20 Lines • Show All 257 Lines • ▼ Show 20 Lines
// CHECK-LABEL: @test_class_f64 // CHECK-LABEL: @test_class_f64
// CHECK: call i1 @llvm.amdgcn.class.f64 // CHECK: call i1 @llvm.amdgcn.class.f64
void test_class_f64(global double* out, double a, int b) void test_class_f64(global double* out, double a, int b)
{ {
*out = __builtin_amdgcn_class(a, b); *out = __builtin_amdgcn_class(a, b);
} }
// CHECK-LABEL: @test_buffer_wbinvl1
// CHECK: call void @llvm.amdgcn.buffer.wbinvl1(
void test_buffer_wbinvl1()
{
__builtin_amdgcn_buffer_wbinvl1();
}
// CHECK-LABEL: @test_s_dcache_inv
// CHECK: call void @llvm.amdgcn.s.dcache.inv(
void test_s_dcache_inv()
{
__builtin_amdgcn_s_dcache_inv();
}
// CHECK-LABEL: @test_s_waitcnt // CHECK-LABEL: @test_s_waitcnt
// CHECK: call void @llvm.amdgcn.s.waitcnt( // CHECK: call void @llvm.amdgcn.s.waitcnt(
void test_s_waitcnt() void test_s_waitcnt()
{ {
__builtin_amdgcn_s_waitcnt(0); __builtin_amdgcn_s_waitcnt(0);
} }
// CHECK-LABEL: @test_s_barrier // CHECK-LABEL: @test_s_barrier
▲ Show 20 Lines • Show All 143 Lines • Show Last 20 Lines