This is an archive of the discontinued LLVM Phabricator instance.

Differential D50320

AMDGPU: Add builtin for s_dcache_inv_vol
ClosedPublic

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

Details

Reviewers
kzhuravl
yaxunl

Diff Detail

Event Timeline

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

LGTM

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

r339109

Revision Contents

PathSize
include/
clang/
Basic/
5 lines
test/
CodeGenOpenCL/
10 lines
SemaOpenCL/
7 lines

Diff 159252

include/clang/Basic/BuiltinsAMDGPU.def

Show First 20 Lines • Show All 95 Lines • ▼ Show 20 Lines
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_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.
//===----------------------------------------------------------------------===//
TARGET_BUILTIN(__builtin_amdgcn_s_dcache_inv_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")
TARGET_BUILTIN(__builtin_amdgcn_sinh, "hh", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_sinh, "hh", "nc", "16-bit-insts")
TARGET_BUILTIN(__builtin_amdgcn_cosh, "hh", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_cosh, "hh", "nc", "16-bit-insts")
▲ Show 20 Lines • Show All 52 Lines • Show Last 20 Lines

test/CodeGenOpenCL/builtins-amdgcn-ci.cl

  • This file was added.
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu hawaii -S -emit-llvm -o - %s | FileCheck %s
// CHECK-LABEL: @test_s_dcache_inv_vol
// CHECK: call void @llvm.amdgcn.s.dcache.inv.vol(
void test_s_dcache_inv_vol()
{
__builtin_amdgcn_s_dcache_inv_vol();
}

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

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