1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
|
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
// RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \
// RUN: -internal-isystem %S/Inputs/include \
// RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
// RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -O1 -o - \
// RUN: -D__HIPCC_RTC__ | FileCheck %s
// Test deprecated functions in the header that should be removed eventually
// CHECK-LABEL: @test_rcpf16_wrapper(
// CHECK-NEXT: entry:
// CHECK-NEXT: [[DIV_I:%.*]] = fdiv contract half 0xH3C00, [[X:%.*]]
// CHECK-NEXT: ret half [[DIV_I]]
//
extern "C" __device__ _Float16 test_rcpf16_wrapper(_Float16 x) {
return __llvm_amdgcn_rcp_f16(x);
}
// CHECK-LABEL: @test_rcp2f16_wrapper(
// CHECK-NEXT: entry:
// CHECK-NEXT: [[DIV_I:%.*]] = fdiv contract <2 x half> <half 0xH3C00, half 0xH3C00>, [[X:%.*]]
// CHECK-NEXT: ret <2 x half> [[DIV_I]]
//
extern "C" __device__ __2f16 test_rcp2f16_wrapper(__2f16 x) {
return __llvm_amdgcn_rcp_2f16(x);
}
|
