1 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
2 // REQUIRES: amdgpu-registered-target
4 // RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
5 // RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \
6 // RUN: -internal-isystem %S/Inputs/include \
7 // RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
8 // RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -O1 -o - \
9 // RUN: -D__HIPCC_RTC__ | FileCheck -check-prefix=DEFAULT %s
11 // Check that we end up with fast math flags set on intrinsic calls
12 // RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
13 // RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \
14 // RUN: -internal-isystem %S/Inputs/include \
15 // RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
16 // RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -O1 -menable-no-infs \
17 // RUN: -menable-no-nans -o - \
18 // RUN: -D__HIPCC_RTC__ | FileCheck -check-prefix=FINITEONLY %s
20 // DEFAULT-LABEL: @test_fma_f16(
21 // DEFAULT-NEXT: entry:
22 // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef half @llvm.fma.f16(half [[X:%.*]], half [[Y:%.*]], half [[Z:%.*]])
23 // DEFAULT-NEXT: ret half [[TMP0]]
25 // FINITEONLY-LABEL: @test_fma_f16(
26 // FINITEONLY-NEXT: entry:
27 // FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef half @llvm.fma.f16(half [[X:%.*]], half [[Y:%.*]], half [[Z:%.*]])
28 // FINITEONLY-NEXT: ret half [[TMP0]]
30 extern "C" __device__ _Float16 test_fma_f16(_Float16 x, _Float16 y,
35 // DEFAULT-LABEL: @test_pow_f16(
36 // DEFAULT-NEXT: entry:
37 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract noundef half @__ocml_pown_f16(half noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR7:[0-9]+]]
38 // DEFAULT-NEXT: ret half [[CALL_I]]
40 // FINITEONLY-LABEL: @test_pow_f16(
41 // FINITEONLY-NEXT: entry:
42 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) half @__ocml_pown_f16(half noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR7:[0-9]+]]
43 // FINITEONLY-NEXT: ret half [[CALL_I]]
45 extern "C" __device__ _Float16 test_pow_f16(_Float16 x, int y) {
49 // DEFAULT-LABEL: @test_fabs_f32(
50 // DEFAULT-NEXT: entry:
51 // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.fabs.f32(float [[X:%.*]])
52 // DEFAULT-NEXT: ret float [[TMP0]]
54 // FINITEONLY-LABEL: @test_fabs_f32(
55 // FINITEONLY-NEXT: entry:
56 // FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.fabs.f32(float [[X:%.*]])
57 // FINITEONLY-NEXT: ret float [[TMP0]]
59 extern "C" __device__ float test_fabs_f32(float x) {
63 // DEFAULT-LABEL: @test_sin_f32(
64 // DEFAULT-NEXT: entry:
65 // DEFAULT-NEXT: [[CALL_I1:%.*]] = tail call contract noundef float @__ocml_sin_f32(float noundef [[X:%.*]]) #[[ATTR8:[0-9]+]]
66 // DEFAULT-NEXT: ret float [[CALL_I1]]
68 // FINITEONLY-LABEL: @test_sin_f32(
69 // FINITEONLY-NEXT: entry:
70 // FINITEONLY-NEXT: [[CALL_I1:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_sin_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR8:[0-9]+]]
71 // FINITEONLY-NEXT: ret float [[CALL_I1]]
73 extern "C" __device__ float test_sin_f32(float x) {
77 // DEFAULT-LABEL: @test_cos_f32(
78 // DEFAULT-NEXT: entry:
79 // DEFAULT-NEXT: [[CALL_I1:%.*]] = tail call contract noundef float @__ocml_cos_f32(float noundef [[X:%.*]]) #[[ATTR8]]
80 // DEFAULT-NEXT: ret float [[CALL_I1]]
82 // FINITEONLY-LABEL: @test_cos_f32(
83 // FINITEONLY-NEXT: entry:
84 // FINITEONLY-NEXT: [[CALL_I1:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_cos_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR8]]
85 // FINITEONLY-NEXT: ret float [[CALL_I1]]
87 extern "C" __device__ float test_cos_f32(float x) {
91 // Check user defined type which can be converted to float and double but not
92 // specializes __numeric_type will not cause ambiguity diagnostics.
93 struct user_bfloat16 {
94 __host__ __device__ user_bfloat16(float);
99 namespace user_namespace {
100 __device__ user_bfloat16 fma(const user_bfloat16 a, const user_bfloat16 b, const user_bfloat16 c) {
104 __global__ void test_fma() {
105 user_bfloat16 a = 1.0f, b = 2.0f;