mirror of
https://github.com/capstone-engine/llvm-capstone.git
synced 2025-02-19 17:31:55 +00:00
[Clang] Convert some tests to opaque pointers (NFC)
This commit is contained in:
Nikita Popov
parent
f098fb69f1
commit
06621ecdaf
@ -1,7 +1,7 @@
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-unknown -w -S -o - -emit-llvm %s | FileCheck %s -check-prefix=NO__ERRNO
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-unknown -w -S -o - -emit-llvm -fmath-errno %s | FileCheck %s -check-prefix=HAS_ERRNO
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-unknown-gnu -w -S -o - -emit-llvm -fmath-errno %s | FileCheck %s --check-prefix=HAS_ERRNO_GNU
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-windows-msvc -w -S -o - -emit-llvm -fmath-errno %s | FileCheck %s --check-prefix=HAS_ERRNO_WIN
|
||||
// RUN: %clang_cc1 -triple x86_64-unknown-unknown -w -S -o - -emit-llvm %s | FileCheck %s -check-prefix=NO__ERRNO
|
||||
// RUN: %clang_cc1 -triple x86_64-unknown-unknown -w -S -o - -emit-llvm -fmath-errno %s | FileCheck %s -check-prefix=HAS_ERRNO
|
||||
// RUN: %clang_cc1 -triple x86_64-unknown-unknown-gnu -w -S -o - -emit-llvm -fmath-errno %s | FileCheck %s --check-prefix=HAS_ERRNO_GNU
|
||||
// RUN: %clang_cc1 -triple x86_64-unknown-windows-msvc -w -S -o - -emit-llvm -fmath-errno %s | FileCheck %s --check-prefix=HAS_ERRNO_WIN
|
||||
|
||||
// Test attributes and codegen of math builtins.
|
||||
|
||||
@ -52,14 +52,14 @@ void foo(double *d, float f, float *fp, long double *l, int *i, const char *c) {
|
||||
|
||||
__builtin_frexp(f,i); __builtin_frexpf(f,i); __builtin_frexpl(f,i); __builtin_frexpf128(f,i);
|
||||
|
||||
// NO__ERRNO: declare double @frexp(double noundef, i32* noundef) [[NOT_READNONE:#[0-9]+]]
|
||||
// NO__ERRNO: declare float @frexpf(float noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare x86_fp80 @frexpl(x86_fp80 noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare fp128 @frexpf128(fp128 noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare double @frexp(double noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare float @frexpf(float noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare x86_fp80 @frexpl(x86_fp80 noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare fp128 @frexpf128(fp128 noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare double @frexp(double noundef, ptr noundef) [[NOT_READNONE:#[0-9]+]]
|
||||
// NO__ERRNO: declare float @frexpf(float noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare x86_fp80 @frexpl(x86_fp80 noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare fp128 @frexpf128(fp128 noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare double @frexp(double noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare float @frexpf(float noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare x86_fp80 @frexpl(x86_fp80 noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare fp128 @frexpf128(fp128 noundef, ptr noundef) [[NOT_READNONE]]
|
||||
|
||||
__builtin_huge_val(); __builtin_huge_valf(); __builtin_huge_vall(); __builtin_huge_valf128();
|
||||
|
||||
@ -88,36 +88,36 @@ void foo(double *d, float f, float *fp, long double *l, int *i, const char *c) {
|
||||
|
||||
__builtin_modf(f,d); __builtin_modff(f,fp); __builtin_modfl(f,l); __builtin_modff128(f,l);
|
||||
|
||||
// NO__ERRNO: declare double @modf(double noundef, double* noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare float @modff(float noundef, float* noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare x86_fp80 @modfl(x86_fp80 noundef, x86_fp80* noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare fp128 @modff128(fp128 noundef, fp128* noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare double @modf(double noundef, double* noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare float @modff(float noundef, float* noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare x86_fp80 @modfl(x86_fp80 noundef, x86_fp80* noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare fp128 @modff128(fp128 noundef, fp128* noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare double @modf(double noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare float @modff(float noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare x86_fp80 @modfl(x86_fp80 noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare fp128 @modff128(fp128 noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare double @modf(double noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare float @modff(float noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare x86_fp80 @modfl(x86_fp80 noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare fp128 @modff128(fp128 noundef, ptr noundef) [[NOT_READNONE]]
|
||||
|
||||
__builtin_nan(c); __builtin_nanf(c); __builtin_nanl(c); __builtin_nanf128(c);
|
||||
|
||||
// NO__ERRNO: declare double @nan(i8* noundef) [[PURE:#[0-9]+]]
|
||||
// NO__ERRNO: declare float @nanf(i8* noundef) [[PURE]]
|
||||
// NO__ERRNO: declare x86_fp80 @nanl(i8* noundef) [[PURE]]
|
||||
// NO__ERRNO: declare fp128 @nanf128(i8* noundef) [[PURE]]
|
||||
// HAS_ERRNO: declare double @nan(i8* noundef) [[PURE:#[0-9]+]]
|
||||
// HAS_ERRNO: declare float @nanf(i8* noundef) [[PURE]]
|
||||
// HAS_ERRNO: declare x86_fp80 @nanl(i8* noundef) [[PURE]]
|
||||
// HAS_ERRNO: declare fp128 @nanf128(i8* noundef) [[PURE]]
|
||||
// NO__ERRNO: declare double @nan(ptr noundef) [[PURE:#[0-9]+]]
|
||||
// NO__ERRNO: declare float @nanf(ptr noundef) [[PURE]]
|
||||
// NO__ERRNO: declare x86_fp80 @nanl(ptr noundef) [[PURE]]
|
||||
// NO__ERRNO: declare fp128 @nanf128(ptr noundef) [[PURE]]
|
||||
// HAS_ERRNO: declare double @nan(ptr noundef) [[PURE:#[0-9]+]]
|
||||
// HAS_ERRNO: declare float @nanf(ptr noundef) [[PURE]]
|
||||
// HAS_ERRNO: declare x86_fp80 @nanl(ptr noundef) [[PURE]]
|
||||
// HAS_ERRNO: declare fp128 @nanf128(ptr noundef) [[PURE]]
|
||||
|
||||
__builtin_nans(c); __builtin_nansf(c); __builtin_nansl(c); __builtin_nansf128(c);
|
||||
|
||||
// NO__ERRNO: declare double @nans(i8* noundef) [[PURE]]
|
||||
// NO__ERRNO: declare float @nansf(i8* noundef) [[PURE]]
|
||||
// NO__ERRNO: declare x86_fp80 @nansl(i8* noundef) [[PURE]]
|
||||
// NO__ERRNO: declare fp128 @nansf128(i8* noundef) [[PURE]]
|
||||
// HAS_ERRNO: declare double @nans(i8* noundef) [[PURE]]
|
||||
// HAS_ERRNO: declare float @nansf(i8* noundef) [[PURE]]
|
||||
// HAS_ERRNO: declare x86_fp80 @nansl(i8* noundef) [[PURE]]
|
||||
// HAS_ERRNO: declare fp128 @nansf128(i8* noundef) [[PURE]]
|
||||
// NO__ERRNO: declare double @nans(ptr noundef) [[PURE]]
|
||||
// NO__ERRNO: declare float @nansf(ptr noundef) [[PURE]]
|
||||
// NO__ERRNO: declare x86_fp80 @nansl(ptr noundef) [[PURE]]
|
||||
// NO__ERRNO: declare fp128 @nansf128(ptr noundef) [[PURE]]
|
||||
// HAS_ERRNO: declare double @nans(ptr noundef) [[PURE]]
|
||||
// HAS_ERRNO: declare float @nansf(ptr noundef) [[PURE]]
|
||||
// HAS_ERRNO: declare x86_fp80 @nansl(ptr noundef) [[PURE]]
|
||||
// HAS_ERRNO: declare fp128 @nansf128(ptr noundef) [[PURE]]
|
||||
|
||||
__builtin_pow(f,f); __builtin_powf(f,f); __builtin_powl(f,f); __builtin_powf128(f,f);
|
||||
|
||||
@ -549,14 +549,14 @@ __builtin_remainder(f,f); __builtin_remainderf(f,f); __builtin_remainderl(f,f);
|
||||
|
||||
__builtin_remquo(f,f,i); __builtin_remquof(f,f,i); __builtin_remquol(f,f,i); __builtin_remquof128(f,f,i);
|
||||
|
||||
// NO__ERRNO: declare double @remquo(double noundef, double noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare float @remquof(float noundef, float noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare x86_fp80 @remquol(x86_fp80 noundef, x86_fp80 noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare fp128 @remquof128(fp128 noundef, fp128 noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare double @remquo(double noundef, double noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare float @remquof(float noundef, float noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare x86_fp80 @remquol(x86_fp80 noundef, x86_fp80 noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare fp128 @remquof128(fp128 noundef, fp128 noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare double @remquo(double noundef, double noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare float @remquof(float noundef, float noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare x86_fp80 @remquol(x86_fp80 noundef, x86_fp80 noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare fp128 @remquof128(fp128 noundef, fp128 noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare double @remquo(double noundef, double noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare float @remquof(float noundef, float noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare x86_fp80 @remquol(x86_fp80 noundef, x86_fp80 noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare fp128 @remquof128(fp128 noundef, fp128 noundef, ptr noundef) [[NOT_READNONE]]
|
||||
|
||||
__builtin_rint(f); __builtin_rintf(f); __builtin_rintl(f); __builtin_rintf128(f);
|
||||
|
||||
|
||||
@ -1,8 +1,8 @@
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-unknown -Wno-implicit-function-declaration -w -S -o - -emit-llvm %s | FileCheck %s --check-prefix=NO__ERRNO
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-unknown -Wno-implicit-function-declaration -w -S -o - -emit-llvm -fmath-errno %s | FileCheck %s --check-prefix=HAS_ERRNO
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-unknown -Wno-implicit-function-declaration -w -S -o - -emit-llvm -ffp-exception-behavior=maytrap %s | FileCheck %s --check-prefix=HAS_MAYTRAP
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-unknown-gnu -Wno-implicit-function-declaration -w -S -o - -emit-llvm -fmath-errno %s | FileCheck %s --check-prefix=HAS_ERRNO_GNU
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-windows-msvc -Wno-implicit-function-declaration -w -S -o - -emit-llvm -fmath-errno %s | FileCheck %s --check-prefix=HAS_ERRNO_WIN
|
||||
// RUN: %clang_cc1 -triple x86_64-unknown-unknown -Wno-implicit-function-declaration -w -S -o - -emit-llvm %s | FileCheck %s --check-prefix=NO__ERRNO
|
||||
// RUN: %clang_cc1 -triple x86_64-unknown-unknown -Wno-implicit-function-declaration -w -S -o - -emit-llvm -fmath-errno %s | FileCheck %s --check-prefix=HAS_ERRNO
|
||||
// RUN: %clang_cc1 -triple x86_64-unknown-unknown -Wno-implicit-function-declaration -w -S -o - -emit-llvm -ffp-exception-behavior=maytrap %s | FileCheck %s --check-prefix=HAS_MAYTRAP
|
||||
// RUN: %clang_cc1 -triple x86_64-unknown-unknown-gnu -Wno-implicit-function-declaration -w -S -o - -emit-llvm -fmath-errno %s | FileCheck %s --check-prefix=HAS_ERRNO_GNU
|
||||
// RUN: %clang_cc1 -triple x86_64-unknown-windows-msvc -Wno-implicit-function-declaration -w -S -o - -emit-llvm -fmath-errno %s | FileCheck %s --check-prefix=HAS_ERRNO_WIN
|
||||
|
||||
// Test attributes and builtin codegen of math library calls.
|
||||
|
||||
@ -57,15 +57,15 @@ void foo(double *d, float f, float *fp, long double *l, int *i, const char *c) {
|
||||
|
||||
frexp(f,i); frexpf(f,i); frexpl(f,i);
|
||||
|
||||
// NO__ERRNO: declare double @frexp(double noundef, i32* noundef) [[NOT_READNONE:#[0-9]+]]
|
||||
// NO__ERRNO: declare float @frexpf(float noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare x86_fp80 @frexpl(x86_fp80 noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare double @frexp(double noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare float @frexpf(float noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare x86_fp80 @frexpl(x86_fp80 noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// HAS_MAYTRAP: declare double @frexp(double noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// HAS_MAYTRAP: declare float @frexpf(float noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// HAS_MAYTRAP: declare x86_fp80 @frexpl(x86_fp80 noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare double @frexp(double noundef, ptr noundef) [[NOT_READNONE:#[0-9]+]]
|
||||
// NO__ERRNO: declare float @frexpf(float noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare x86_fp80 @frexpl(x86_fp80 noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare double @frexp(double noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare float @frexpf(float noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare x86_fp80 @frexpl(x86_fp80 noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_MAYTRAP: declare double @frexp(double noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_MAYTRAP: declare float @frexpf(float noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_MAYTRAP: declare x86_fp80 @frexpl(x86_fp80 noundef, ptr noundef) [[NOT_READNONE]]
|
||||
|
||||
ldexp(f,f); ldexpf(f,f); ldexpl(f,f);
|
||||
|
||||
@ -81,27 +81,27 @@ void foo(double *d, float f, float *fp, long double *l, int *i, const char *c) {
|
||||
|
||||
modf(f,d); modff(f,fp); modfl(f,l);
|
||||
|
||||
// NO__ERRNO: declare double @modf(double noundef, double* noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare float @modff(float noundef, float* noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare x86_fp80 @modfl(x86_fp80 noundef, x86_fp80* noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare double @modf(double noundef, double* noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare float @modff(float noundef, float* noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare x86_fp80 @modfl(x86_fp80 noundef, x86_fp80* noundef) [[NOT_READNONE]]
|
||||
// HAS_MAYTRAP: declare double @modf(double noundef, double* noundef) [[NOT_READNONE]]
|
||||
// HAS_MAYTRAP: declare float @modff(float noundef, float* noundef) [[NOT_READNONE]]
|
||||
// HAS_MAYTRAP: declare x86_fp80 @modfl(x86_fp80 noundef, x86_fp80* noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare double @modf(double noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare float @modff(float noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare x86_fp80 @modfl(x86_fp80 noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare double @modf(double noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare float @modff(float noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare x86_fp80 @modfl(x86_fp80 noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_MAYTRAP: declare double @modf(double noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_MAYTRAP: declare float @modff(float noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_MAYTRAP: declare x86_fp80 @modfl(x86_fp80 noundef, ptr noundef) [[NOT_READNONE]]
|
||||
|
||||
nan(c); nanf(c); nanl(c);
|
||||
|
||||
// NO__ERRNO: declare double @nan(i8* noundef) [[READONLY:#[0-9]+]]
|
||||
// NO__ERRNO: declare float @nanf(i8* noundef) [[READONLY]]
|
||||
// NO__ERRNO: declare x86_fp80 @nanl(i8* noundef) [[READONLY]]
|
||||
// HAS_ERRNO: declare double @nan(i8* noundef) [[READONLY:#[0-9]+]]
|
||||
// HAS_ERRNO: declare float @nanf(i8* noundef) [[READONLY]]
|
||||
// HAS_ERRNO: declare x86_fp80 @nanl(i8* noundef) [[READONLY]]
|
||||
// HAS_MAYTRAP: declare double @nan(i8* noundef) [[READONLY:#[0-9]+]]
|
||||
// HAS_MAYTRAP: declare float @nanf(i8* noundef) [[READONLY]]
|
||||
// HAS_MAYTRAP: declare x86_fp80 @nanl(i8* noundef) [[READONLY]]
|
||||
// NO__ERRNO: declare double @nan(ptr noundef) [[READONLY:#[0-9]+]]
|
||||
// NO__ERRNO: declare float @nanf(ptr noundef) [[READONLY]]
|
||||
// NO__ERRNO: declare x86_fp80 @nanl(ptr noundef) [[READONLY]]
|
||||
// HAS_ERRNO: declare double @nan(ptr noundef) [[READONLY:#[0-9]+]]
|
||||
// HAS_ERRNO: declare float @nanf(ptr noundef) [[READONLY]]
|
||||
// HAS_ERRNO: declare x86_fp80 @nanl(ptr noundef) [[READONLY]]
|
||||
// HAS_MAYTRAP: declare double @nan(ptr noundef) [[READONLY:#[0-9]+]]
|
||||
// HAS_MAYTRAP: declare float @nanf(ptr noundef) [[READONLY]]
|
||||
// HAS_MAYTRAP: declare x86_fp80 @nanl(ptr noundef) [[READONLY]]
|
||||
|
||||
pow(f,f); powf(f,f); powl(f,f);
|
||||
|
||||
@ -564,15 +564,15 @@ void foo(double *d, float f, float *fp, long double *l, int *i, const char *c) {
|
||||
|
||||
remquo(f,f,i); remquof(f,f,i); remquol(f,f,i);
|
||||
|
||||
// NO__ERRNO: declare double @remquo(double noundef, double noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare float @remquof(float noundef, float noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare x86_fp80 @remquol(x86_fp80 noundef, x86_fp80 noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare double @remquo(double noundef, double noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare float @remquof(float noundef, float noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare x86_fp80 @remquol(x86_fp80 noundef, x86_fp80 noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// HAS_MAYTRAP: declare double @remquo(double noundef, double noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// HAS_MAYTRAP: declare float @remquof(float noundef, float noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// HAS_MAYTRAP: declare x86_fp80 @remquol(x86_fp80 noundef, x86_fp80 noundef, i32* noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare double @remquo(double noundef, double noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare float @remquof(float noundef, float noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// NO__ERRNO: declare x86_fp80 @remquol(x86_fp80 noundef, x86_fp80 noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare double @remquo(double noundef, double noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare float @remquof(float noundef, float noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_ERRNO: declare x86_fp80 @remquol(x86_fp80 noundef, x86_fp80 noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_MAYTRAP: declare double @remquo(double noundef, double noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_MAYTRAP: declare float @remquof(float noundef, float noundef, ptr noundef) [[NOT_READNONE]]
|
||||
// HAS_MAYTRAP: declare x86_fp80 @remquol(x86_fp80 noundef, x86_fp80 noundef, ptr noundef) [[NOT_READNONE]]
|
||||
|
||||
rint(f); rintf(f); rintl(f);
|
||||
|
||||
|
||||
@ -1,14 +1,14 @@
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple i686 %s -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64 %s -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple arm %s -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple mips %s -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple mipsel %s -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple powerpc %s -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple powerpc64 %s -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple s390x %s -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple sparc %s -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple sparcv9 %s -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple thumb %s -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -triple i686 %s -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -triple x86_64 %s -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -triple arm %s -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -triple mips %s -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -triple mipsel %s -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -triple powerpc %s -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -triple powerpc64 %s -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -triple s390x %s -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -triple sparc %s -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -triple sparcv9 %s -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -triple thumb %s -emit-llvm -o - | FileCheck %s
|
||||
|
||||
int mout0;
|
||||
int min1;
|
||||
@ -17,7 +17,7 @@ int marray[2];
|
||||
// CHECK: @single_m
|
||||
void single_m(void)
|
||||
{
|
||||
// CHECK: call void asm "foo $1,$0", "=*m,*m[[CLOBBERS:[a-zA-Z0-9@%{},~_$ ]*\"]](i32* elementtype(i32) {{[a-zA-Z0-9@%]+}}, i32* elementtype(i32) {{[a-zA-Z0-9@%]+}})
|
||||
// CHECK: call void asm "foo $1,$0", "=*m,*m[[CLOBBERS:[a-zA-Z0-9@%{},~_$ ]*\"]](ptr elementtype(i32) {{[a-zA-Z0-9@%]+}}, ptr elementtype(i32) {{[a-zA-Z0-9@%]+}})
|
||||
asm("foo %1,%0" : "=m" (mout0) : "m" (min1));
|
||||
}
|
||||
|
||||
@ -130,7 +130,7 @@ void single_X(void)
|
||||
asm("foo %1,%0" : "=r" (out0) : "X" (min1));
|
||||
// CHECK: call i32 asm "foo $1,$0", "=r,X[[CLOBBERS]](i32 1)
|
||||
asm("foo %1,%0" : "=r" (out0) : "X" (1));
|
||||
// CHECK: call i32 asm "foo $1,$0", "=r,X[[CLOBBERS]](i32* getelementptr inbounds ([2 x i32], [2 x i32]* {{[a-zA-Z0-9@%]+}}, i{{32|64}} 0, i{{32|64}} 0))
|
||||
// CHECK: call i32 asm "foo $1,$0", "=r,X[[CLOBBERS]](ptr {{[a-zA-Z0-9@%]+}})
|
||||
asm("foo %1,%0" : "=r" (out0) : "X" (marray));
|
||||
// CHECK: call i32 asm "foo $1,$0", "=r,X[[CLOBBERS]](double {{[0-9.eE+-]+}})
|
||||
asm("foo %1,%0" : "=r" (out0) : "X" (1.0e+01));
|
||||
@ -143,14 +143,14 @@ void single_p(void)
|
||||
{
|
||||
register int out0 = 0;
|
||||
// Constraint converted differently on different platforms moved to platform-specific.
|
||||
// : call i32 asm "foo $1,$0", "=r,im[[CLOBBERS]](i32* getelementptr inbounds ([2 x i32], [2 x i32]* {{[a-zA-Z0-9@%]+}}, i{{32|64}} 0, i{{32|64}} 0))
|
||||
// : call i32 asm "foo $1,$0", "=r,im[[CLOBBERS]](ptr getelementptr inbounds ([2 x i32], ptr {{[a-zA-Z0-9@%]+}}, i{{32|64}} 0, i{{32|64}} 0))
|
||||
asm("foo %1,%0" : "=r" (out0) : "p" (marray));
|
||||
}
|
||||
|
||||
// CHECK: @multi_m
|
||||
void multi_m(void)
|
||||
{
|
||||
// CHECK: call void asm "foo $1,$0", "=*m|r,m|r[[CLOBBERS]](i32* elementtype(i32) {{[a-zA-Z0-9@%]+}}, i32 {{[a-zA-Z0-9@%]+}})
|
||||
// CHECK: call void asm "foo $1,$0", "=*m|r,m|r[[CLOBBERS]](ptr elementtype(i32) {{[a-zA-Z0-9@%]+}}, i32 {{[a-zA-Z0-9@%]+}})
|
||||
asm("foo %1,%0" : "=m,r" (mout0) : "m,r" (min1));
|
||||
}
|
||||
|
||||
@ -263,7 +263,7 @@ void multi_X(void)
|
||||
asm("foo %1,%0" : "=r,r" (out0) : "r,X" (min1));
|
||||
// CHECK: call i32 asm "foo $1,$0", "=r|r,r|X[[CLOBBERS]](i32 1)
|
||||
asm("foo %1,%0" : "=r,r" (out0) : "r,X" (1));
|
||||
// CHECK: call i32 asm "foo $1,$0", "=r|r,r|X[[CLOBBERS]](i32* getelementptr inbounds ([2 x i32], [2 x i32]* {{[a-zA-Z0-9@%]+}}, i{{32|64}} 0, i{{32|64}} 0))
|
||||
// CHECK: call i32 asm "foo $1,$0", "=r|r,r|X[[CLOBBERS]](ptr {{[a-zA-Z0-9@%]+}})
|
||||
asm("foo %1,%0" : "=r,r" (out0) : "r,X" (marray));
|
||||
// CHECK: call i32 asm "foo $1,$0", "=r|r,r|X[[CLOBBERS]](double {{[0-9.eE+-]+}})
|
||||
asm("foo %1,%0" : "=r,r" (out0) : "r,X" (1.0e+01));
|
||||
@ -276,6 +276,6 @@ void multi_p(void)
|
||||
{
|
||||
register int out0 = 0;
|
||||
// Constraint converted differently on different platforms moved to platform-specific.
|
||||
// : call i32 asm "foo $1,$0", "=r|r,r|im[[CLOBBERS]](i32* getelementptr inbounds ([2 x i32], [2 x i32]* {{[a-zA-Z0-9@%]+}}, {{i[0-9]*}} 0, {{i[0-9]*}} 0))
|
||||
// : call i32 asm "foo $1,$0", "=r|r,r|im[[CLOBBERS]](ptr getelementptr inbounds ([2 x i32], ptr {{[a-zA-Z0-9@%]+}}, {{i[0-9]*}} 0, {{i[0-9]*}} 0))
|
||||
asm("foo %1,%0" : "=r,r" (out0) : "r,p" (marray));
|
||||
}
|
||||
|
||||
@ -1,63 +1,63 @@
|
||||
// RUN: echo -n "GPU binary would be here." > %t
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s \
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
|
||||
// RUN: -target-sdk-version=8.0 -fcuda-include-gpubinary %t -o - \
|
||||
// RUN: | FileCheck -allow-deprecated-dag-overlap %s \
|
||||
// RUN: --check-prefixes=ALL,LNX,NORDC,CUDA,CUDANORDC,CUDA-OLD
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s \
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
|
||||
// RUN: -target-sdk-version=8.0 -fcuda-include-gpubinary %t \
|
||||
// RUN: -o - -DNOGLOBALS \
|
||||
// RUN: | FileCheck -allow-deprecated-dag-overlap %s \
|
||||
// RUN: -check-prefixes=NOGLOBALS,CUDANOGLOBALS
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s \
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
|
||||
// RUN: -target-sdk-version=8.0 -fgpu-rdc -fcuda-include-gpubinary %t \
|
||||
// RUN: -o - \
|
||||
// RUN: | FileCheck -allow-deprecated-dag-overlap %s \
|
||||
// RUN: --check-prefixes=ALL,LNX,RDC,CUDA,CUDARDC,CUDA-OLD
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s \
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
|
||||
// RUN: -target-sdk-version=8.0 -o - \
|
||||
// RUN: | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=NOGPUBIN
|
||||
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s \
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
|
||||
// RUN: -target-sdk-version=9.2 -fcuda-include-gpubinary %t -o - \
|
||||
// RUN: | FileCheck %s -allow-deprecated-dag-overlap \
|
||||
// RUN: --check-prefixes=ALL,LNX,NORDC,CUDA,CUDANORDC,CUDA-NEW
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s \
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
|
||||
// RUN: -target-sdk-version=9.2 -fcuda-include-gpubinary %t -o - -DNOGLOBALS \
|
||||
// RUN: | FileCheck -allow-deprecated-dag-overlap %s \
|
||||
// RUN: --check-prefixes=NOGLOBALS,CUDANOGLOBALS
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s \
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
|
||||
// RUN: -target-sdk-version=9.2 -fgpu-rdc -fcuda-include-gpubinary %t -o - \
|
||||
// RUN: | FileCheck %s -allow-deprecated-dag-overlap \
|
||||
// RUN: --check-prefixes=ALL,LNX,RDC,CUDA,CUDARDC,CUDA-NEW
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s -std=c++17 \
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -std=c++17 \
|
||||
// RUN: -target-sdk-version=9.2 -fcuda-include-gpubinary %t -o - \
|
||||
// RUN: | FileCheck %s -allow-deprecated-dag-overlap \
|
||||
// RUN: --check-prefixes=ALL,LNX,NORDC,CUDA,CUDANORDC,CUDA-NEW,LNX_17,NORDC17
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s -std=c++17 \
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -std=c++17 \
|
||||
// RUN: -target-sdk-version=9.2 -fgpu-rdc -fcuda-include-gpubinary %t -o - \
|
||||
// RUN: | FileCheck %s -allow-deprecated-dag-overlap \
|
||||
// RUN: --check-prefixes=ALL,LNX,RDC,CUDA,CUDARDC,CUDA-NEW,LNX_17,RDC17
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s \
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
|
||||
// RUN: -target-sdk-version=9.2 -o - \
|
||||
// RUN: | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=NOGPUBIN
|
||||
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s \
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
|
||||
// RUN: -fcuda-include-gpubinary %t -o - -x hip\
|
||||
// RUN: | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,LNX,NORDC,HIP,HIPEF
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s \
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
|
||||
// RUN: -fcuda-include-gpubinary %t -o - -DNOGLOBALS -x hip \
|
||||
// RUN: | FileCheck -allow-deprecated-dag-overlap %s -check-prefixes=NOGLOBALS,HIPNOGLOBALS
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s \
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
|
||||
// RUN: -fgpu-rdc -fcuda-include-gpubinary %t -o - -x hip \
|
||||
// RUN: | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,LNX,RDC,HIP,HIPEF
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s -o - -x hip\
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -o - -x hip\
|
||||
// RUN: | FileCheck -allow-deprecated-dag-overlap %s -check-prefixes=ALL,LNX,NORDC,HIP,HIPNEF
|
||||
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-pc-windows-msvc -aux-triple amdgcn -emit-llvm %s \
|
||||
// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -aux-triple amdgcn -emit-llvm %s \
|
||||
// RUN: -fcuda-include-gpubinary %t -o - -x hip\
|
||||
// RUN: | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,WIN
|
||||
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-pc-windows-msvc -aux-triple amdgcn -emit-llvm %s \
|
||||
// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -aux-triple amdgcn -emit-llvm %s \
|
||||
// RUN: -o - -x hip\
|
||||
// RUN: | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,WIN,HIP,HIPNEF
|
||||
|
||||
@ -166,26 +166,26 @@ __device__ void device_use() {
|
||||
// CUDARDC-SAME: section "__nv_relfatbin", align 8
|
||||
// * constant struct that wraps GPU binary
|
||||
// ALL: @__[[PREFIX:cuda|hip]]_fatbin_wrapper = internal constant
|
||||
// LNX-SAME: { i32, i32, i8*, i8* }
|
||||
// LNX-SAME: { i32, i32, ptr, ptr }
|
||||
// CUDA-SAME: { i32 1180844977, i32 1,
|
||||
// HIP-SAME: { i32 1212764230, i32 1,
|
||||
// CUDA-SAME: i8* getelementptr inbounds ({{.*}}@[[FATBIN]], i64 0, i64 0),
|
||||
// HIPEF-SAME: i8* getelementptr inbounds ({{.*}}@[[FATBIN]], i64 0, i64 0),
|
||||
// HIPNEF-SAME: i8* @[[FATBIN]],
|
||||
// LNX-SAME: i8* null }
|
||||
// CUDA-SAME: ptr @[[FATBIN]],
|
||||
// HIPEF-SAME: ptr @[[FATBIN]],
|
||||
// HIPNEF-SAME: ptr @[[FATBIN]],
|
||||
// LNX-SAME: ptr null }
|
||||
// CUDA-SAME: section ".nvFatBinSegment"
|
||||
// HIP-SAME: section ".hipFatBinSegment"
|
||||
// * variable to save GPU binary handle after initialization
|
||||
// CUDANORDC: @__[[PREFIX]]_gpubin_handle = internal global i8** null
|
||||
// HIPNEF: @__[[PREFIX]]_gpubin_handle = linkonce hidden global i8** null
|
||||
// CUDANORDC: @__[[PREFIX]]_gpubin_handle = internal global ptr null
|
||||
// HIPNEF: @__[[PREFIX]]_gpubin_handle = linkonce hidden global ptr null
|
||||
// * constant unnamed string with NVModuleID
|
||||
// CUDARDC: [[MODULE_ID_GLOBAL:@.*]] = private constant
|
||||
// CUDARDC-SAME: c"[[MODULE_ID:.+]]\00", section "__nv_module_id", align 32
|
||||
// * Make sure our constructor was added to global ctor list.
|
||||
// LNX: @llvm.global_ctors = appending global {{.*}}@__[[PREFIX]]_module_ctor
|
||||
// * Alias to global symbol containing the NVModuleID.
|
||||
// CUDARDC: @__fatbinwrap[[MODULE_ID]] ={{.*}} alias { i32, i32, i8*, i8* }
|
||||
// CUDARDC-SAME: { i32, i32, i8*, i8* }* @__[[PREFIX]]_fatbin_wrapper
|
||||
// CUDARDC: @__fatbinwrap[[MODULE_ID]] ={{.*}} alias { i32, i32, ptr, ptr }
|
||||
// CUDARDC-SAME: ptr @__[[PREFIX]]_fatbin_wrapper
|
||||
|
||||
// Test that we build the correct number of calls to cudaSetupArgument followed
|
||||
// by a call to cudaLaunch.
|
||||
@ -225,13 +225,13 @@ void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); }
|
||||
|
||||
// Test that we've built a function to register kernels and global vars.
|
||||
// ALL: define internal void @__[[PREFIX]]_register_globals
|
||||
// ALL: call{{.*}}[[PREFIX]]RegisterFunction(i8** %0, {{.*}}kernelfunc{{[^,]*}}, {{[^@]*}}@0
|
||||
// ALL-DAG: call void {{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}device_var{{[^,]*}}, {{[^@]*}}@1, {{.*}}i32 0, {{i32|i64}} 4, i32 0, i32 0
|
||||
// ALL-DAG: call void {{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}constant_var{{[^,]*}}, {{[^@]*}}@2, {{.*}}i32 0, {{i32|i64}} 4, i32 1, i32 0
|
||||
// ALL-DAG: call void {{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}ext_device_var_def{{[^,]*}}, {{[^@]*}}@3, {{.*}}i32 0, {{i32|i64}} 4, i32 0, i32 0
|
||||
// ALL-DAG: call void {{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}ext_constant_var_def{{[^,]*}}, {{[^@]*}}@4, {{.*}}i32 0, {{i32|i64}} 4, i32 1, i32 0
|
||||
// LNX_17-DAG: [[PREFIX]]RegisterVar(i8** %0, {{.*}}inline_var
|
||||
// LNX_17-NOT: [[PREFIX]]RegisterVar(i8** %0, {{.*}}inline_var2
|
||||
// ALL: call{{.*}}[[PREFIX]]RegisterFunction(ptr %0, {{.*}}kernelfunc{{[^,]*}}, {{[^@]*}}@0
|
||||
// ALL-DAG: call void {{.*}}[[PREFIX]]RegisterVar(ptr %0, {{.*}}device_var{{[^,]*}}, {{[^@]*}}@1, {{.*}}i32 0, {{i32|i64}} 4, i32 0, i32 0
|
||||
// ALL-DAG: call void {{.*}}[[PREFIX]]RegisterVar(ptr %0, {{.*}}constant_var{{[^,]*}}, {{[^@]*}}@2, {{.*}}i32 0, {{i32|i64}} 4, i32 1, i32 0
|
||||
// ALL-DAG: call void {{.*}}[[PREFIX]]RegisterVar(ptr %0, {{.*}}ext_device_var_def{{[^,]*}}, {{[^@]*}}@3, {{.*}}i32 0, {{i32|i64}} 4, i32 0, i32 0
|
||||
// ALL-DAG: call void {{.*}}[[PREFIX]]RegisterVar(ptr %0, {{.*}}ext_constant_var_def{{[^,]*}}, {{[^@]*}}@4, {{.*}}i32 0, {{i32|i64}} 4, i32 1, i32 0
|
||||
// LNX_17-DAG: [[PREFIX]]RegisterVar(ptr %0, {{.*}}inline_var
|
||||
// LNX_17-NOT: [[PREFIX]]RegisterVar(ptr %0, {{.*}}inline_var2
|
||||
// ALL: ret void
|
||||
|
||||
// Test that we've built a constructor.
|
||||
@ -239,8 +239,8 @@ void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); }
|
||||
|
||||
// In separate mode it calls __[[PREFIX]]RegisterFatBinary(&__[[PREFIX]]_fatbin_wrapper)
|
||||
// HIP only register fat binary once.
|
||||
// HIP: load i8**, i8*** @__hip_gpubin_handle
|
||||
// HIP-NEXT: icmp eq i8** {{.*}}, null
|
||||
// HIP: load ptr, ptr @__hip_gpubin_handle
|
||||
// HIP-NEXT: icmp eq ptr {{.*}}, null
|
||||
// HIP-NEXT: br i1 {{.*}}, label %if, label %exit
|
||||
// HIP: if:
|
||||
// CUDANORDC: call{{.*}}[[PREFIX]]RegisterFatBinary{{.*}}__[[PREFIX]]_fatbin_wrapper
|
||||
@ -253,12 +253,12 @@ void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); }
|
||||
// .. and then calls __[[PREFIX]]_register_globals
|
||||
// HIP-NEXT: br label %exit
|
||||
// HIP: exit:
|
||||
// HIP-NEXT: load i8**, i8*** @__hip_gpubin_handle
|
||||
// HIP-NEXT: load ptr, ptr @__hip_gpubin_handle
|
||||
// CUDANORDC-NEXT: call void @__[[PREFIX]]_register_globals
|
||||
// HIP-NEXT: call void @__[[PREFIX]]_register_globals
|
||||
// * In separate mode we also register a destructor.
|
||||
// CUDANORDC-NEXT: call i32 @atexit(void ()* @__[[PREFIX]]_module_dtor)
|
||||
// HIP-NEXT: call i32 @atexit(void ()* @__[[PREFIX]]_module_dtor)
|
||||
// CUDANORDC-NEXT: call i32 @atexit(ptr @__[[PREFIX]]_module_dtor)
|
||||
// HIP-NEXT: call i32 @atexit(ptr @__[[PREFIX]]_module_dtor)
|
||||
|
||||
// With relocatable device code we call __[[PREFIX]]RegisterLinkedBinary%NVModuleID%
|
||||
// CUDARDC: call{{.*}}__[[PREFIX]]RegisterLinkedBinary[[MODULE_ID]](
|
||||
@ -271,11 +271,11 @@ void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); }
|
||||
// CUDANORDC: load{{.*}}__[[PREFIX]]_gpubin_handle
|
||||
// HIP: load{{.*}}__[[PREFIX]]_gpubin_handle
|
||||
// CUDANORDC-NEXT: call void @__[[PREFIX]]UnregisterFatBinary
|
||||
// HIP-NEXT: icmp ne i8** {{.*}}, null
|
||||
// HIP-NEXT: icmp ne ptr {{.*}}, null
|
||||
// HIP-NEXT: br i1 {{.*}}, label %if, label %exit
|
||||
// HIP: if:
|
||||
// HIP-NEXT: call void @__[[PREFIX]]UnregisterFatBinary
|
||||
// HIP-NEXT: store i8** null, i8*** @__hip_gpubin_handle
|
||||
// HIP-NEXT: store ptr null, ptr @__hip_gpubin_handle
|
||||
// HIP-NEXT: br label %exit
|
||||
// HIP: exit:
|
||||
|
||||
|
||||
@ -1,19 +1,19 @@
|
||||
// REQUIRES: x86-registered-target
|
||||
// REQUIRES: amdgpu-registered-target
|
||||
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa -fcuda-is-device \
|
||||
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
|
||||
// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o %t.nocuid.dev -x hip %s
|
||||
// RUN: cat %t.nocuid.dev | FileCheck -check-prefixes=DEV,INT-DEV %s
|
||||
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-gnu-linux \
|
||||
// RUN: %clang_cc1 -triple x86_64-gnu-linux \
|
||||
// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o %t.nocuid.host -x hip %s
|
||||
// RUN: cat %t.nocuid.host | FileCheck -check-prefixes=HOST,INT-HOST %s
|
||||
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa -fcuda-is-device -cuid=abc \
|
||||
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -cuid=abc \
|
||||
// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s > %t.dev
|
||||
// RUN: cat %t.dev | FileCheck -check-prefixes=DEV,EXT-DEV %s
|
||||
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-gnu-linux -cuid=abc \
|
||||
// RUN: %clang_cc1 -triple x86_64-gnu-linux -cuid=abc \
|
||||
// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s > %t.host
|
||||
// RUN: cat %t.host | FileCheck -check-prefixes=HOST,EXT-HOST %s
|
||||
|
||||
@ -25,25 +25,25 @@
|
||||
|
||||
// Negative tests.
|
||||
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa -fcuda-is-device \
|
||||
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
|
||||
// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s | FileCheck \
|
||||
// RUN: -check-prefix=DEV-NEG %s
|
||||
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-gnu-linux \
|
||||
// RUN: %clang_cc1 -triple x86_64-gnu-linux \
|
||||
// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s | FileCheck \
|
||||
// RUN: -check-prefix=HOST-NEG %s
|
||||
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa -fcuda-is-device -cuid=abc \
|
||||
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -cuid=abc \
|
||||
// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s > %t.dev
|
||||
// RUN: cat %t.dev | FileCheck -check-prefix=DEV-NEG %s
|
||||
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-gnu-linux -cuid=abc \
|
||||
// RUN: %clang_cc1 -triple x86_64-gnu-linux -cuid=abc \
|
||||
// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s > %t.host
|
||||
// RUN: cat %t.host | FileCheck -check-prefix=HOST-NEG %s
|
||||
|
||||
// Check postfix for CUDA.
|
||||
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple nvptx -fcuda-is-device -cuid=abc \
|
||||
// RUN: %clang_cc1 -triple nvptx -fcuda-is-device -cuid=abc \
|
||||
// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - %s | FileCheck \
|
||||
// RUN: -check-prefixes=CUDA %s
|
||||
|
||||
@ -128,8 +128,8 @@ void foo() {
|
||||
decltype(u) tmp;
|
||||
}
|
||||
|
||||
// HOST-DAG: __hipRegisterVar({{.*}}@_ZL1x {{.*}}@[[DEVNAMEX]]
|
||||
// HOST-DAG: __hipRegisterVar({{.*}}@_ZL1y {{.*}}@[[DEVNAMEY]]
|
||||
// HOST-DAG: __hipRegisterVar({{.*}}@_ZL1x, {{.*}}@[[DEVNAMEX]]
|
||||
// HOST-DAG: __hipRegisterVar({{.*}}@_ZL1y, {{.*}}@[[DEVNAMEY]]
|
||||
// HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZL2x2
|
||||
// HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZZ6kernelPiPPKiE1w
|
||||
// HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZZ6devfunPPKiE1p
|
||||
|
||||
@ -1,13 +1,13 @@
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-unknown -ftrivial-auto-var-init=pattern -ftrivial-auto-var-init-stop-after=1 %s -emit-llvm -o - | FileCheck %s -check-prefix=PATTERN-STOP-AFTER-1-SCALAR
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-unknown -ftrivial-auto-var-init=pattern -ftrivial-auto-var-init-stop-after=2 %s -emit-llvm -o - | FileCheck %s -check-prefix=PATTERN-STOP-AFTER-2-ARRAY
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-unknown -ftrivial-auto-var-init=pattern -ftrivial-auto-var-init-stop-after=3 %s -emit-llvm -o - | FileCheck %s -check-prefix=PATTERN-STOP-AFTER-3-VLA
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-unknown -ftrivial-auto-var-init=pattern -ftrivial-auto-var-init-stop-after=4 %s -emit-llvm -o - | FileCheck %s -check-prefix=PATTERN-STOP-AFTER-4-POINTER
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-unknown -ftrivial-auto-var-init=pattern -ftrivial-auto-var-init-stop-after=5 %s -emit-llvm -o - | FileCheck %s -check-prefix=PATTERN-STOP-AFTER-5-BUILTIN
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-unknown -ftrivial-auto-var-init=zero -ftrivial-auto-var-init-stop-after=1 %s -emit-llvm -o - | FileCheck %s -check-prefix=ZERO-STOP-AFTER-1-SCALAR
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-unknown -ftrivial-auto-var-init=zero -ftrivial-auto-var-init-stop-after=2 %s -emit-llvm -o - | FileCheck %s -check-prefix=ZERO-STOP-AFTER-2-ARRAY
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-unknown -ftrivial-auto-var-init=zero -ftrivial-auto-var-init-stop-after=3 %s -emit-llvm -o - | FileCheck %s -check-prefix=ZERO-STOP-AFTER-3-VLA
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-unknown -ftrivial-auto-var-init=zero -ftrivial-auto-var-init-stop-after=4 %s -emit-llvm -o - | FileCheck %s -check-prefix=ZERO-STOP-AFTER-4-POINTER
|
||||
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-unknown -ftrivial-auto-var-init=zero -ftrivial-auto-var-init-stop-after=5 %s -emit-llvm -o - | FileCheck %s -check-prefix=ZERO-STOP-AFTER-5-BUILTIN
|
||||
// RUN: %clang_cc1 -triple x86_64-unknown-unknown -ftrivial-auto-var-init=pattern -ftrivial-auto-var-init-stop-after=1 %s -emit-llvm -o - | FileCheck %s -check-prefix=PATTERN-STOP-AFTER-1-SCALAR
|
||||
// RUN: %clang_cc1 -triple x86_64-unknown-unknown -ftrivial-auto-var-init=pattern -ftrivial-auto-var-init-stop-after=2 %s -emit-llvm -o - | FileCheck %s -check-prefix=PATTERN-STOP-AFTER-2-ARRAY
|
||||
// RUN: %clang_cc1 -triple x86_64-unknown-unknown -ftrivial-auto-var-init=pattern -ftrivial-auto-var-init-stop-after=3 %s -emit-llvm -o - | FileCheck %s -check-prefix=PATTERN-STOP-AFTER-3-VLA
|
||||
// RUN: %clang_cc1 -triple x86_64-unknown-unknown -ftrivial-auto-var-init=pattern -ftrivial-auto-var-init-stop-after=4 %s -emit-llvm -o - | FileCheck %s -check-prefix=PATTERN-STOP-AFTER-4-POINTER
|
||||
// RUN: %clang_cc1 -triple x86_64-unknown-unknown -ftrivial-auto-var-init=pattern -ftrivial-auto-var-init-stop-after=5 %s -emit-llvm -o - | FileCheck %s -check-prefix=PATTERN-STOP-AFTER-5-BUILTIN
|
||||
// RUN: %clang_cc1 -triple x86_64-unknown-unknown -ftrivial-auto-var-init=zero -ftrivial-auto-var-init-stop-after=1 %s -emit-llvm -o - | FileCheck %s -check-prefix=ZERO-STOP-AFTER-1-SCALAR
|
||||
// RUN: %clang_cc1 -triple x86_64-unknown-unknown -ftrivial-auto-var-init=zero -ftrivial-auto-var-init-stop-after=2 %s -emit-llvm -o - | FileCheck %s -check-prefix=ZERO-STOP-AFTER-2-ARRAY
|
||||
// RUN: %clang_cc1 -triple x86_64-unknown-unknown -ftrivial-auto-var-init=zero -ftrivial-auto-var-init-stop-after=3 %s -emit-llvm -o - | FileCheck %s -check-prefix=ZERO-STOP-AFTER-3-VLA
|
||||
// RUN: %clang_cc1 -triple x86_64-unknown-unknown -ftrivial-auto-var-init=zero -ftrivial-auto-var-init-stop-after=4 %s -emit-llvm -o - | FileCheck %s -check-prefix=ZERO-STOP-AFTER-4-POINTER
|
||||
// RUN: %clang_cc1 -triple x86_64-unknown-unknown -ftrivial-auto-var-init=zero -ftrivial-auto-var-init-stop-after=5 %s -emit-llvm -o - | FileCheck %s -check-prefix=ZERO-STOP-AFTER-5-BUILTIN
|
||||
|
||||
#define ARRLEN 10
|
||||
|
||||
@ -27,32 +27,29 @@ int foo(unsigned n) {
|
||||
void *p;
|
||||
// builtin
|
||||
p = __builtin_alloca(sizeof(unsigned long long) * n);
|
||||
// PATTERN-STOP-AFTER-1-SCALAR: store i64 -6148914691236517206, i64* %a, align 8
|
||||
// PATTERN-STOP-AFTER-1-SCALAR-NOT: call void @llvm.memset.p0i8.i64(i8* align 16 %0, i8 -86, i64 80, i1 false)
|
||||
// PATTERN-STOP-AFTER-2-ARRAY: %0 = bitcast [10 x %struct.S]* %arr to i8*
|
||||
// PATTERN-STOP-AFTER-2-ARRAY-NEXT: call void @llvm.memset.p0i8.i64(i8* align 16 %0, i8 -86, i64 80, i1 false)
|
||||
// PATTERN-STOP-AFTER-1-SCALAR: store i64 -6148914691236517206, ptr %a, align 8
|
||||
// PATTERN-STOP-AFTER-1-SCALAR-NOT: call void @llvm.memset.p0.i64(ptr align 16 %0, i8 -86, i64 80, i1 false)
|
||||
// PATTERN-STOP-AFTER-2-ARRAY: call void @llvm.memset.p0.i64(ptr align 16 %arr, i8 -86, i64 80, i1 false)
|
||||
// PATTERN-STOP-AFTER-2-ARRAY-NOT: vla-init.loop:
|
||||
// PATTERN-STOP-AFTER-3-VLA: vla-init.loop:
|
||||
// PATTERN-STOP-AFTER-3-VLA-NEXT: %vla.cur = phi i8* [ %vla.begin, %vla-setup.loop ], [ %vla.next, %vla-init.loop ]
|
||||
// PATTERN-STOP-AFTER-3-VLA-NEXT-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 %vla.cur, i8* align 4 bitcast ({ i32, i8, [3 x i8] }* @__const._Z3fooj.vla to i8*), i64 8, i1 false)
|
||||
// PATTERN-STOP-AFTER-3-VLA-NOT: store i8* inttoptr (i64 -6148914691236517206 to i8*), i8** %p, align 8
|
||||
// PATTERN-STOP-AFTER-4-POINTER: store i8* inttoptr (i64 -6148914691236517206 to i8*), i8** %p, align 8
|
||||
// PATTERN-STOP-AFTER-4-POINTER-NOT: call void @llvm.memset.p0i8.i64(i8* align 16 %6, i8 -86, i64 %mul, i1 false)
|
||||
// PATTERN-STOP-AFTER-5-BUILTIN: call void @llvm.memset.p0i8.i64(i8* align 16 %6, i8 -86, i64 %mul, i1 false)
|
||||
// PATTERN-STOP-AFTER-3-VLA-NEXT: %vla.cur = phi ptr [ %vla, %vla-setup.loop ], [ %vla.next, %vla-init.loop ]
|
||||
// PATTERN-STOP-AFTER-3-VLA-NEXT-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 %vla.cur, ptr align 4 @__const._Z3fooj.vla, i64 8, i1 false)
|
||||
// PATTERN-STOP-AFTER-3-VLA-NOT: store ptr inttoptr (i64 -6148914691236517206 to ptr), ptr %p, align 8
|
||||
// PATTERN-STOP-AFTER-4-POINTER: store ptr inttoptr (i64 -6148914691236517206 to ptr), ptr %p, align 8
|
||||
// PATTERN-STOP-AFTER-4-POINTER-NOT: call void @llvm.memset.p0.i64(ptr align 16 %5, i8 -86, i64 %mul, i1 false)
|
||||
// PATTERN-STOP-AFTER-5-BUILTIN: call void @llvm.memset.p0.i64(ptr align 16 %5, i8 -86, i64 %mul, i1 false)
|
||||
// PATTERN-STOP-AFTER-5-BUILTIN-MESSAGES: warning: -ftrivial-auto-var-init-stop-after=5 has been enabled to limit the number of times ftrivial-auto-var-init=pattern gets applied.
|
||||
|
||||
// ZERO-STOP-AFTER-1-SCALAR: store i64 0, i64* %a, align 8
|
||||
// ZERO-STOP-AFTER-1-SCALAR-NOT: call void @llvm.memset.p0i8.i64(i8* align 16 %0, i8 0, i64 80, i1 false)
|
||||
// ZERO-STOP-AFTER-2-ARRAY: %0 = bitcast [10 x %struct.S]* %arr to i8*
|
||||
// ZERO-STOP-AFTER-2-ARRAY-NEXT: call void @llvm.memset.p0i8.i64(i8* align 16 %0, i8 0, i64 80, i1 false)
|
||||
// ZERO-STOP-AFTER-2-ARRAY-NOT: %call void @llvm.memset.p0i8.i64(i8* align 16 %5, i8 0, i64 %4, i1 false)
|
||||
// ZERO-STOP-AFTER-3-VLA: %5 = bitcast %struct.S* %vla to i8*
|
||||
// ZERO-STOP-AFTER-3-VLA-NEXT: call void @llvm.memset.p0i8.i64(i8* align 16 %5, i8 0, i64 %4, i1 false)
|
||||
// ZERO-STOP-AFTER-3-VLA-NOT: store i8* null, i8** %p, align 8
|
||||
// ZERO-STOP-AFTER-4-POINTER: store i8* null, i8** %p, align 8
|
||||
// ZERO-STOP-AFTER-4-POINTER-NOT: call void @llvm.memset.p0i8.i64(i8* align 16 %7, i8 0, i64 %mul, i1 false)
|
||||
// ZERO-STOP-AFTER-5-BUILTIN: %7 = alloca i8, i64 %mul, align 16
|
||||
// ZERO-STOP-AFTER-5-BUILTIN-NEXT: call void @llvm.memset.p0i8.i64(i8* align 16 %7, i8 0, i64 %mul, i1 false)
|
||||
// ZERO-STOP-AFTER-1-SCALAR: store i64 0, ptr %a, align 8
|
||||
// ZERO-STOP-AFTER-1-SCALAR-NOT: call void @llvm.memset.p0.i64(ptr align 16 %arr, i8 0, i64 80, i1 false)
|
||||
// ZERO-STOP-AFTER-2-ARRAY: call void @llvm.memset.p0.i64(ptr align 16 %arr, i8 0, i64 80, i1 false)
|
||||
// ZERO-STOP-AFTER-2-ARRAY-NOT: %call void @llvm.memset.p0.i64(ptr align 16 %3, i8 0, i64 %2, i1 false)
|
||||
// ZERO-STOP-AFTER-3-VLA: call void @llvm.memset.p0.i64(ptr align 16 %vla, i8 0, i64 %3, i1 false)
|
||||
// ZERO-STOP-AFTER-3-VLA-NOT: store ptr null, ptr %p, align 8
|
||||
// ZERO-STOP-AFTER-4-POINTER: store ptr null, ptr %p, align 8
|
||||
// ZERO-STOP-AFTER-4-POINTER-NOT: call void @llvm.memset.p0.i64(ptr align 16 %4, i8 0, i64 %mul, i1 false)
|
||||
// ZERO-STOP-AFTER-5-BUILTIN: %5 = alloca i8, i64 %mul, align 16
|
||||
// ZERO-STOP-AFTER-5-BUILTIN-NEXT: call void @llvm.memset.p0.i64(ptr align 16 %5, i8 0, i64 %mul, i1 false)
|
||||
// ZERO-STOP-AFTER-5-BUILTIN-MESSAGES: warnings: -ftrivial-auto-var-init-stop-after=5 has been enabled to limit the number of times ftrivial-auto-var-init=zero gets applied.
|
||||
return 0;
|
||||
}
|
||||
|
||||
Loading…
x
Reference in New Issue
Block a user