Author: Nikita Popov Date: 2023-02-17T15:08:50+01:00 New Revision: 06621ecdaf416d87968dcaf57a634c83cf0dc1b8
URL: https://github.com/llvm/llvm-project/commit/06621ecdaf416d87968dcaf57a634c83cf0dc1b8 DIFF: https://github.com/llvm/llvm-project/commit/06621ecdaf416d87968dcaf57a634c83cf0dc1b8.diff LOG: [Clang] Convert some tests to opaque pointers (NFC) Added: Modified: clang/test/CodeGen/math-builtins.c clang/test/CodeGen/math-libcalls.c clang/test/CodeGen/mult-alt-generic.c clang/test/CodeGenCUDA/device-stub.cu clang/test/CodeGenCUDA/static-device-var-rdc.cu clang/test/CodeGenCXX/auto-var-init-stop-after.cpp Removed: ################################################################################ diff --git a/clang/test/CodeGen/math-builtins.c b/clang/test/CodeGen/math-builtins.c index 559421a4882c5..04738cce01073 100644 --- a/clang/test/CodeGen/math-builtins.c +++ b/clang/test/CodeGen/math-builtins.c @@ -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); diff --git a/clang/test/CodeGen/math-libcalls.c b/clang/test/CodeGen/math-libcalls.c index 0e61f92f8c751..fa8f49d8a2c9f 100644 --- a/clang/test/CodeGen/math-libcalls.c +++ b/clang/test/CodeGen/math-libcalls.c @@ -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); diff --git a/clang/test/CodeGen/mult-alt-generic.c b/clang/test/CodeGen/mult-alt-generic.c index 236acf27cd00f..581200b55da8a 100644 --- a/clang/test/CodeGen/mult-alt-generic.c +++ b/clang/test/CodeGen/mult-alt-generic.c @@ -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 diff erently on diff erent 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 diff erently on diff erent 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)); } diff --git a/clang/test/CodeGenCUDA/device-stub.cu b/clang/test/CodeGenCUDA/device-stub.cu index 8605df8a37702..d7a7b1bb9fe95 100644 --- a/clang/test/CodeGenCUDA/device-stub.cu +++ b/clang/test/CodeGenCUDA/device-stub.cu @@ -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: diff --git a/clang/test/CodeGenCUDA/static-device-var-rdc.cu b/clang/test/CodeGenCUDA/static-device-var-rdc.cu index a9cca4e9212ab..16ec413397235 100644 --- a/clang/test/CodeGenCUDA/static-device-var-rdc.cu +++ b/clang/test/CodeGenCUDA/static-device-var-rdc.cu @@ -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 diff --git a/clang/test/CodeGenCXX/auto-var-init-stop-after.cpp b/clang/test/CodeGenCXX/auto-var-init-stop-after.cpp index b87af1647fb7d..a782692d0127e 100644 --- a/clang/test/CodeGenCXX/auto-var-init-stop-after.cpp +++ b/clang/test/CodeGenCXX/auto-var-init-stop-after.cpp @@ -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; } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
