This is an archive of the discontinued LLVM Phabricator instance.

Differential D50321

AMDGPU: Add builtin for s_dcache_wb
ClosedPublic

Authored by arsenm on Aug 5 2018, 11:48 PM.

Details

Reviewers
yaxunl
kzhuravl

Diff Detail

Event Timeline

arsenm created this revision.Aug 5 2018, 11:48 PM
Herald added subscribers: t-tye, tpr, dstuttard and 2 others. · View Herald TranscriptAug 5 2018, 11:48 PM
kzhuravl accepted this revision.Aug 6 2018, 10:43 AM

LGTM

This revision is now accepted and ready to land.Aug 6 2018, 10:43 AM
arsenm closed this revision.Aug 7 2018, 12:50 AM

r339110

Revision Contents

PathSize
include/
clang/
Basic/
1 line
test/
CodeGenOpenCL/
7 lines
SemaOpenCL/
8 lines

Diff 159253

include/clang/Basic/BuiltinsAMDGPU.def

Show First 20 Lines • Show All 115 Lines • ▼ Show 20 Lines
TARGET_BUILTIN(__builtin_amdgcn_cosh, "hh", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_cosh, "hh", "nc", "16-bit-insts")
TARGET_BUILTIN(__builtin_amdgcn_ldexph, "hhi", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_ldexph, "hhi", "nc", "16-bit-insts")
TARGET_BUILTIN(__builtin_amdgcn_frexp_manth, "hh", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_frexp_manth, "hh", "nc", "16-bit-insts")
TARGET_BUILTIN(__builtin_amdgcn_frexp_exph, "sh", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_frexp_exph, "sh", "nc", "16-bit-insts")
TARGET_BUILTIN(__builtin_amdgcn_fracth, "hh", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_fracth, "hh", "nc", "16-bit-insts")
TARGET_BUILTIN(__builtin_amdgcn_classh, "bhi", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_classh, "bhi", "nc", "16-bit-insts")
TARGET_BUILTIN(__builtin_amdgcn_s_memrealtime, "LUi", "n", "s-memrealtime") TARGET_BUILTIN(__builtin_amdgcn_s_memrealtime, "LUi", "n", "s-memrealtime")
TARGET_BUILTIN(__builtin_amdgcn_mov_dpp, "iiIiIiIiIb", "nc", "dpp") TARGET_BUILTIN(__builtin_amdgcn_mov_dpp, "iiIiIiIiIb", "nc", "dpp")
TARGET_BUILTIN(__builtin_amdgcn_s_dcache_wb, "v", "n", "vi-insts")
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// GFX9+ only builtins. // GFX9+ only builtins.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
TARGET_BUILTIN(__builtin_amdgcn_fmed3h, "hhhh", "nc", "gfx9-insts") TARGET_BUILTIN(__builtin_amdgcn_fmed3h, "hhhh", "nc", "gfx9-insts")
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
Show All 37 Lines

test/CodeGenOpenCL/builtins-amdgcn-vi.cl

Show First 20 Lines • Show All 76 Lines • ▼ Show 20 Lines
// CHECK-LABEL: @test_s_memrealtime // CHECK-LABEL: @test_s_memrealtime
// CHECK: call i64 @llvm.amdgcn.s.memrealtime() // CHECK: call i64 @llvm.amdgcn.s.memrealtime()
void test_s_memrealtime(global ulong* out) void test_s_memrealtime(global ulong* out)
{ {
*out = __builtin_amdgcn_s_memrealtime(); *out = __builtin_amdgcn_s_memrealtime();
} }
// CHECK-LABEL: @test_s_dcache_wb()
// CHECK: call void @llvm.amdgcn.s.dcache.wb()
void test_s_dcache_wb()
{
__builtin_amdgcn_s_dcache_wb();
}
// 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-LABEL: @test_ds_fadd
Show All 16 Lines

test/SemaOpenCL/builtins-amdgcn-error-vi.cl

  • This file was added.
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s
// RUN: %clang_cc1 -triple amdgcn-- -target-cpu hawaii -verify -S -o - %s
void test_vi_s_dcache_wb()
{
__builtin_amdgcn_s_dcache_wb(); // expected-error {{'__builtin_amdgcn_s_dcache_wb' needs target feature vi-insts}}
}