This is an archive of the discontinued LLVM Phabricator instance.

Differential D50493

AMDGPU: Add another missing builtin
ClosedPublic

Authored by arsenm on Aug 8 2018, 11:58 PM.

Details

Reviewers
kzhuravl
b-sumner

Diff Detail

Event Timeline

arsenm created this revision.Aug 8 2018, 11:58 PM
Herald added subscribers: t-tye, tpr, dstuttard and 3 others. · View Herald TranscriptAug 8 2018, 11:58 PM
kzhuravl accepted this revision.Aug 9 2018, 10:15 AM

LGTM

This revision is now accepted and ready to land.Aug 9 2018, 10:15 AM
arsenm closed this revision.Aug 9 2018, 3:19 PM

r339395

Revision Contents

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

Diff 159860

include/clang/Basic/BuiltinsAMDGPU.def

Show First 20 Lines • Show All 98 Lines • ▼ Show 20 Lines
BUILTIN(__builtin_amdgcn_ds_faddf, "ff*3fIiIiIb", "n") BUILTIN(__builtin_amdgcn_ds_faddf, "ff*3fIiIiIb", "n")
BUILTIN(__builtin_amdgcn_ds_fminf, "ff*3fIiIiIb", "n") BUILTIN(__builtin_amdgcn_ds_fminf, "ff*3fIiIiIb", "n")
BUILTIN(__builtin_amdgcn_ds_fmaxf, "ff*3fIiIiIb", "n") BUILTIN(__builtin_amdgcn_ds_fmaxf, "ff*3fIiIiIb", "n")
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// CI+ only builtins. // CI+ only builtins.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
TARGET_BUILTIN(__builtin_amdgcn_s_dcache_inv_vol, "v", "n", "ci-insts") TARGET_BUILTIN(__builtin_amdgcn_s_dcache_inv_vol, "v", "n", "ci-insts")
TARGET_BUILTIN(__builtin_amdgcn_buffer_wbinvl1_vol, "v", "n", "ci-insts")
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// 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 55 Lines • Show Last 20 Lines

test/CodeGenOpenCL/builtins-amdgcn-ci.cl

// REQUIRES: amdgpu-registered-target // REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu hawaii -S -emit-llvm -o - %s | FileCheck %s // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu hawaii -S -emit-llvm -o - %s | FileCheck %s
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu fiji -S -emit-llvm -o - %s | FileCheck %s // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu fiji -S -emit-llvm -o - %s | FileCheck %s
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck %s // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck %s
// CHECK-LABEL: @test_s_dcache_inv_vol // CHECK-LABEL: @test_s_dcache_inv_vol
// CHECK: call void @llvm.amdgcn.s.dcache.inv.vol( // CHECK: call void @llvm.amdgcn.s.dcache.inv.vol(
void test_s_dcache_inv_vol() void test_s_dcache_inv_vol()
{ {
__builtin_amdgcn_s_dcache_inv_vol(); __builtin_amdgcn_s_dcache_inv_vol();
} }
// CHECK-LABEL: @test_buffer_wbinvl1_vol
// CHECK: call void @llvm.amdgcn.buffer.wbinvl1.vol()
void test_buffer_wbinvl1_vol()
{
__builtin_amdgcn_buffer_wbinvl1_vol();
}

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

// REQUIRES: amdgpu-registered-target // REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s // RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s
void test_ci_s_dcache_inv_vol() void test_ci_biltins()
{ {
__builtin_amdgcn_s_dcache_inv_vol(); // expected-error {{'__builtin_amdgcn_s_dcache_inv_vol' needs target feature ci-insts}} __builtin_amdgcn_s_dcache_inv_vol(); // expected-error {{'__builtin_amdgcn_s_dcache_inv_vol' needs target feature ci-insts}}
__builtin_amdgcn_buffer_wbinvl1_vol(); // expected-error {{'__builtin_amdgcn_buffer_wbinvl1_vol' needs target feature ci-insts}}
} }