mirror of
https://github.com/llvm/llvm-project.git
synced 2025-04-19 07:46:49 +00:00
[Inliner] Also propagate noundef and align ret attributes during inlining
Both of these can potentially be lost otherwise.
This commit is contained in:
Noah Goldstein
parent
2d037f5aed
commit
2da4960f20
@ -1,9 +1,9 @@
|
||||
// RUN: %clang_cc1 -x c -ffreestanding %s -triple=x86_64-unknown-unknown -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -x c++ -std=c++11 -ffreestanding %s -triple=x86_64-unknown-unknown -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -x c -ffreestanding %s -triple=x86_64-unknown-unknown -no-enable-noundef-analysis -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -x c++ -std=c++11 -ffreestanding %s -triple=x86_64-unknown-unknown -no-enable-noundef-analysis -emit-llvm -o - | FileCheck %s
|
||||
|
||||
// PR33722
|
||||
// RUN: %clang_cc1 -x c -ffreestanding %s -triple x86_64-unknown-unknown -fms-extensions -fms-compatibility-version=19.00 -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -x c++ -ffreestanding %s -triple x86_64-unknown-unknown -fms-extensions -fms-compatibility-version=19.00 -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -x c -ffreestanding %s -triple x86_64-unknown-unknown -fms-extensions -fms-compatibility-version=19.00 -no-enable-noundef-analysis -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -x c++ -ffreestanding %s -triple x86_64-unknown-unknown -fms-extensions -fms-compatibility-version=19.00 -no-enable-noundef-analysis -emit-llvm -o - | FileCheck %s
|
||||
|
||||
#include <x86intrin.h>
|
||||
|
||||
|
||||
@ -1,7 +1,7 @@
|
||||
// RUN: %clang_cc1 -x c -ffreestanding %s -triple=x86_64-apple-darwin -target-feature +popcnt -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK-POPCNT
|
||||
// RUN: %clang_cc1 -x c++ -std=c++11 -ffreestanding %s -triple=x86_64-apple-darwin -target-feature +popcnt -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK-POPCNT
|
||||
// RUN: %clang_cc1 -x c -ffreestanding %s -triple=x86_64-apple-darwin -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -x c++ -std=c++11 -ffreestanding %s -triple=x86_64-apple-darwin -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -x c -ffreestanding %s -triple=x86_64-apple-darwin -target-feature +popcnt -no-enable-noundef-analysis -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK-POPCNT
|
||||
// RUN: %clang_cc1 -x c++ -std=c++11 -ffreestanding %s -triple=x86_64-apple-darwin -target-feature +popcnt -no-enable-noundef-analysis -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK-POPCNT
|
||||
// RUN: %clang_cc1 -x c -ffreestanding %s -triple=x86_64-apple-darwin -no-enable-noundef-analysis -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -x c++ -std=c++11 -ffreestanding %s -triple=x86_64-apple-darwin -no-enable-noundef-analysis -emit-llvm -o - | FileCheck %s
|
||||
|
||||
#include <x86intrin.h>
|
||||
|
||||
|
||||
@ -1,16 +1,16 @@
|
||||
// RUN: %clang_cc1 -x c -ffreestanding -triple i686--linux -emit-llvm %s -o - | FileCheck %s --check-prefixes CHECK,CHECK-32BIT-LONG
|
||||
// RUN: %clang_cc1 -x c -ffreestanding -triple x86_64--linux -emit-llvm %s -o - | FileCheck %s --check-prefixes CHECK,CHECK-64BIT-LONG
|
||||
// RUN: %clang_cc1 -x c -fms-extensions -fms-compatibility -ffreestanding %s -triple=i686-windows-msvc -target-feature +sse2 -emit-llvm -o - -Wall -Werror | FileCheck %s --check-prefixes CHECK,CHECK-32BIT-LONG
|
||||
// RUN: %clang_cc1 -x c -fms-extensions -fms-compatibility -ffreestanding %s -triple=x86_64-windows-msvc -target-feature +sse2 -emit-llvm -o - -Wall -Werror | FileCheck %s --check-prefixes CHECK,CHECK-32BIT-LONG
|
||||
// RUN: %clang_cc1 -x c -fms-extensions -fms-compatibility -fms-compatibility-version=17.00 -ffreestanding %s -triple=i686-windows-msvc -target-feature +sse2 -emit-llvm -o - -Wall -Werror | FileCheck %s --check-prefixes CHECK,CHECK-32BIT-LONG
|
||||
// RUN: %clang_cc1 -x c -fms-extensions -fms-compatibility -fms-compatibility-version=17.00 -ffreestanding %s -triple=x86_64-windows-msvc -target-feature +sse2 -emit-llvm -o - -Wall -Werror | FileCheck %s --check-prefixes CHECK,CHECK-32BIT-LONG
|
||||
// RUN: %clang_cc1 -x c -ffreestanding -triple i686--linux -no-enable-noundef-analysis -emit-llvm %s -o - | FileCheck %s --check-prefixes CHECK,CHECK-32BIT-LONG
|
||||
// RUN: %clang_cc1 -x c -ffreestanding -triple x86_64--linux -no-enable-noundef-analysis -emit-llvm %s -o - | FileCheck %s --check-prefixes CHECK,CHECK-64BIT-LONG
|
||||
// RUN: %clang_cc1 -x c -fms-extensions -fms-compatibility -ffreestanding %s -triple=i686-windows-msvc -target-feature +sse2 -no-enable-noundef-analysis -emit-llvm -o - -Wall -Werror | FileCheck %s --check-prefixes CHECK,CHECK-32BIT-LONG
|
||||
// RUN: %clang_cc1 -x c -fms-extensions -fms-compatibility -ffreestanding %s -triple=x86_64-windows-msvc -target-feature +sse2 -no-enable-noundef-analysis -emit-llvm -o - -Wall -Werror | FileCheck %s --check-prefixes CHECK,CHECK-32BIT-LONG
|
||||
// RUN: %clang_cc1 -x c -fms-extensions -fms-compatibility -fms-compatibility-version=17.00 -ffreestanding %s -triple=i686-windows-msvc -target-feature +sse2 -no-enable-noundef-analysis -emit-llvm -o - -Wall -Werror | FileCheck %s --check-prefixes CHECK,CHECK-32BIT-LONG
|
||||
// RUN: %clang_cc1 -x c -fms-extensions -fms-compatibility -fms-compatibility-version=17.00 -ffreestanding %s -triple=x86_64-windows-msvc -target-feature +sse2 -no-enable-noundef-analysis -emit-llvm -o - -Wall -Werror | FileCheck %s --check-prefixes CHECK,CHECK-32BIT-LONG
|
||||
|
||||
// RUN: %clang_cc1 -x c++ -std=c++11 -ffreestanding -triple i686--linux -emit-llvm %s -o - | FileCheck %s --check-prefixes CHECK,CHECK-32BIT-LONG
|
||||
// RUN: %clang_cc1 -x c++ -std=c++11 -ffreestanding -triple x86_64--linux -emit-llvm %s -o - | FileCheck %s --check-prefixes CHECK,CHECK-64BIT-LONG
|
||||
// RUN: %clang_cc1 -x c++ -std=c++11 -fms-extensions -fms-compatibility -ffreestanding %s -triple=i686-windows-msvc -target-feature +sse2 -emit-llvm -o - -Wall -Werror | FileCheck %s --check-prefixes CHECK,CHECK-32BIT-LONG
|
||||
// RUN: %clang_cc1 -x c++ -std=c++11 -fms-extensions -fms-compatibility -ffreestanding %s -triple=x86_64-windows-msvc -target-feature +sse2 -emit-llvm -o - -Wall -Werror | FileCheck %s --check-prefixes CHECK,CHECK-32BIT-LONG
|
||||
// RUN: %clang_cc1 -x c++ -std=c++11 -fms-extensions -fms-compatibility -fms-compatibility-version=17.00 -ffreestanding %s -triple=i686-windows-msvc -target-feature +sse2 -emit-llvm -o - -Wall -Werror | FileCheck %s --check-prefixes CHECK,CHECK-32BIT-LONG
|
||||
// RUN: %clang_cc1 -x c++ -std=c++11 -fms-extensions -fms-compatibility -fms-compatibility-version=17.00 -ffreestanding %s -triple=x86_64-windows-msvc -target-feature +sse2 -emit-llvm -o - -Wall -Werror | FileCheck %s --check-prefixes CHECK,CHECK-32BIT-LONG
|
||||
// RUN: %clang_cc1 -x c++ -std=c++11 -ffreestanding -triple i686--linux -no-enable-noundef-analysis -emit-llvm %s -o - | FileCheck %s --check-prefixes CHECK,CHECK-32BIT-LONG
|
||||
// RUN: %clang_cc1 -x c++ -std=c++11 -ffreestanding -triple x86_64--linux -no-enable-noundef-analysis -emit-llvm %s -o - | FileCheck %s --check-prefixes CHECK,CHECK-64BIT-LONG
|
||||
// RUN: %clang_cc1 -x c++ -std=c++11 -fms-extensions -fms-compatibility -ffreestanding %s -triple=i686-windows-msvc -target-feature +sse2 -no-enable-noundef-analysis -emit-llvm -o - -Wall -Werror | FileCheck %s --check-prefixes CHECK,CHECK-32BIT-LONG
|
||||
// RUN: %clang_cc1 -x c++ -std=c++11 -fms-extensions -fms-compatibility -ffreestanding %s -triple=x86_64-windows-msvc -target-feature +sse2 -no-enable-noundef-analysis -emit-llvm -o - -Wall -Werror | FileCheck %s --check-prefixes CHECK,CHECK-32BIT-LONG
|
||||
// RUN: %clang_cc1 -x c++ -std=c++11 -fms-extensions -fms-compatibility -fms-compatibility-version=17.00 -ffreestanding %s -triple=i686-windows-msvc -target-feature +sse2 -no-enable-noundef-analysis -emit-llvm -o - -Wall -Werror | FileCheck %s --check-prefixes CHECK,CHECK-32BIT-LONG
|
||||
// RUN: %clang_cc1 -x c++ -std=c++11 -fms-extensions -fms-compatibility -fms-compatibility-version=17.00 -ffreestanding %s -triple=x86_64-windows-msvc -target-feature +sse2 -no-enable-noundef-analysis -emit-llvm -o - -Wall -Werror | FileCheck %s --check-prefixes CHECK,CHECK-32BIT-LONG
|
||||
|
||||
#include <x86intrin.h>
|
||||
|
||||
|
||||
@ -1,5 +1,5 @@
|
||||
// RUN: %clang_cc1 -x c -ffreestanding %s -triple=x86_64-apple-darwin -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -x c++ -std=c++11 -ffreestanding %s -triple=x86_64-apple-darwin -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -x c -ffreestanding %s -triple=x86_64-apple-darwin -no-enable-noundef-analysis -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -x c++ -std=c++11 -ffreestanding %s -triple=x86_64-apple-darwin -no-enable-noundef-analysis -emit-llvm -o - | FileCheck %s
|
||||
|
||||
#include <x86intrin.h>
|
||||
|
||||
|
||||
@ -206,7 +206,7 @@ EXTERN_C void test_st64b(void)
|
||||
// CHECK-CXX-NEXT: [[TMP14:%.*]] = load i64, ptr [[TMP13]], align 8
|
||||
// CHECK-CXX-NEXT: [[TMP15:%.*]] = getelementptr i64, ptr [[AGG_TMP]], i32 7
|
||||
// CHECK-CXX-NEXT: [[TMP16:%.*]] = load i64, ptr [[TMP15]], align 8
|
||||
// CHECK-CXX-NEXT: [[TMP17:%.*]] = call i64 @llvm.aarch64.st64bv(ptr [[TMP1]], i64 [[TMP2]], i64 [[TMP4]], i64 [[TMP6]], i64 [[TMP8]], i64 [[TMP10]], i64 [[TMP12]], i64 [[TMP14]], i64 [[TMP16]])
|
||||
// CHECK-CXX-NEXT: [[TMP17:%.*]] = call noundef i64 @llvm.aarch64.st64bv(ptr [[TMP1]], i64 [[TMP2]], i64 [[TMP4]], i64 [[TMP6]], i64 [[TMP8]], i64 [[TMP10]], i64 [[TMP12]], i64 [[TMP14]], i64 [[TMP16]])
|
||||
// CHECK-CXX-NEXT: store i64 [[TMP17]], ptr @status, align 8
|
||||
// CHECK-CXX-NEXT: ret void
|
||||
//
|
||||
@ -269,7 +269,7 @@ EXTERN_C void test_st64bv(void)
|
||||
// CHECK-CXX-NEXT: [[TMP14:%.*]] = load i64, ptr [[TMP13]], align 8
|
||||
// CHECK-CXX-NEXT: [[TMP15:%.*]] = getelementptr i64, ptr [[AGG_TMP]], i32 7
|
||||
// CHECK-CXX-NEXT: [[TMP16:%.*]] = load i64, ptr [[TMP15]], align 8
|
||||
// CHECK-CXX-NEXT: [[TMP17:%.*]] = call i64 @llvm.aarch64.st64bv0(ptr [[TMP1]], i64 [[TMP2]], i64 [[TMP4]], i64 [[TMP6]], i64 [[TMP8]], i64 [[TMP10]], i64 [[TMP12]], i64 [[TMP14]], i64 [[TMP16]])
|
||||
// CHECK-CXX-NEXT: [[TMP17:%.*]] = call noundef i64 @llvm.aarch64.st64bv0(ptr [[TMP1]], i64 [[TMP2]], i64 [[TMP4]], i64 [[TMP6]], i64 [[TMP8]], i64 [[TMP10]], i64 [[TMP12]], i64 [[TMP14]], i64 [[TMP16]])
|
||||
// CHECK-CXX-NEXT: store i64 [[TMP17]], ptr @status, align 8
|
||||
// CHECK-CXX-NEXT: ret void
|
||||
//
|
||||
|
||||
@ -46,7 +46,7 @@ float fp_contract_3(float a, float b, float c) {
|
||||
// CHECK: %[[M:.+]] = fmul contract float %a, %b
|
||||
// CHECK-NEXT: fadd contract float %[[M]], %c
|
||||
// STRICT: %[[M:.+]] = tail call contract float @llvm.experimental.constrained.fmul.f32(float %a, float %b, metadata !"round.tonearest", metadata !"fpexcept.strict")
|
||||
// STRICT-NEXT: tail call contract float @llvm.experimental.constrained.fadd.f32(float %[[M]], float %c, metadata !"round.tonearest", metadata !"fpexcept.strict")
|
||||
// STRICT-NEXT: tail call contract noundef float @llvm.experimental.constrained.fadd.f32(float %[[M]], float %c, metadata !"round.tonearest", metadata !"fpexcept.strict")
|
||||
return template_muladd<float>(a, b, c);
|
||||
}
|
||||
|
||||
|
||||
@ -31,7 +31,7 @@ T template_muladd(T a, T b, T c) {
|
||||
|
||||
float fp_contract_3(float a, float b, float c) {
|
||||
// CHECK: _Z13fp_contract_3fff
|
||||
// CHECK: tail call float @llvm.fmuladd
|
||||
// CHECK: tail call noundef float @llvm.fmuladd
|
||||
return template_muladd<float>(a, b, c);
|
||||
}
|
||||
|
||||
|
||||
@ -31,7 +31,7 @@ T template_muladd(T a, T b, T c) {
|
||||
|
||||
float fp_contract_3(float a, float b, float c) {
|
||||
// CHECK: _Z13fp_contract_3fff
|
||||
// CHECK: tail call float @llvm.fmuladd
|
||||
// CHECK: tail call noundef float @llvm.fmuladd
|
||||
return template_muladd<float>(a, b, c);
|
||||
}
|
||||
|
||||
|
||||
@ -6,21 +6,21 @@
|
||||
__attribute__((global))
|
||||
void kernel(int *out) {
|
||||
int i = 0;
|
||||
out[i++] = threadIdx.x; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
|
||||
out[i++] = threadIdx.y; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.y()
|
||||
out[i++] = threadIdx.z; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.z()
|
||||
out[i++] = threadIdx.x; // CHECK: call noundef i32 @llvm.nvvm.read.ptx.sreg.tid.x()
|
||||
out[i++] = threadIdx.y; // CHECK: call noundef i32 @llvm.nvvm.read.ptx.sreg.tid.y()
|
||||
out[i++] = threadIdx.z; // CHECK: call noundef i32 @llvm.nvvm.read.ptx.sreg.tid.z()
|
||||
|
||||
out[i++] = blockIdx.x; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
|
||||
out[i++] = blockIdx.y; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()
|
||||
out[i++] = blockIdx.z; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.z()
|
||||
out[i++] = blockIdx.x; // CHECK: call noundef i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
|
||||
out[i++] = blockIdx.y; // CHECK: call noundef i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()
|
||||
out[i++] = blockIdx.z; // CHECK: call noundef i32 @llvm.nvvm.read.ptx.sreg.ctaid.z()
|
||||
|
||||
out[i++] = blockDim.x; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
||||
out[i++] = blockDim.y; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.y()
|
||||
out[i++] = blockDim.z; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.z()
|
||||
out[i++] = blockDim.x; // CHECK: call noundef i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
|
||||
out[i++] = blockDim.y; // CHECK: call noundef i32 @llvm.nvvm.read.ptx.sreg.ntid.y()
|
||||
out[i++] = blockDim.z; // CHECK: call noundef i32 @llvm.nvvm.read.ptx.sreg.ntid.z()
|
||||
|
||||
out[i++] = gridDim.x; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()
|
||||
out[i++] = gridDim.y; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.y()
|
||||
out[i++] = gridDim.z; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.z()
|
||||
out[i++] = gridDim.x; // CHECK: call noundef i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()
|
||||
out[i++] = gridDim.y; // CHECK: call noundef i32 @llvm.nvvm.read.ptx.sreg.nctaid.y()
|
||||
out[i++] = gridDim.z; // CHECK: call noundef i32 @llvm.nvvm.read.ptx.sreg.nctaid.z()
|
||||
|
||||
out[i++] = warpSize; // CHECK: store i32 32,
|
||||
|
||||
|
||||
@ -18,12 +18,12 @@
|
||||
|
||||
// DEFAULT-LABEL: @test_fma_f16(
|
||||
// DEFAULT-NEXT: entry:
|
||||
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract half @llvm.fma.f16(half [[X:%.*]], half [[Y:%.*]], half [[Z:%.*]])
|
||||
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef half @llvm.fma.f16(half [[X:%.*]], half [[Y:%.*]], half [[Z:%.*]])
|
||||
// DEFAULT-NEXT: ret half [[TMP0]]
|
||||
//
|
||||
// FINITEONLY-LABEL: @test_fma_f16(
|
||||
// FINITEONLY-NEXT: entry:
|
||||
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract half @llvm.fma.f16(half [[X:%.*]], half [[Y:%.*]], half [[Z:%.*]])
|
||||
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef half @llvm.fma.f16(half [[X:%.*]], half [[Y:%.*]], half [[Z:%.*]])
|
||||
// FINITEONLY-NEXT: ret half [[TMP0]]
|
||||
//
|
||||
extern "C" __device__ _Float16 test_fma_f16(_Float16 x, _Float16 y,
|
||||
@ -33,12 +33,12 @@ extern "C" __device__ _Float16 test_fma_f16(_Float16 x, _Float16 y,
|
||||
|
||||
// DEFAULT-LABEL: @test_pow_f16(
|
||||
// DEFAULT-NEXT: entry:
|
||||
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract half @__ocml_pown_f16(half noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR7:[0-9]+]]
|
||||
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract noundef half @__ocml_pown_f16(half noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR7:[0-9]+]]
|
||||
// DEFAULT-NEXT: ret half [[CALL_I]]
|
||||
//
|
||||
// FINITEONLY-LABEL: @test_pow_f16(
|
||||
// FINITEONLY-NEXT: entry:
|
||||
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) half @__ocml_pown_f16(half noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR7:[0-9]+]]
|
||||
// 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]+]]
|
||||
// FINITEONLY-NEXT: ret half [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ _Float16 test_pow_f16(_Float16 x, int y) {
|
||||
@ -47,12 +47,12 @@ extern "C" __device__ _Float16 test_pow_f16(_Float16 x, int y) {
|
||||
|
||||
// DEFAULT-LABEL: @test_fabs_f32(
|
||||
// DEFAULT-NEXT: entry:
|
||||
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.fabs.f32(float [[X:%.*]])
|
||||
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.fabs.f32(float [[X:%.*]])
|
||||
// DEFAULT-NEXT: ret float [[TMP0]]
|
||||
//
|
||||
// FINITEONLY-LABEL: @test_fabs_f32(
|
||||
// FINITEONLY-NEXT: entry:
|
||||
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.fabs.f32(float [[X:%.*]])
|
||||
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.fabs.f32(float [[X:%.*]])
|
||||
// FINITEONLY-NEXT: ret float [[TMP0]]
|
||||
//
|
||||
extern "C" __device__ float test_fabs_f32(float x) {
|
||||
@ -61,12 +61,12 @@ extern "C" __device__ float test_fabs_f32(float x) {
|
||||
|
||||
// DEFAULT-LABEL: @test_sin_f32(
|
||||
// DEFAULT-NEXT: entry:
|
||||
// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_sin_f32(float noundef [[X:%.*]]) #[[ATTR8:[0-9]+]]
|
||||
// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_sin_f32(float noundef [[X:%.*]]) #[[ATTR8:[0-9]+]]
|
||||
// DEFAULT-NEXT: ret float [[CALL_I_I]]
|
||||
//
|
||||
// FINITEONLY-LABEL: @test_sin_f32(
|
||||
// FINITEONLY-NEXT: entry:
|
||||
// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_sin_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR8:[0-9]+]]
|
||||
// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_sin_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR8:[0-9]+]]
|
||||
// FINITEONLY-NEXT: ret float [[CALL_I_I]]
|
||||
//
|
||||
extern "C" __device__ float test_sin_f32(float x) {
|
||||
@ -75,12 +75,12 @@ extern "C" __device__ float test_sin_f32(float x) {
|
||||
|
||||
// DEFAULT-LABEL: @test_cos_f32(
|
||||
// DEFAULT-NEXT: entry:
|
||||
// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_cos_f32(float noundef [[X:%.*]]) #[[ATTR8]]
|
||||
// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_cos_f32(float noundef [[X:%.*]]) #[[ATTR8]]
|
||||
// DEFAULT-NEXT: ret float [[CALL_I_I]]
|
||||
//
|
||||
// FINITEONLY-LABEL: @test_cos_f32(
|
||||
// FINITEONLY-NEXT: entry:
|
||||
// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_cos_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR8]]
|
||||
// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_cos_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR8]]
|
||||
// FINITEONLY-NEXT: ret float [[CALL_I_I]]
|
||||
//
|
||||
extern "C" __device__ float test_cos_f32(float x) {
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
@ -9,7 +9,7 @@
|
||||
|
||||
// CHECK-LABEL: @test___fadd_rd(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_add_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3:[0-9]+]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_add_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3:[0-9]+]]
|
||||
// CHECK-NEXT: ret float [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ float test___fadd_rd(float x, float y) {
|
||||
@ -18,7 +18,7 @@ extern "C" __device__ float test___fadd_rd(float x, float y) {
|
||||
|
||||
// CHECK-LABEL: @test___fadd_rn(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_add_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_add_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret float [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ float test___fadd_rn(float x, float y) {
|
||||
@ -27,7 +27,7 @@ extern "C" __device__ float test___fadd_rn(float x, float y) {
|
||||
|
||||
// CHECK-LABEL: @test___fadd_ru(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_add_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_add_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret float [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ float test___fadd_ru(float x, float y) {
|
||||
@ -36,7 +36,7 @@ extern "C" __device__ float test___fadd_ru(float x, float y) {
|
||||
|
||||
// CHECK-LABEL: @test___fadd_rz(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_add_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_add_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret float [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ float test___fadd_rz(float x, float y) {
|
||||
@ -45,7 +45,7 @@ extern "C" __device__ float test___fadd_rz(float x, float y) {
|
||||
|
||||
// CHECK-LABEL: @test__fmaf_rd(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_fma_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret float [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ float test__fmaf_rd(float x, float y, float z) {
|
||||
@ -54,7 +54,7 @@ extern "C" __device__ float test__fmaf_rd(float x, float y, float z) {
|
||||
|
||||
// CHECK-LABEL: @test__fmaf_rn(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_fma_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret float [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ float test__fmaf_rn(float x, float y, float z) {
|
||||
@ -63,7 +63,7 @@ extern "C" __device__ float test__fmaf_rn(float x, float y, float z) {
|
||||
|
||||
// CHECK-LABEL: @test__fmaf_ru(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_fma_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret float [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ float test__fmaf_ru(float x, float y, float z) {
|
||||
@ -72,7 +72,7 @@ extern "C" __device__ float test__fmaf_ru(float x, float y, float z) {
|
||||
|
||||
// CHECK-LABEL: @test__fmaf_rz(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_fma_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret float [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ float test__fmaf_rz(float x, float y, float z) {
|
||||
@ -81,7 +81,7 @@ extern "C" __device__ float test__fmaf_rz(float x, float y, float z) {
|
||||
|
||||
// CHECK-LABEL: @test___fmul_rd(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_mul_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_mul_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret float [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ float test___fmul_rd(float x, float y) {
|
||||
@ -90,7 +90,7 @@ extern "C" __device__ float test___fmul_rd(float x, float y) {
|
||||
|
||||
// CHECK-LABEL: @test___fmul_rn(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_mul_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_mul_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret float [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ float test___fmul_rn(float x, float y) {
|
||||
@ -99,7 +99,7 @@ extern "C" __device__ float test___fmul_rn(float x, float y) {
|
||||
|
||||
// CHECK-LABEL: @test___fmul_ru(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_mul_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_mul_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret float [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ float test___fmul_ru(float x, float y) {
|
||||
@ -108,7 +108,7 @@ extern "C" __device__ float test___fmul_ru(float x, float y) {
|
||||
|
||||
// CHECK-LABEL: @test___fmul_rz(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_mul_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_mul_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret float [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ float test___fmul_rz(float x, float y) {
|
||||
@ -117,7 +117,7 @@ extern "C" __device__ float test___fmul_rz(float x, float y) {
|
||||
|
||||
// CHECK-LABEL: @test___frcp_rd(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_div_rtn_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_div_rtn_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret float [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ float test___frcp_rd(float x) {
|
||||
@ -126,7 +126,7 @@ extern "C" __device__ float test___frcp_rd(float x) {
|
||||
|
||||
// CHECK-LABEL: @test___frcp_rn(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_div_rte_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_div_rte_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret float [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ float test___frcp_rn(float x) {
|
||||
@ -135,7 +135,7 @@ extern "C" __device__ float test___frcp_rn(float x) {
|
||||
|
||||
// CHECK-LABEL: @test___frcp_ru(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_div_rtp_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_div_rtp_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret float [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ float test___frcp_ru(float x) {
|
||||
@ -144,7 +144,7 @@ extern "C" __device__ float test___frcp_ru(float x) {
|
||||
|
||||
// CHECK-LABEL: @test___frcp_rz(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_div_rtz_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_div_rtz_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret float [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ float test___frcp_rz(float x) {
|
||||
@ -153,7 +153,7 @@ extern "C" __device__ float test___frcp_rz(float x) {
|
||||
|
||||
// CHECK-LABEL: @test___fsqrt_rd(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_rtn_f32(float noundef [[X:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sqrt_rtn_f32(float noundef [[X:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret float [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ float test___fsqrt_rd(float x) {
|
||||
@ -162,7 +162,7 @@ extern "C" __device__ float test___fsqrt_rd(float x) {
|
||||
|
||||
// CHECK-LABEL: @test___fsqrt_rn(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_rte_f32(float noundef [[X:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sqrt_rte_f32(float noundef [[X:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret float [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ float test___fsqrt_rn(float x) {
|
||||
@ -171,7 +171,7 @@ extern "C" __device__ float test___fsqrt_rn(float x) {
|
||||
|
||||
// CHECK-LABEL: @test___fsqrt_ru(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_rtp_f32(float noundef [[X:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sqrt_rtp_f32(float noundef [[X:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret float [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ float test___fsqrt_ru(float x) {
|
||||
@ -180,7 +180,7 @@ extern "C" __device__ float test___fsqrt_ru(float x) {
|
||||
|
||||
// CHECK-LABEL: @test___fsqrt_rz(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_rtz_f32(float noundef [[X:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sqrt_rtz_f32(float noundef [[X:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret float [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ float test___fsqrt_rz(float x) {
|
||||
@ -189,7 +189,7 @@ extern "C" __device__ float test___fsqrt_rz(float x) {
|
||||
|
||||
// CHECK-LABEL: @test___fsub_rd(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sub_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sub_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret float [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ float test___fsub_rd(float x, float y) {
|
||||
@ -198,7 +198,7 @@ extern "C" __device__ float test___fsub_rd(float x, float y) {
|
||||
|
||||
// CHECK-LABEL: @test___fsub_rn(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sub_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sub_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret float [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ float test___fsub_rn(float x, float y) {
|
||||
@ -207,7 +207,7 @@ extern "C" __device__ float test___fsub_rn(float x, float y) {
|
||||
|
||||
// CHECK-LABEL: @test___fsub_ru(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sub_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sub_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret float [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ float test___fsub_ru(float x, float y) {
|
||||
@ -216,7 +216,7 @@ extern "C" __device__ float test___fsub_ru(float x, float y) {
|
||||
|
||||
// CHECK-LABEL: @test___fsub_rz(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sub_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sub_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret float [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ float test___fsub_rz(float x, float y) {
|
||||
@ -225,7 +225,7 @@ extern "C" __device__ float test___fsub_rz(float x, float y) {
|
||||
|
||||
// CHECK-LABEL: @test___dadd_rd(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_add_rtn_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_add_rtn_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret double [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ double test___dadd_rd(double x, double y) {
|
||||
@ -234,7 +234,7 @@ extern "C" __device__ double test___dadd_rd(double x, double y) {
|
||||
|
||||
// CHECK-LABEL: @test___dadd_rn(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_add_rte_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_add_rte_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret double [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ double test___dadd_rn(double x, double y) {
|
||||
@ -243,7 +243,7 @@ extern "C" __device__ double test___dadd_rn(double x, double y) {
|
||||
|
||||
// CHECK-LABEL: @test___dadd_ru(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_add_rtp_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_add_rtp_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret double [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ double test___dadd_ru(double x, double y) {
|
||||
@ -252,7 +252,7 @@ extern "C" __device__ double test___dadd_ru(double x, double y) {
|
||||
|
||||
// CHECK-LABEL: @test___dadd_rz(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_add_rtz_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_add_rtz_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret double [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ double test___dadd_rz(double x, double y) {
|
||||
@ -261,7 +261,7 @@ extern "C" __device__ double test___dadd_rz(double x, double y) {
|
||||
|
||||
// CHECK-LABEL: @test___dmul_rd(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_mul_rtn_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_mul_rtn_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret double [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ double test___dmul_rd(double x, double y) {
|
||||
@ -270,7 +270,7 @@ extern "C" __device__ double test___dmul_rd(double x, double y) {
|
||||
|
||||
// CHECK-LABEL: @test___dmul_rn(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_mul_rte_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_mul_rte_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret double [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ double test___dmul_rn(double x, double y) {
|
||||
@ -279,7 +279,7 @@ extern "C" __device__ double test___dmul_rn(double x, double y) {
|
||||
|
||||
// CHECK-LABEL: @test___dmul_ru(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_mul_rtp_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_mul_rtp_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret double [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ double test___dmul_ru(double x, double y) {
|
||||
@ -288,7 +288,7 @@ extern "C" __device__ double test___dmul_ru(double x, double y) {
|
||||
|
||||
// CHECK-LABEL: @test___dmul_rz(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_mul_rtz_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_mul_rtz_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret double [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ double test___dmul_rz(double x, double y) {
|
||||
@ -298,7 +298,7 @@ extern "C" __device__ double test___dmul_rz(double x, double y) {
|
||||
// CHECK-LABEL: @test___drcp_rd(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CONV:%.*]] = fpext float [[X:%.*]] to double
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_div_rtn_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_div_rtn_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
|
||||
// CHECK-NEXT: ret float [[CONV1]]
|
||||
//
|
||||
@ -309,7 +309,7 @@ extern "C" __device__ float test___drcp_rd(float x) {
|
||||
// CHECK-LABEL: @test___drcp_rn(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CONV:%.*]] = fpext float [[X:%.*]] to double
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_div_rte_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_div_rte_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
|
||||
// CHECK-NEXT: ret float [[CONV1]]
|
||||
//
|
||||
@ -320,7 +320,7 @@ extern "C" __device__ float test___drcp_rn(float x) {
|
||||
// CHECK-LABEL: @test___drcp_ru(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CONV:%.*]] = fpext float [[X:%.*]] to double
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_div_rtp_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_div_rtp_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
|
||||
// CHECK-NEXT: ret float [[CONV1]]
|
||||
//
|
||||
@ -331,7 +331,7 @@ extern "C" __device__ float test___drcp_ru(float x) {
|
||||
// CHECK-LABEL: @test___drcp_rz(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CONV:%.*]] = fpext float [[X:%.*]] to double
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_div_rtz_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_div_rtz_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
|
||||
// CHECK-NEXT: ret float [[CONV1]]
|
||||
//
|
||||
@ -342,7 +342,7 @@ extern "C" __device__ float test___drcp_rz(float x) {
|
||||
// CHECK-LABEL: @test___dsqrt_rd(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CONV:%.*]] = fpext float [[X:%.*]] to double
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_rtn_f64(double noundef [[CONV]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_sqrt_rtn_f64(double noundef [[CONV]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
|
||||
// CHECK-NEXT: ret float [[CONV1]]
|
||||
//
|
||||
@ -353,7 +353,7 @@ extern "C" __device__ float test___dsqrt_rd(float x) {
|
||||
// CHECK-LABEL: @test___dsqrt_rn(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CONV:%.*]] = fpext float [[X:%.*]] to double
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_rte_f64(double noundef [[CONV]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_sqrt_rte_f64(double noundef [[CONV]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
|
||||
// CHECK-NEXT: ret float [[CONV1]]
|
||||
//
|
||||
@ -364,7 +364,7 @@ extern "C" __device__ float test___dsqrt_rn(float x) {
|
||||
// CHECK-LABEL: @test___dsqrt_ru(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CONV:%.*]] = fpext float [[X:%.*]] to double
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_rtp_f64(double noundef [[CONV]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_sqrt_rtp_f64(double noundef [[CONV]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
|
||||
// CHECK-NEXT: ret float [[CONV1]]
|
||||
//
|
||||
@ -375,7 +375,7 @@ extern "C" __device__ float test___dsqrt_ru(float x) {
|
||||
// CHECK-LABEL: @test___dsqrt_rz(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CONV:%.*]] = fpext float [[X:%.*]] to double
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_rtz_f64(double noundef [[CONV]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_sqrt_rtz_f64(double noundef [[CONV]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
|
||||
// CHECK-NEXT: ret float [[CONV1]]
|
||||
//
|
||||
@ -385,7 +385,7 @@ extern "C" __device__ float test___dsqrt_rz(float x) {
|
||||
|
||||
// CHECK-LABEL: @test__fma_rd(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_rtn_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_fma_rtn_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret double [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ double test__fma_rd(double x, double y, double z) {
|
||||
@ -394,7 +394,7 @@ extern "C" __device__ double test__fma_rd(double x, double y, double z) {
|
||||
|
||||
// CHECK-LABEL: @test__fma_rn(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_rte_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_fma_rte_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret double [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ double test__fma_rn(double x, double y, double z) {
|
||||
@ -403,7 +403,7 @@ extern "C" __device__ double test__fma_rn(double x, double y, double z) {
|
||||
|
||||
// CHECK-LABEL: @test__fma_ru(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_rtp_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_fma_rtp_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret double [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ double test__fma_ru(double x, double y, double z) {
|
||||
@ -412,7 +412,7 @@ extern "C" __device__ double test__fma_ru(double x, double y, double z) {
|
||||
|
||||
// CHECK-LABEL: @test__fma_rz(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_rtz_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_fma_rtz_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: ret double [[CALL_I]]
|
||||
//
|
||||
extern "C" __device__ double test__fma_rz(double x, double y, double z) {
|
||||
|
||||
@ -113,21 +113,21 @@
|
||||
// CHECK-CPP-NEXT: [[__X_ADDR_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I]] to ptr
|
||||
// CHECK-CPP-NEXT: store double [[TMP0]], ptr [[__X_ADDR_ASCAST_I]], align 8
|
||||
// CHECK-CPP-NEXT: [[TMP1:%.*]] = load double, ptr [[__X_ADDR_ASCAST_I]], align 8
|
||||
// CHECK-CPP-NEXT: [[CALL_I:%.*]] = call double @__ocml_sin_f64(double noundef [[TMP1]]) #[[ATTR3:[0-9]+]]
|
||||
// CHECK-CPP-NEXT: [[CALL_I:%.*]] = call noundef double @__ocml_sin_f64(double noundef [[TMP1]]) #[[ATTR3:[0-9]+]]
|
||||
// CHECK-CPP-NEXT: store double [[CALL_I]], ptr [[L1_ASCAST]], align 8
|
||||
// CHECK-CPP-NEXT: [[TMP2:%.*]] = load double, ptr [[X_ADDR_ASCAST]], align 8
|
||||
// CHECK-CPP-NEXT: [[RETVAL_ASCAST_I6:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I4]] to ptr
|
||||
// CHECK-CPP-NEXT: [[__X_ADDR_ASCAST_I7:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I5]] to ptr
|
||||
// CHECK-CPP-NEXT: store double [[TMP2]], ptr [[__X_ADDR_ASCAST_I7]], align 8
|
||||
// CHECK-CPP-NEXT: [[TMP3:%.*]] = load double, ptr [[__X_ADDR_ASCAST_I7]], align 8
|
||||
// CHECK-CPP-NEXT: [[CALL_I8:%.*]] = call double @__ocml_cos_f64(double noundef [[TMP3]]) #[[ATTR3]]
|
||||
// CHECK-CPP-NEXT: [[CALL_I8:%.*]] = call noundef double @__ocml_cos_f64(double noundef [[TMP3]]) #[[ATTR3]]
|
||||
// CHECK-CPP-NEXT: store double [[CALL_I8]], ptr [[L2_ASCAST]], align 8
|
||||
// CHECK-CPP-NEXT: [[TMP4:%.*]] = load double, ptr [[X_ADDR_ASCAST]], align 8
|
||||
// CHECK-CPP-NEXT: [[RETVAL_ASCAST_I11:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I9]] to ptr
|
||||
// CHECK-CPP-NEXT: [[__X_ADDR_ASCAST_I12:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I10]] to ptr
|
||||
// CHECK-CPP-NEXT: store double [[TMP4]], ptr [[__X_ADDR_ASCAST_I12]], align 8
|
||||
// CHECK-CPP-NEXT: [[TMP5:%.*]] = load double, ptr [[__X_ADDR_ASCAST_I12]], align 8
|
||||
// CHECK-CPP-NEXT: [[TMP6:%.*]] = call double @llvm.fabs.f64(double [[TMP5]])
|
||||
// CHECK-CPP-NEXT: [[TMP6:%.*]] = call noundef double @llvm.fabs.f64(double [[TMP5]])
|
||||
// CHECK-CPP-NEXT: store double [[TMP6]], ptr [[L3_ASCAST]], align 8
|
||||
// CHECK-CPP-NEXT: [[TMP7:%.*]] = load double, ptr [[X_ADDR_ASCAST]], align 8
|
||||
// CHECK-CPP-NEXT: [[TMP8:%.*]] = load double, ptr [[Y_ADDR_ASCAST]], align 8
|
||||
@ -142,7 +142,7 @@
|
||||
// CHECK-CPP-NEXT: [[TMP10:%.*]] = load double, ptr [[__X_ADDR_ASCAST_I16]], align 8
|
||||
// CHECK-CPP-NEXT: [[TMP11:%.*]] = load double, ptr [[__Y_ADDR_ASCAST_I]], align 8
|
||||
// CHECK-CPP-NEXT: [[TMP12:%.*]] = load double, ptr [[__Z_ADDR_ASCAST_I]], align 8
|
||||
// CHECK-CPP-NEXT: [[TMP13:%.*]] = call double @llvm.fma.f64(double [[TMP10]], double [[TMP11]], double [[TMP12]])
|
||||
// CHECK-CPP-NEXT: [[TMP13:%.*]] = call noundef double @llvm.fma.f64(double [[TMP10]], double [[TMP11]], double [[TMP12]])
|
||||
// CHECK-CPP-NEXT: store double [[TMP13]], ptr [[L4_ASCAST]], align 8
|
||||
// CHECK-CPP-NEXT: ret void
|
||||
//
|
||||
@ -278,7 +278,7 @@ void test_math_f64(double x, double y, double z) {
|
||||
// CHECK-CPP-NEXT: [[__X_ADDR_ASCAST_I22:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I20]] to ptr
|
||||
// CHECK-CPP-NEXT: store float [[TMP1]], ptr [[__X_ADDR_ASCAST_I22]], align 4
|
||||
// CHECK-CPP-NEXT: [[TMP2:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I22]], align 4
|
||||
// CHECK-CPP-NEXT: [[CALL_I23:%.*]] = call float @__ocml_sin_f32(float noundef [[TMP2]]) #[[ATTR3]]
|
||||
// CHECK-CPP-NEXT: [[CALL_I23:%.*]] = call noundef float @__ocml_sin_f32(float noundef [[TMP2]]) #[[ATTR3]]
|
||||
// CHECK-CPP-NEXT: store float [[CALL_I23]], ptr [[L1_ASCAST]], align 4
|
||||
// CHECK-CPP-NEXT: [[TMP3:%.*]] = load float, ptr [[X_ADDR_ASCAST]], align 4
|
||||
// CHECK-CPP-NEXT: [[RETVAL_ASCAST_I6:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I4]] to ptr
|
||||
@ -289,7 +289,7 @@ void test_math_f64(double x, double y, double z) {
|
||||
// CHECK-CPP-NEXT: [[__X_ADDR_ASCAST_I27:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I25]] to ptr
|
||||
// CHECK-CPP-NEXT: store float [[TMP4]], ptr [[__X_ADDR_ASCAST_I27]], align 4
|
||||
// CHECK-CPP-NEXT: [[TMP5:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I27]], align 4
|
||||
// CHECK-CPP-NEXT: [[CALL_I:%.*]] = call float @__ocml_cos_f32(float noundef [[TMP5]]) #[[ATTR3]]
|
||||
// CHECK-CPP-NEXT: [[CALL_I:%.*]] = call noundef float @__ocml_cos_f32(float noundef [[TMP5]]) #[[ATTR3]]
|
||||
// CHECK-CPP-NEXT: store float [[CALL_I]], ptr [[L2_ASCAST]], align 4
|
||||
// CHECK-CPP-NEXT: [[TMP6:%.*]] = load float, ptr [[X_ADDR_ASCAST]], align 4
|
||||
// CHECK-CPP-NEXT: [[RETVAL_ASCAST_I11:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I9]] to ptr
|
||||
@ -300,7 +300,7 @@ void test_math_f64(double x, double y, double z) {
|
||||
// CHECK-CPP-NEXT: [[__X_ADDR_ASCAST_I31:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I29]] to ptr
|
||||
// CHECK-CPP-NEXT: store float [[TMP7]], ptr [[__X_ADDR_ASCAST_I31]], align 4
|
||||
// CHECK-CPP-NEXT: [[TMP8:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I31]], align 4
|
||||
// CHECK-CPP-NEXT: [[TMP9:%.*]] = call float @llvm.fabs.f32(float [[TMP8]])
|
||||
// CHECK-CPP-NEXT: [[TMP9:%.*]] = call noundef float @llvm.fabs.f32(float [[TMP8]])
|
||||
// CHECK-CPP-NEXT: store float [[TMP9]], ptr [[L3_ASCAST]], align 4
|
||||
// CHECK-CPP-NEXT: [[TMP10:%.*]] = load float, ptr [[X_ADDR_ASCAST]], align 4
|
||||
// CHECK-CPP-NEXT: [[TMP11:%.*]] = load float, ptr [[Y_ADDR_ASCAST]], align 4
|
||||
@ -325,7 +325,7 @@ void test_math_f64(double x, double y, double z) {
|
||||
// CHECK-CPP-NEXT: [[TMP16:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I37]], align 4
|
||||
// CHECK-CPP-NEXT: [[TMP17:%.*]] = load float, ptr [[__Y_ADDR_ASCAST_I38]], align 4
|
||||
// CHECK-CPP-NEXT: [[TMP18:%.*]] = load float, ptr [[__Z_ADDR_ASCAST_I39]], align 4
|
||||
// CHECK-CPP-NEXT: [[TMP19:%.*]] = call float @llvm.fma.f32(float [[TMP16]], float [[TMP17]], float [[TMP18]])
|
||||
// CHECK-CPP-NEXT: [[TMP19:%.*]] = call noundef float @llvm.fma.f32(float [[TMP16]], float [[TMP17]], float [[TMP18]])
|
||||
// CHECK-CPP-NEXT: store float [[TMP19]], ptr [[L4_ASCAST]], align 4
|
||||
// CHECK-CPP-NEXT: ret void
|
||||
//
|
||||
@ -437,21 +437,21 @@ void test_math_f32(float x, float y, float z) {
|
||||
// CHECK-CPP-NEXT: [[__X_ADDR_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I]] to ptr
|
||||
// CHECK-CPP-NEXT: store float [[TMP0]], ptr [[__X_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-CPP-NEXT: [[TMP1:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-CPP-NEXT: [[CALL_I:%.*]] = call float @__ocml_sin_f32(float noundef [[TMP1]]) #[[ATTR3]]
|
||||
// CHECK-CPP-NEXT: [[CALL_I:%.*]] = call noundef float @__ocml_sin_f32(float noundef [[TMP1]]) #[[ATTR3]]
|
||||
// CHECK-CPP-NEXT: store float [[CALL_I]], ptr [[L1_ASCAST]], align 4
|
||||
// CHECK-CPP-NEXT: [[TMP2:%.*]] = load float, ptr [[X_ADDR_ASCAST]], align 4
|
||||
// CHECK-CPP-NEXT: [[RETVAL_ASCAST_I6:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I4]] to ptr
|
||||
// CHECK-CPP-NEXT: [[__X_ADDR_ASCAST_I7:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I5]] to ptr
|
||||
// CHECK-CPP-NEXT: store float [[TMP2]], ptr [[__X_ADDR_ASCAST_I7]], align 4
|
||||
// CHECK-CPP-NEXT: [[TMP3:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I7]], align 4
|
||||
// CHECK-CPP-NEXT: [[CALL_I8:%.*]] = call float @__ocml_cos_f32(float noundef [[TMP3]]) #[[ATTR3]]
|
||||
// CHECK-CPP-NEXT: [[CALL_I8:%.*]] = call noundef float @__ocml_cos_f32(float noundef [[TMP3]]) #[[ATTR3]]
|
||||
// CHECK-CPP-NEXT: store float [[CALL_I8]], ptr [[L2_ASCAST]], align 4
|
||||
// CHECK-CPP-NEXT: [[TMP4:%.*]] = load float, ptr [[X_ADDR_ASCAST]], align 4
|
||||
// CHECK-CPP-NEXT: [[RETVAL_ASCAST_I11:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I9]] to ptr
|
||||
// CHECK-CPP-NEXT: [[__X_ADDR_ASCAST_I12:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I10]] to ptr
|
||||
// CHECK-CPP-NEXT: store float [[TMP4]], ptr [[__X_ADDR_ASCAST_I12]], align 4
|
||||
// CHECK-CPP-NEXT: [[TMP5:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I12]], align 4
|
||||
// CHECK-CPP-NEXT: [[TMP6:%.*]] = call float @llvm.fabs.f32(float [[TMP5]])
|
||||
// CHECK-CPP-NEXT: [[TMP6:%.*]] = call noundef float @llvm.fabs.f32(float [[TMP5]])
|
||||
// CHECK-CPP-NEXT: store float [[TMP6]], ptr [[L3_ASCAST]], align 4
|
||||
// CHECK-CPP-NEXT: [[TMP7:%.*]] = load float, ptr [[X_ADDR_ASCAST]], align 4
|
||||
// CHECK-CPP-NEXT: [[TMP8:%.*]] = load float, ptr [[Y_ADDR_ASCAST]], align 4
|
||||
@ -466,7 +466,7 @@ void test_math_f32(float x, float y, float z) {
|
||||
// CHECK-CPP-NEXT: [[TMP10:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I16]], align 4
|
||||
// CHECK-CPP-NEXT: [[TMP11:%.*]] = load float, ptr [[__Y_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-CPP-NEXT: [[TMP12:%.*]] = load float, ptr [[__Z_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-CPP-NEXT: [[TMP13:%.*]] = call float @llvm.fma.f32(float [[TMP10]], float [[TMP11]], float [[TMP12]])
|
||||
// CHECK-CPP-NEXT: [[TMP13:%.*]] = call noundef float @llvm.fma.f32(float [[TMP10]], float [[TMP11]], float [[TMP12]])
|
||||
// CHECK-CPP-NEXT: store float [[TMP13]], ptr [[L4_ASCAST]], align 4
|
||||
// CHECK-CPP-NEXT: ret void
|
||||
//
|
||||
|
||||
@ -44,7 +44,7 @@ const float constexpr_fmaxf_f32 = fmaxf(2.0f, -4.0f);
|
||||
// CHECK-NEXT: [[__X_ADDR_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I]] to ptr
|
||||
// CHECK-NEXT: store float -2.000000e+00, ptr [[__X_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-NEXT: [[TMP1:%.*]] = call float @llvm.fabs.f32(float [[TMP0]])
|
||||
// CHECK-NEXT: [[TMP1:%.*]] = call noundef float @llvm.fabs.f32(float [[TMP0]])
|
||||
// CHECK-NEXT: store float [[TMP1]], ptr addrspacecast (ptr addrspace(1) @_ZL19constexpr_fabsf_f32 to ptr), align 4
|
||||
// CHECK-NEXT: ret void
|
||||
//
|
||||
@ -64,7 +64,7 @@ const float constexpr_fmaxf_f32 = fmaxf(2.0f, -4.0f);
|
||||
// CHECK-NEXT: [[__X_ADDR_ASCAST_I_I:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I_I]] to ptr
|
||||
// CHECK-NEXT: store float [[TMP0]], ptr [[__X_ADDR_ASCAST_I_I]], align 4
|
||||
// CHECK-NEXT: [[TMP1:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I_I]], align 4
|
||||
// CHECK-NEXT: [[TMP2:%.*]] = call float @llvm.fabs.f32(float [[TMP1]])
|
||||
// CHECK-NEXT: [[TMP2:%.*]] = call noundef float @llvm.fabs.f32(float [[TMP1]])
|
||||
// CHECK-NEXT: store float [[TMP2]], ptr addrspacecast (ptr addrspace(1) @_ZL18constexpr_fabs_f32 to ptr), align 4
|
||||
// CHECK-NEXT: ret void
|
||||
//
|
||||
@ -78,7 +78,7 @@ const float constexpr_fmaxf_f32 = fmaxf(2.0f, -4.0f);
|
||||
// CHECK-NEXT: [[__X_ADDR_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I]] to ptr
|
||||
// CHECK-NEXT: store float -2.000000e+00, ptr [[__X_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = call float @__ocml_sin_f32(float noundef [[TMP0]]) #[[ATTR3:[0-9]+]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = call noundef float @__ocml_sin_f32(float noundef [[TMP0]]) #[[ATTR3:[0-9]+]]
|
||||
// CHECK-NEXT: store float [[CALL_I]], ptr addrspacecast (ptr addrspace(1) @_ZL18constexpr_sinf_f32 to ptr), align 4
|
||||
// CHECK-NEXT: ret void
|
||||
//
|
||||
@ -98,7 +98,7 @@ const float constexpr_fmaxf_f32 = fmaxf(2.0f, -4.0f);
|
||||
// CHECK-NEXT: [[__X_ADDR_ASCAST_I_I:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I_I]] to ptr
|
||||
// CHECK-NEXT: store float [[TMP0]], ptr [[__X_ADDR_ASCAST_I_I]], align 4
|
||||
// CHECK-NEXT: [[TMP1:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I_I]], align 4
|
||||
// CHECK-NEXT: [[CALL_I_I:%.*]] = call float @__ocml_sin_f32(float noundef [[TMP1]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I_I:%.*]] = call noundef float @__ocml_sin_f32(float noundef [[TMP1]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: store float [[CALL_I_I]], ptr addrspacecast (ptr addrspace(1) @_ZL17constexpr_sin_f32 to ptr), align 4
|
||||
// CHECK-NEXT: ret void
|
||||
//
|
||||
@ -112,7 +112,7 @@ const float constexpr_fmaxf_f32 = fmaxf(2.0f, -4.0f);
|
||||
// CHECK-NEXT: [[__X_ADDR_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I]] to ptr
|
||||
// CHECK-NEXT: store float -2.000000e+00, ptr [[__X_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = call float @__ocml_cos_f32(float noundef [[TMP0]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I:%.*]] = call noundef float @__ocml_cos_f32(float noundef [[TMP0]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: store float [[CALL_I]], ptr addrspacecast (ptr addrspace(1) @_ZL18constexpr_cosf_f32 to ptr), align 4
|
||||
// CHECK-NEXT: ret void
|
||||
//
|
||||
@ -132,7 +132,7 @@ const float constexpr_fmaxf_f32 = fmaxf(2.0f, -4.0f);
|
||||
// CHECK-NEXT: [[__X_ADDR_ASCAST_I_I:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I_I]] to ptr
|
||||
// CHECK-NEXT: store float [[TMP0]], ptr [[__X_ADDR_ASCAST_I_I]], align 4
|
||||
// CHECK-NEXT: [[TMP1:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I_I]], align 4
|
||||
// CHECK-NEXT: [[CALL_I_I:%.*]] = call float @__ocml_cos_f32(float noundef [[TMP1]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: [[CALL_I_I:%.*]] = call noundef float @__ocml_cos_f32(float noundef [[TMP1]]) #[[ATTR3]]
|
||||
// CHECK-NEXT: store float [[CALL_I_I]], ptr addrspacecast (ptr addrspace(1) @_ZL17constexpr_cos_f32 to ptr), align 4
|
||||
// CHECK-NEXT: ret void
|
||||
//
|
||||
@ -154,7 +154,7 @@ const float constexpr_fmaxf_f32 = fmaxf(2.0f, -4.0f);
|
||||
// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-NEXT: [[TMP1:%.*]] = load float, ptr [[__Y_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-NEXT: [[TMP2:%.*]] = load float, ptr [[__Z_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-NEXT: [[TMP3:%.*]] = call float @llvm.fma.f32(float [[TMP0]], float [[TMP1]], float [[TMP2]])
|
||||
// CHECK-NEXT: [[TMP3:%.*]] = call noundef float @llvm.fma.f32(float [[TMP0]], float [[TMP1]], float [[TMP2]])
|
||||
// CHECK-NEXT: store float [[TMP3]], ptr addrspacecast (ptr addrspace(1) @_ZL18constexpr_fmaf_f32 to ptr), align 4
|
||||
// CHECK-NEXT: ret void
|
||||
//
|
||||
@ -190,7 +190,7 @@ const float constexpr_fmaxf_f32 = fmaxf(2.0f, -4.0f);
|
||||
// CHECK-NEXT: [[TMP3:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I_I]], align 4
|
||||
// CHECK-NEXT: [[TMP4:%.*]] = load float, ptr [[__Y_ADDR_ASCAST_I_I]], align 4
|
||||
// CHECK-NEXT: [[TMP5:%.*]] = load float, ptr [[__Z_ADDR_ASCAST_I_I]], align 4
|
||||
// CHECK-NEXT: [[TMP6:%.*]] = call float @llvm.fma.f32(float [[TMP3]], float [[TMP4]], float [[TMP5]])
|
||||
// CHECK-NEXT: [[TMP6:%.*]] = call noundef float @llvm.fma.f32(float [[TMP3]], float [[TMP4]], float [[TMP5]])
|
||||
// CHECK-NEXT: store float [[TMP6]], ptr addrspacecast (ptr addrspace(1) @_ZL17constexpr_fma_f32 to ptr), align 4
|
||||
// CHECK-NEXT: ret void
|
||||
//
|
||||
@ -208,7 +208,7 @@ const float constexpr_fmaxf_f32 = fmaxf(2.0f, -4.0f);
|
||||
// CHECK-NEXT: store float -4.000000e+00, ptr [[__Y_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-NEXT: [[TMP1:%.*]] = load float, ptr [[__Y_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-NEXT: [[TMP2:%.*]] = call float @llvm.minnum.f32(float [[TMP0]], float [[TMP1]])
|
||||
// CHECK-NEXT: [[TMP2:%.*]] = call noundef float @llvm.minnum.f32(float [[TMP0]], float [[TMP1]])
|
||||
// CHECK-NEXT: store float [[TMP2]], ptr addrspacecast (ptr addrspace(1) @_ZL17constexpr_min_f32 to ptr), align 4
|
||||
// CHECK-NEXT: ret void
|
||||
//
|
||||
@ -226,7 +226,7 @@ const float constexpr_fmaxf_f32 = fmaxf(2.0f, -4.0f);
|
||||
// CHECK-NEXT: store float -4.000000e+00, ptr [[__Y_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-NEXT: [[TMP1:%.*]] = load float, ptr [[__Y_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-NEXT: [[TMP2:%.*]] = call float @llvm.maxnum.f32(float [[TMP0]], float [[TMP1]])
|
||||
// CHECK-NEXT: [[TMP2:%.*]] = call noundef float @llvm.maxnum.f32(float [[TMP0]], float [[TMP1]])
|
||||
// CHECK-NEXT: store float [[TMP2]], ptr addrspacecast (ptr addrspace(1) @_ZL17constexpr_max_f32 to ptr), align 4
|
||||
// CHECK-NEXT: ret void
|
||||
//
|
||||
@ -260,7 +260,7 @@ const float constexpr_fmaxf_f32 = fmaxf(2.0f, -4.0f);
|
||||
// CHECK-NEXT: store float -4.000000e+00, ptr [[__Y_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-NEXT: [[TMP1:%.*]] = load float, ptr [[__Y_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-NEXT: [[TMP2:%.*]] = call float @llvm.minnum.f32(float [[TMP0]], float [[TMP1]])
|
||||
// CHECK-NEXT: [[TMP2:%.*]] = call noundef float @llvm.minnum.f32(float [[TMP0]], float [[TMP1]])
|
||||
// CHECK-NEXT: store float [[TMP2]], ptr addrspacecast (ptr addrspace(1) @_ZL19constexpr_fminf_f32 to ptr), align 4
|
||||
// CHECK-NEXT: ret void
|
||||
//
|
||||
@ -278,7 +278,7 @@ const float constexpr_fmaxf_f32 = fmaxf(2.0f, -4.0f);
|
||||
// CHECK-NEXT: store float -4.000000e+00, ptr [[__Y_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-NEXT: [[TMP1:%.*]] = load float, ptr [[__Y_ADDR_ASCAST_I]], align 4
|
||||
// CHECK-NEXT: [[TMP2:%.*]] = call float @llvm.maxnum.f32(float [[TMP0]], float [[TMP1]])
|
||||
// CHECK-NEXT: [[TMP2:%.*]] = call noundef float @llvm.maxnum.f32(float [[TMP0]], float [[TMP1]])
|
||||
// CHECK-NEXT: store float [[TMP2]], ptr addrspacecast (ptr addrspace(1) @_ZL19constexpr_fmaxf_f32 to ptr), align 4
|
||||
// CHECK-NEXT: ret void
|
||||
//
|
||||
|
||||
@ -111,7 +111,7 @@ __device__ void test_numeric_type() {
|
||||
// instead of fma(_Float16, _Float16, _Float16).
|
||||
|
||||
// CXX14-LABEL: define{{.*}}@_Z8test_fma
|
||||
// CXX14: call contract half @llvm.fma.f16
|
||||
// CXX14: call contract noundef half @llvm.fma.f16
|
||||
__device__ double test_fma(_Float16 h, int i) {
|
||||
return fma(h, h, i);
|
||||
}
|
||||
@ -149,12 +149,12 @@ __device__ double test_isnan() {
|
||||
double d = 5.0;
|
||||
float f = 5.0;
|
||||
|
||||
// AMD_INT_RETURN: call i1 @llvm.is.fpclass.f32(float {{.*}}, i32 3)
|
||||
// AMD_BOOL_RETURN: call i1 @llvm.is.fpclass.f32(float {{.*}}, i32 3)
|
||||
// AMD_INT_RETURN: call noundef i1 @llvm.is.fpclass.f32(float {{.*}}, i32 3)
|
||||
// AMD_BOOL_RETURN: call noundef i1 @llvm.is.fpclass.f32(float {{.*}}, i32 3)
|
||||
r += isnan(f);
|
||||
|
||||
// AMD_INT_RETURN: call i1 @llvm.is.fpclass.f64(double {{.*}}, i32 3)
|
||||
// AMD_BOOL_RETURN: call i1 @llvm.is.fpclass.f64(double {{.*}}, i32 3)
|
||||
// AMD_INT_RETURN: call noundef i1 @llvm.is.fpclass.f64(double {{.*}}, i32 3)
|
||||
// AMD_BOOL_RETURN: call noundef i1 @llvm.is.fpclass.f64(double {{.*}}, i32 3)
|
||||
r += isnan(d);
|
||||
|
||||
return r ;
|
||||
|
||||
@ -12,15 +12,15 @@
|
||||
void test_sqrt(double a1) {
|
||||
#pragma omp target
|
||||
{
|
||||
// CHECK-YES: call double @__nv_sqrt(double
|
||||
// CHECK-YES: call noundef double @__nv_sqrt(double
|
||||
double l1 = sqrt(a1);
|
||||
// CHECK-YES: call double @__nv_pow(double
|
||||
// CHECK-YES: call noundef double @__nv_pow(double
|
||||
double l2 = pow(a1, a1);
|
||||
// CHECK-YES: call double @__nv_modf(double
|
||||
// CHECK-YES: call noundef double @__nv_modf(double
|
||||
double l3 = modf(a1 + 3.5, &a1);
|
||||
// CHECK-YES: call double @__nv_fabs(double
|
||||
// CHECK-YES: call noundef double @__nv_fabs(double
|
||||
double l4 = fabs(a1);
|
||||
// CHECK-YES: call i32 @__nv_abs(i32
|
||||
// CHECK-YES: call noundef i32 @__nv_abs(i32
|
||||
double l5 = abs((int)a1);
|
||||
}
|
||||
}
|
||||
|
||||
@ -12,15 +12,15 @@
|
||||
void test_sqrt(double a1) {
|
||||
#pragma omp target
|
||||
{
|
||||
// CHECK-YES: call double @__nv_sqrt(double
|
||||
// CHECK-YES: call noundef double @__nv_sqrt(double
|
||||
double l1 = sqrt(a1);
|
||||
// CHECK-YES: call double @__nv_pow(double
|
||||
// CHECK-YES: call noundef double @__nv_pow(double
|
||||
double l2 = pow(a1, a1);
|
||||
// CHECK-YES: call double @__nv_modf(double
|
||||
// CHECK-YES: call noundef double @__nv_modf(double
|
||||
double l3 = modf(a1 + 3.5, &a1);
|
||||
// CHECK-YES: call double @__nv_fabs(double
|
||||
// CHECK-YES: call noundef double @__nv_fabs(double
|
||||
double l4 = fabs(a1);
|
||||
// CHECK-YES: call i32 @__nv_abs(i32
|
||||
// CHECK-YES: call noundef i32 @__nv_abs(i32
|
||||
double l5 = abs((int)a1);
|
||||
}
|
||||
}
|
||||
|
||||
@ -4,9 +4,9 @@
|
||||
// REQUIRES: nvptx-registered-target
|
||||
|
||||
// RUN: %clang_cc1 -x c -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
|
||||
// RUN: %clang_cc1 -x c -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -internal-isystem %S/Inputs/include -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -x c -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -internal-isystem %S/Inputs/include -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK-C
|
||||
// RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
|
||||
// RUN: %clang_cc1 -x c++ -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -internal-isystem %S/Inputs/include -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -x c++ -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -internal-isystem %S/Inputs/include -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK-CXX
|
||||
|
||||
#ifdef __cplusplus
|
||||
#include <cstdlib>
|
||||
@ -19,15 +19,20 @@
|
||||
void test_sqrt(double a1) {
|
||||
#pragma omp target
|
||||
{
|
||||
// CHECK: call double @__nv_sqrt(double
|
||||
// CHECK-C: call double @__nv_sqrt(double
|
||||
// CHECK-CXX: call noundef double @__nv_sqrt(double
|
||||
double l1 = sqrt(a1);
|
||||
// CHECK: call double @__nv_pow(double
|
||||
// CHECK-C: call double @__nv_pow(double
|
||||
// CHECK-CXX: call noundef double @__nv_pow(double
|
||||
double l2 = pow(a1, a1);
|
||||
// CHECK: call double @__nv_modf(double
|
||||
// CHECK-C: call double @__nv_modf(double
|
||||
// CHECK-CXX: call noundef double @__nv_modf(double
|
||||
double l3 = modf(a1 + 3.5, &a1);
|
||||
// CHECK: call double @__nv_fabs(double
|
||||
// CHECK-C: call double @__nv_fabs(double
|
||||
// CHECK-CXX: call noundef double @__nv_fabs(double
|
||||
double l4 = fabs(a1);
|
||||
// CHECK: call i32 @__nv_abs(i32
|
||||
// CHECK-C: call i32 @__nv_abs(i32
|
||||
// CHECK-CXX: call noundef i32 @__nv_abs(i32
|
||||
double l5 = abs((int)a1);
|
||||
}
|
||||
}
|
||||
|
||||
@ -12,15 +12,15 @@
|
||||
void test_sqrt(double a1) {
|
||||
#pragma omp target
|
||||
{
|
||||
// CHECK-YES: call double @__nv_sqrt(double
|
||||
// CHECK-YES: call noundef double @__nv_sqrt(double
|
||||
double l1 = sqrt(a1);
|
||||
// CHECK-YES: call double @__nv_pow(double
|
||||
// CHECK-YES: call noundef double @__nv_pow(double
|
||||
double l2 = pow(a1, a1);
|
||||
// CHECK-YES: call double @__nv_modf(double
|
||||
// CHECK-YES: call noundef double @__nv_modf(double
|
||||
double l3 = modf(a1 + 3.5, &a1);
|
||||
// CHECK-YES: call double @__nv_fabs(double
|
||||
// CHECK-YES: call noundef double @__nv_fabs(double
|
||||
double l4 = fabs(a1);
|
||||
// CHECK-YES: call i32 @__nv_abs(i32
|
||||
// CHECK-YES: call noundef i32 @__nv_abs(i32
|
||||
double l5 = abs((int)a1);
|
||||
}
|
||||
}
|
||||
|
||||
@ -12,15 +12,15 @@
|
||||
void test_sqrt(double a1) {
|
||||
#pragma omp target
|
||||
{
|
||||
// CHECK-YES: call double @__nv_sqrt(double
|
||||
// CHECK-YES: call noundef double @__nv_sqrt(double
|
||||
double l1 = sqrt(a1);
|
||||
// CHECK-YES: call double @__nv_pow(double
|
||||
// CHECK-YES: call noundef double @__nv_pow(double
|
||||
double l2 = pow(a1, a1);
|
||||
// CHECK-YES: call double @__nv_modf(double
|
||||
// CHECK-YES: call noundef double @__nv_modf(double
|
||||
double l3 = modf(a1 + 3.5, &a1);
|
||||
// CHECK-YES: call double @__nv_fabs(double
|
||||
// CHECK-YES: call noundef double @__nv_fabs(double
|
||||
double l4 = fabs(a1);
|
||||
// CHECK-YES: call i32 @__nv_abs(i32
|
||||
// CHECK-YES: call noundef i32 @__nv_abs(i32
|
||||
double l5 = abs((int)a1);
|
||||
}
|
||||
}
|
||||
|
||||
@ -7,13 +7,13 @@
|
||||
// 4 calls to modf(f), all translated to __nv_modf calls:
|
||||
|
||||
// CHECK-NOT: _Z.modf
|
||||
// CHECK: call double @__nv_modf(double
|
||||
// CHECK: call noundef double @__nv_modf(double
|
||||
// CHECK-NOT: _Z.modf
|
||||
// CHECK: call float @__nv_modff(float
|
||||
// CHECK: call noundef float @__nv_modff(float
|
||||
// CHECK-NOT: _Z.modf
|
||||
// CHECK: call double @__nv_modf(double
|
||||
// CHECK: call noundef double @__nv_modf(double
|
||||
// CHECK-NOT: _Z.modf
|
||||
// CHECK: call float @__nv_modff(float
|
||||
// CHECK: call noundef float @__nv_modff(float
|
||||
// CHECK-NOT: _Z.modf
|
||||
|
||||
template<typename T>
|
||||
|
||||
@ -9,11 +9,11 @@
|
||||
|
||||
double math(float f, double d) {
|
||||
double r = 0;
|
||||
// SLOW: call float @__nv_sinf(float
|
||||
// FAST: call fast nofpclass(nan inf) float @__nv_fast_sinf(float
|
||||
// SLOW: call noundef float @__nv_sinf(float
|
||||
// FAST: call fast noundef nofpclass(nan inf) float @__nv_fast_sinf(float
|
||||
r += sin(f);
|
||||
// SLOW: call double @__nv_sin(double
|
||||
// FAST: call fast nofpclass(nan inf) double @__nv_sin(double
|
||||
// SLOW: call noundef double @__nv_sin(double
|
||||
// FAST: call fast noundef nofpclass(nan inf) double @__nv_sin(double
|
||||
r += sin(d);
|
||||
return r;
|
||||
}
|
||||
|
||||
@ -8,22 +8,22 @@
|
||||
|
||||
// CHECK-NOT: _Z.sin
|
||||
// CHECK-NOT: _Z.cos
|
||||
// CHECK: call double @__nv_sin(double
|
||||
// CHECK: call noundef double @__nv_sin(double
|
||||
// CHECK-NOT: _Z.sin
|
||||
// CHECK-NOT: _Z.cos
|
||||
// CHECK: call float @__nv_sinf(float
|
||||
// CHECK: call noundef float @__nv_sinf(float
|
||||
// CHECK-NOT: _Z.sin
|
||||
// CHECK-NOT: _Z.cos
|
||||
// CHECK: call double @__nv_sin(double
|
||||
// CHECK: call noundef double @__nv_sin(double
|
||||
// CHECK-NOT: _Z.sin
|
||||
// CHECK-NOT: _Z.cos
|
||||
// CHECK: call double @__nv_cos(double
|
||||
// CHECK: call noundef double @__nv_cos(double
|
||||
// CHECK-NOT: _Z.sin
|
||||
// CHECK-NOT: _Z.cos
|
||||
// CHECK: call float @__nv_sinf(float
|
||||
// CHECK: call noundef float @__nv_sinf(float
|
||||
// CHECK-NOT: _Z.sin
|
||||
// CHECK-NOT: _Z.cos
|
||||
// CHECK: call float @__nv_cosf(float
|
||||
// CHECK: call noundef float @__nv_cosf(float
|
||||
// CHECK-NOT: _Z.sin
|
||||
// CHECK-NOT: _Z.cos
|
||||
|
||||
|
||||
@ -20,17 +20,17 @@
|
||||
|
||||
double math(float f, double d) {
|
||||
double r = 0;
|
||||
// INT_RETURN: call i32 @__nv_isnanf(float
|
||||
// INT_RETURN: call noundef i32 @__nv_isnanf(float
|
||||
// AMD_INT_RETURN_SAFE: call i1 @llvm.is.fpclass.f32(float{{.*}}, i32 3)
|
||||
// AMD_INT_RETURN_FAST: sitofp i32 {{.*}} to double
|
||||
// BOOL_RETURN: call i32 @__nv_isnanf(float
|
||||
// BOOL_RETURN: call noundef i32 @__nv_isnanf(float
|
||||
// AMD_BOOL_RETURN_SAFE: call i1 @llvm.is.fpclass.f32(float{{.*}}, i32 3)
|
||||
// AMD_BOOL_RETURN_FAST: icmp ne i32 {{.*}}, 0
|
||||
r += std::isnan(f);
|
||||
// INT_RETURN: call i32 @__nv_isnand(double
|
||||
// INT_RETURN: call noundef i32 @__nv_isnand(double
|
||||
// AMD_INT_RETURN_SAFE: call i1 @llvm.is.fpclass.f64(double{{.*}}, i32 3)
|
||||
// AMD_INT_RETURN_FAST: sitofp i32 {{.*}} to double
|
||||
// BOOL_RETURN: call i32 @__nv_isnand(double
|
||||
// BOOL_RETURN: call noundef i32 @__nv_isnand(double
|
||||
// AMD_BOOL_RETURN_SAFE: call i1 @llvm.is.fpclass.f64(double{{.*}}, i32 3)
|
||||
// AMD_BOOL_RETURN_FAST: icmp ne i32 {{.*}}, 0
|
||||
r += std::isnan(d);
|
||||
|
||||
@ -1358,6 +1358,8 @@ static AttrBuilder IdentifyValidUBGeneratingAttributes(CallBase &CB) {
|
||||
Valid.addDereferenceableOrNullAttr(DerefOrNullBytes);
|
||||
if (CB.hasRetAttr(Attribute::NoAlias))
|
||||
Valid.addAttribute(Attribute::NoAlias);
|
||||
if (CB.hasRetAttr(Attribute::NoUndef))
|
||||
Valid.addAttribute(Attribute::NoUndef);
|
||||
return Valid;
|
||||
}
|
||||
|
||||
@ -1367,6 +1369,8 @@ static AttrBuilder IdentifyValidPoisonGeneratingAttributes(CallBase &CB) {
|
||||
AttrBuilder Valid(CB.getContext());
|
||||
if (CB.hasRetAttr(Attribute::NonNull))
|
||||
Valid.addAttribute(Attribute::NonNull);
|
||||
if (CB.hasRetAttr(Attribute::Alignment))
|
||||
Valid.addAlignmentAttr(CB.getRetAlign());
|
||||
return Valid;
|
||||
}
|
||||
|
||||
@ -1455,6 +1459,8 @@ static void AddReturnAttributes(CallBase &CB, ValueToValueMapTy &VMap) {
|
||||
// any new poison at @use will trigger UB anyways.
|
||||
// In case 3, we can never propagate nonnull because it may create UB due to
|
||||
// the noundef on @bar.
|
||||
if (ValidPG.getAlignment().valueOrOne() < AL.getRetAlignment().valueOrOne())
|
||||
ValidPG.removeAttribute(Attribute::Alignment);
|
||||
if (ValidPG.hasAttributes()) {
|
||||
// Three checks.
|
||||
// If the callsite has `noundef`, then a poison due to violating the
|
||||
|
||||
@ -30,7 +30,7 @@ define ptr @caller0() {
|
||||
|
||||
define ptr @caller1() {
|
||||
; CHECK-LABEL: define ptr @caller1() {
|
||||
; CHECK-NEXT: [[R_I:%.*]] = call ptr @foo()
|
||||
; CHECK-NEXT: [[R_I:%.*]] = call align 16 ptr @foo()
|
||||
; CHECK-NEXT: ret ptr [[R_I]]
|
||||
;
|
||||
%r = call align(16) ptr @callee0123()
|
||||
@ -39,7 +39,7 @@ define ptr @caller1() {
|
||||
|
||||
define ptr @caller2() {
|
||||
; CHECK-LABEL: define ptr @caller2() {
|
||||
; CHECK-NEXT: [[R_I:%.*]] = call ptr @foo()
|
||||
; CHECK-NEXT: [[R_I:%.*]] = call noundef ptr @foo()
|
||||
; CHECK-NEXT: ret ptr [[R_I]]
|
||||
;
|
||||
%r = call noundef ptr @callee0123()
|
||||
@ -57,7 +57,7 @@ define ptr @caller3() {
|
||||
|
||||
define ptr @caller_0123_dornull() {
|
||||
; CHECK-LABEL: define ptr @caller_0123_dornull() {
|
||||
; CHECK-NEXT: [[R_I:%.*]] = call dereferenceable_or_null(16) ptr @foo()
|
||||
; CHECK-NEXT: [[R_I:%.*]] = call noundef align 32 dereferenceable_or_null(16) ptr @foo()
|
||||
; CHECK-NEXT: ret ptr [[R_I]]
|
||||
;
|
||||
%r = call noundef align(32) dereferenceable_or_null(16) ptr @callee0123()
|
||||
@ -66,7 +66,7 @@ define ptr @caller_0123_dornull() {
|
||||
|
||||
define ptr @caller_0123_d() {
|
||||
; CHECK-LABEL: define ptr @caller_0123_d() {
|
||||
; CHECK-NEXT: [[R_I:%.*]] = call dereferenceable(16) ptr @foo()
|
||||
; CHECK-NEXT: [[R_I:%.*]] = call noundef align 32 dereferenceable(16) ptr @foo()
|
||||
; CHECK-NEXT: ret ptr [[R_I]]
|
||||
;
|
||||
%r = call noundef align(32) dereferenceable(16) ptr @callee0123()
|
||||
@ -105,7 +105,7 @@ define ptr @callee5() {
|
||||
|
||||
define ptr @caller5_fail() {
|
||||
; CHECK-LABEL: define ptr @caller5_fail() {
|
||||
; CHECK-NEXT: [[R_I:%.*]] = call align 64 ptr @foo()
|
||||
; CHECK-NEXT: [[R_I:%.*]] = call noundef align 64 ptr @foo()
|
||||
; CHECK-NEXT: ret ptr [[R_I]]
|
||||
;
|
||||
%r = call noundef align(32) ptr @callee5()
|
||||
@ -114,7 +114,7 @@ define ptr @caller5_fail() {
|
||||
|
||||
define ptr @caller5_okay() {
|
||||
; CHECK-LABEL: define ptr @caller5_okay() {
|
||||
; CHECK-NEXT: [[R_I:%.*]] = call align 64 ptr @foo()
|
||||
; CHECK-NEXT: [[R_I:%.*]] = call noundef align 128 ptr @foo()
|
||||
; CHECK-NEXT: ret ptr [[R_I]]
|
||||
;
|
||||
%r = call noundef align(128) ptr @callee5()
|
||||
@ -249,7 +249,7 @@ define ptr @caller10_fail_maybe_poison() {
|
||||
|
||||
define ptr @caller10_okay_will_be_ub() {
|
||||
; CHECK-LABEL: define ptr @caller10_okay_will_be_ub() {
|
||||
; CHECK-NEXT: [[R_I:%.*]] = call nonnull ptr @foo()
|
||||
; CHECK-NEXT: [[R_I:%.*]] = call noundef nonnull ptr @foo()
|
||||
; CHECK-NEXT: call void @use.ptr(ptr [[R_I]])
|
||||
; CHECK-NEXT: ret ptr [[R_I]]
|
||||
;
|
||||
|
||||
Loading…
x
Reference in New Issue
Block a user