| // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 2 |
| // REQUIRES: amdgpu-registered-target |
| // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -emit-llvm -mprintf-kind=buffered -fcuda-is-device \ |
| // RUN: -o - %s | FileCheck --enable-var-scope %s |
| // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -emit-llvm -ffp-exception-behavior=strict -mprintf-kind=buffered -fcuda-is-device \ |
| // RUN: -o - %s | FileCheck --enable-var-scope --check-prefix=CHECK_CONSTRAINED %s |
| |
| #define __device__ __attribute__((device)) |
| #define __shared__ __attribute__((shared)) |
| #define __constant__ __attribute__((constant)) |
| |
| extern "C" __device__ int printf(const char *format, ...); |
| |
| // CHECK-LABEL: define dso_local noundef i32 @_Z4foo1v |
| // CHECK-NEXT: entry: |
| // CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5) |
| // CHECK-NEXT: [[S:%.*]] = alloca ptr, align 8, addrspace(5) |
| // CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr |
| // CHECK-NEXT: [[S_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S]] to ptr |
| // CHECK-NEXT: store ptr addrspacecast (ptr addrspace(4) @.str to ptr), ptr [[S_ASCAST]], align 8 |
| // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[S_ASCAST]], align 8 |
| // CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[S_ASCAST]], align 8 |
| // CHECK-NEXT: [[TMP2:%.*]] = icmp eq ptr [[TMP0]], null |
| // CHECK-NEXT: br i1 [[TMP2]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]] |
| // CHECK: strlen.while: |
| // CHECK-NEXT: [[TMP3:%.*]] = phi ptr [ [[TMP0]], [[ENTRY:%.*]] ], [ [[TMP4:%.*]], [[STRLEN_WHILE]] ] |
| // CHECK-NEXT: [[TMP4]] = getelementptr i8, ptr [[TMP3]], i64 1 |
| // CHECK-NEXT: [[TMP5:%.*]] = load i8, ptr [[TMP3]], align 1 |
| // CHECK-NEXT: [[TMP6:%.*]] = icmp eq i8 [[TMP5]], 0 |
| // CHECK-NEXT: br i1 [[TMP6]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]] |
| // CHECK: strlen.while.done: |
| // CHECK-NEXT: [[TMP7:%.*]] = ptrtoint ptr [[TMP0]] to i64 |
| // CHECK-NEXT: [[TMP8:%.*]] = ptrtoint ptr [[TMP3]] to i64 |
| // CHECK-NEXT: [[TMP9:%.*]] = sub i64 [[TMP8]], [[TMP7]] |
| // CHECK-NEXT: [[TMP10:%.*]] = add i64 [[TMP9]], 1 |
| // CHECK-NEXT: br label [[STRLEN_JOIN]] |
| // CHECK: strlen.join: |
| // CHECK-NEXT: [[TMP11:%.*]] = phi i64 [ [[TMP10]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ] |
| // CHECK-NEXT: [[TMP12:%.*]] = add i64 [[TMP11]], 7 |
| // CHECK-NEXT: [[TMP13:%.*]] = and i64 [[TMP12]], 4294967288 |
| // CHECK-NEXT: [[TMP14:%.*]] = add i64 [[TMP13]], 52 |
| // CHECK-NEXT: [[TMP15:%.*]] = trunc i64 [[TMP14]] to i32 |
| // CHECK-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP15]]) |
| // CHECK-NEXT: [[TMP16:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null |
| // CHECK-NEXT: br i1 [[TMP16]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]] |
| // CHECK: end.block: |
| // CHECK-NEXT: [[TMP17:%.*]] = xor i1 [[TMP16]], true |
| // CHECK-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP17]] to i32 |
| // CHECK-NEXT: ret i32 [[PRINTF_RESULT]] |
| // CHECK: argpush.block: |
| // CHECK-NEXT: [[TMP18:%.*]] = shl i32 [[TMP15]], 2 |
| // CHECK-NEXT: [[TMP19:%.*]] = or i32 [[TMP18]], 2 |
| // CHECK-NEXT: store i32 [[TMP19]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4 |
| // CHECK-NEXT: [[TMP20:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4 |
| // CHECK-NEXT: store i64 1107004088646384690, ptr addrspace(1) [[TMP20]], align 8 |
| // CHECK-NEXT: [[TMP21:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP20]], i32 8 |
| // CHECK-NEXT: store i64 8, ptr addrspace(1) [[TMP21]], align 8 |
| // CHECK-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP21]], i32 8 |
| // CHECK-NEXT: store double 3.141590e+00, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8 |
| // CHECK-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8 |
| // CHECK-NEXT: store i64 8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], align 8 |
| // CHECK-NEXT: [[PRINTBUFFNEXTPTR2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], i32 8 |
| // CHECK-NEXT: store i64 4, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], align 8 |
| // CHECK-NEXT: [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], i32 8 |
| // CHECK-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[PRINTBUFFNEXTPTR3]], ptr align 1 [[TMP0]], i64 [[TMP11]], i1 false) |
| // CHECK-NEXT: [[PRINTBUFFNEXTPTR4:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], i64 [[TMP13]] |
| // CHECK-NEXT: store ptr [[TMP1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], align 8 |
| // CHECK-NEXT: [[PRINTBUFFNEXTPTR5:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], i32 8 |
| // CHECK-NEXT: br label [[END_BLOCK]] |
| // |
| // CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z4foo1v |
| // CHECK_CONSTRAINED-NEXT: entry: |
| // CHECK_CONSTRAINED-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5) |
| // CHECK_CONSTRAINED-NEXT: [[S:%.*]] = alloca ptr, align 8, addrspace(5) |
| // CHECK_CONSTRAINED-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr |
| // CHECK_CONSTRAINED-NEXT: [[S_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S]] to ptr |
| // CHECK_CONSTRAINED-NEXT: store ptr addrspacecast (ptr addrspace(4) @.str to ptr), ptr [[S_ASCAST]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[TMP0:%.*]] = load ptr, ptr [[S_ASCAST]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[TMP1:%.*]] = load ptr, ptr [[S_ASCAST]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[TMP2:%.*]] = icmp eq ptr [[TMP0]], null |
| // CHECK_CONSTRAINED-NEXT: br i1 [[TMP2]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]] |
| // CHECK_CONSTRAINED: strlen.while: |
| // CHECK_CONSTRAINED-NEXT: [[TMP3:%.*]] = phi ptr [ [[TMP0]], [[ENTRY:%.*]] ], [ [[TMP4:%.*]], [[STRLEN_WHILE]] ] |
| // CHECK_CONSTRAINED-NEXT: [[TMP4]] = getelementptr i8, ptr [[TMP3]], i64 1 |
| // CHECK_CONSTRAINED-NEXT: [[TMP5:%.*]] = load i8, ptr [[TMP3]], align 1 |
| // CHECK_CONSTRAINED-NEXT: [[TMP6:%.*]] = icmp eq i8 [[TMP5]], 0 |
| // CHECK_CONSTRAINED-NEXT: br i1 [[TMP6]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]] |
| // CHECK_CONSTRAINED: strlen.while.done: |
| // CHECK_CONSTRAINED-NEXT: [[TMP7:%.*]] = ptrtoint ptr [[TMP0]] to i64 |
| // CHECK_CONSTRAINED-NEXT: [[TMP8:%.*]] = ptrtoint ptr [[TMP3]] to i64 |
| // CHECK_CONSTRAINED-NEXT: [[TMP9:%.*]] = sub i64 [[TMP8]], [[TMP7]] |
| // CHECK_CONSTRAINED-NEXT: [[TMP10:%.*]] = add i64 [[TMP9]], 1 |
| // CHECK_CONSTRAINED-NEXT: br label [[STRLEN_JOIN]] |
| // CHECK_CONSTRAINED: strlen.join: |
| // CHECK_CONSTRAINED-NEXT: [[TMP11:%.*]] = phi i64 [ [[TMP10]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ] |
| // CHECK_CONSTRAINED-NEXT: [[TMP12:%.*]] = add i64 [[TMP11]], 7 |
| // CHECK_CONSTRAINED-NEXT: [[TMP13:%.*]] = and i64 [[TMP12]], 4294967288 |
| // CHECK_CONSTRAINED-NEXT: [[TMP14:%.*]] = add i64 [[TMP13]], 52 |
| // CHECK_CONSTRAINED-NEXT: [[TMP15:%.*]] = trunc i64 [[TMP14]] to i32 |
| // CHECK_CONSTRAINED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP15]]) |
| // CHECK_CONSTRAINED-NEXT: [[TMP16:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null |
| // CHECK_CONSTRAINED-NEXT: br i1 [[TMP16]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]] |
| // CHECK_CONSTRAINED: end.block: |
| // CHECK_CONSTRAINED-NEXT: [[TMP17:%.*]] = xor i1 [[TMP16]], true |
| // CHECK_CONSTRAINED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP17]] to i32 |
| // CHECK_CONSTRAINED-NEXT: ret i32 [[PRINTF_RESULT]] |
| // CHECK_CONSTRAINED: argpush.block: |
| // CHECK_CONSTRAINED-NEXT: [[TMP18:%.*]] = shl i32 [[TMP15]], 2 |
| // CHECK_CONSTRAINED-NEXT: [[TMP19:%.*]] = or i32 [[TMP18]], 2 |
| // CHECK_CONSTRAINED-NEXT: store i32 [[TMP19]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4 |
| // CHECK_CONSTRAINED-NEXT: [[TMP20:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4 |
| // CHECK_CONSTRAINED-NEXT: store i64 1107004088646384690, ptr addrspace(1) [[TMP20]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[TMP21:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP20]], i32 8 |
| // CHECK_CONSTRAINED-NEXT: store i64 8, ptr addrspace(1) [[TMP21]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP21]], i32 8 |
| // CHECK_CONSTRAINED-NEXT: store double 3.141590e+00, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8 |
| // CHECK_CONSTRAINED-NEXT: store i64 8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], i32 8 |
| // CHECK_CONSTRAINED-NEXT: store i64 4, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], i32 8 |
| // CHECK_CONSTRAINED-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[PRINTBUFFNEXTPTR3]], ptr align 1 [[TMP0]], i64 [[TMP11]], i1 false) |
| // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR4:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], i64 [[TMP13]] |
| // CHECK_CONSTRAINED-NEXT: store ptr [[TMP1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR5:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], i32 8 |
| // CHECK_CONSTRAINED-NEXT: br label [[END_BLOCK]] |
| // |
| __device__ int foo1() { |
| const char *s = "hello world"; |
| return printf("%.*f %*.*s %p\n", 8, 3.14159, 8, 4, s, s); |
| } |
| |
| __device__ char *dstr; |
| __device__ const |
| // CHECK-LABEL: define dso_local noundef i32 @_Z4foo2v |
| // CHECK-NEXT: entry: |
| // CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5) |
| // CHECK-NEXT: [[LCVAL:%.*]] = alloca i32, align 4, addrspace(5) |
| // CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr |
| // CHECK-NEXT: [[LCVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[LCVAL]] to ptr |
| // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspacecast (ptr addrspace(1) @dstr to ptr), align 8 |
| // CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr addrspacecast (ptr addrspace(1) @dstr to ptr), align 8 |
| // CHECK-NEXT: [[TMP2:%.*]] = icmp eq ptr [[TMP0]], null |
| // CHECK-NEXT: br i1 [[TMP2]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]] |
| // CHECK: strlen.while: |
| // CHECK-NEXT: [[TMP3:%.*]] = phi ptr [ [[TMP0]], [[ENTRY:%.*]] ], [ [[TMP4:%.*]], [[STRLEN_WHILE]] ] |
| // CHECK-NEXT: [[TMP4]] = getelementptr i8, ptr [[TMP3]], i64 1 |
| // CHECK-NEXT: [[TMP5:%.*]] = load i8, ptr [[TMP3]], align 1 |
| // CHECK-NEXT: [[TMP6:%.*]] = icmp eq i8 [[TMP5]], 0 |
| // CHECK-NEXT: br i1 [[TMP6]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]] |
| // CHECK: strlen.while.done: |
| // CHECK-NEXT: [[TMP7:%.*]] = ptrtoint ptr [[TMP0]] to i64 |
| // CHECK-NEXT: [[TMP8:%.*]] = ptrtoint ptr [[TMP3]] to i64 |
| // CHECK-NEXT: [[TMP9:%.*]] = sub i64 [[TMP8]], [[TMP7]] |
| // CHECK-NEXT: [[TMP10:%.*]] = add i64 [[TMP9]], 1 |
| // CHECK-NEXT: br label [[STRLEN_JOIN]] |
| // CHECK: strlen.join: |
| // CHECK-NEXT: [[TMP11:%.*]] = phi i64 [ [[TMP10]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ] |
| // CHECK-NEXT: [[TMP12:%.*]] = add i64 [[TMP11]], 7 |
| // CHECK-NEXT: [[TMP13:%.*]] = and i64 [[TMP12]], 4294967288 |
| // CHECK-NEXT: [[TMP14:%.*]] = add i64 [[TMP13]], 36 |
| // CHECK-NEXT: [[TMP15:%.*]] = trunc i64 [[TMP14]] to i32 |
| // CHECK-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP15]]) |
| // CHECK-NEXT: [[TMP16:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null |
| // CHECK-NEXT: br i1 [[TMP16]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]] |
| // CHECK: end.block: |
| // CHECK-NEXT: [[TMP17:%.*]] = xor i1 [[TMP16]], true |
| // CHECK-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP17]] to i32 |
| // CHECK-NEXT: ret i32 [[PRINTF_RESULT]] |
| // CHECK: argpush.block: |
| // CHECK-NEXT: [[TMP18:%.*]] = shl i32 [[TMP15]], 2 |
| // CHECK-NEXT: [[TMP19:%.*]] = or i32 [[TMP18]], 2 |
| // CHECK-NEXT: store i32 [[TMP19]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4 |
| // CHECK-NEXT: [[TMP20:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4 |
| // CHECK-NEXT: store i64 7257695813269076350, ptr addrspace(1) [[TMP20]], align 8 |
| // CHECK-NEXT: [[TMP21:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP20]], i32 8 |
| // CHECK-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[TMP21]], ptr align 1 [[TMP0]], i64 [[TMP11]], i1 false) |
| // CHECK-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP21]], i64 [[TMP13]] |
| // CHECK-NEXT: store ptr [[TMP1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8 |
| // CHECK-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8 |
| // CHECK-NEXT: store ptr addrspacecast (ptr addrspace(3) @_ZZ4foo2vE5shval to ptr), ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], align 8 |
| // CHECK-NEXT: [[PRINTBUFFNEXTPTR2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], i32 8 |
| // CHECK-NEXT: store ptr [[LCVAL_ASCAST]], ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], align 8 |
| // CHECK-NEXT: [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], i32 8 |
| // CHECK-NEXT: br label [[END_BLOCK]] |
| // |
| // CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z4foo2v |
| // CHECK_CONSTRAINED-NEXT: entry: |
| // CHECK_CONSTRAINED-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5) |
| // CHECK_CONSTRAINED-NEXT: [[LCVAL:%.*]] = alloca i32, align 4, addrspace(5) |
| // CHECK_CONSTRAINED-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr |
| // CHECK_CONSTRAINED-NEXT: [[LCVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[LCVAL]] to ptr |
| // CHECK_CONSTRAINED-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspacecast (ptr addrspace(1) @dstr to ptr), align 8 |
| // CHECK_CONSTRAINED-NEXT: [[TMP1:%.*]] = load ptr, ptr addrspacecast (ptr addrspace(1) @dstr to ptr), align 8 |
| // CHECK_CONSTRAINED-NEXT: [[TMP2:%.*]] = icmp eq ptr [[TMP0]], null |
| // CHECK_CONSTRAINED-NEXT: br i1 [[TMP2]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]] |
| // CHECK_CONSTRAINED: strlen.while: |
| // CHECK_CONSTRAINED-NEXT: [[TMP3:%.*]] = phi ptr [ [[TMP0]], [[ENTRY:%.*]] ], [ [[TMP4:%.*]], [[STRLEN_WHILE]] ] |
| // CHECK_CONSTRAINED-NEXT: [[TMP4]] = getelementptr i8, ptr [[TMP3]], i64 1 |
| // CHECK_CONSTRAINED-NEXT: [[TMP5:%.*]] = load i8, ptr [[TMP3]], align 1 |
| // CHECK_CONSTRAINED-NEXT: [[TMP6:%.*]] = icmp eq i8 [[TMP5]], 0 |
| // CHECK_CONSTRAINED-NEXT: br i1 [[TMP6]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]] |
| // CHECK_CONSTRAINED: strlen.while.done: |
| // CHECK_CONSTRAINED-NEXT: [[TMP7:%.*]] = ptrtoint ptr [[TMP0]] to i64 |
| // CHECK_CONSTRAINED-NEXT: [[TMP8:%.*]] = ptrtoint ptr [[TMP3]] to i64 |
| // CHECK_CONSTRAINED-NEXT: [[TMP9:%.*]] = sub i64 [[TMP8]], [[TMP7]] |
| // CHECK_CONSTRAINED-NEXT: [[TMP10:%.*]] = add i64 [[TMP9]], 1 |
| // CHECK_CONSTRAINED-NEXT: br label [[STRLEN_JOIN]] |
| // CHECK_CONSTRAINED: strlen.join: |
| // CHECK_CONSTRAINED-NEXT: [[TMP11:%.*]] = phi i64 [ [[TMP10]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ] |
| // CHECK_CONSTRAINED-NEXT: [[TMP12:%.*]] = add i64 [[TMP11]], 7 |
| // CHECK_CONSTRAINED-NEXT: [[TMP13:%.*]] = and i64 [[TMP12]], 4294967288 |
| // CHECK_CONSTRAINED-NEXT: [[TMP14:%.*]] = add i64 [[TMP13]], 36 |
| // CHECK_CONSTRAINED-NEXT: [[TMP15:%.*]] = trunc i64 [[TMP14]] to i32 |
| // CHECK_CONSTRAINED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP15]]) |
| // CHECK_CONSTRAINED-NEXT: [[TMP16:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null |
| // CHECK_CONSTRAINED-NEXT: br i1 [[TMP16]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]] |
| // CHECK_CONSTRAINED: end.block: |
| // CHECK_CONSTRAINED-NEXT: [[TMP17:%.*]] = xor i1 [[TMP16]], true |
| // CHECK_CONSTRAINED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP17]] to i32 |
| // CHECK_CONSTRAINED-NEXT: ret i32 [[PRINTF_RESULT]] |
| // CHECK_CONSTRAINED: argpush.block: |
| // CHECK_CONSTRAINED-NEXT: [[TMP18:%.*]] = shl i32 [[TMP15]], 2 |
| // CHECK_CONSTRAINED-NEXT: [[TMP19:%.*]] = or i32 [[TMP18]], 2 |
| // CHECK_CONSTRAINED-NEXT: store i32 [[TMP19]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4 |
| // CHECK_CONSTRAINED-NEXT: [[TMP20:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4 |
| // CHECK_CONSTRAINED-NEXT: store i64 7257695813269076350, ptr addrspace(1) [[TMP20]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[TMP21:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP20]], i32 8 |
| // CHECK_CONSTRAINED-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[TMP21]], ptr align 1 [[TMP0]], i64 [[TMP11]], i1 false) |
| // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP21]], i64 [[TMP13]] |
| // CHECK_CONSTRAINED-NEXT: store ptr [[TMP1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8 |
| // CHECK_CONSTRAINED-NEXT: store ptr addrspacecast (ptr addrspace(3) @_ZZ4foo2vE5shval to ptr), ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], i32 8 |
| // CHECK_CONSTRAINED-NEXT: store ptr [[LCVAL_ASCAST]], ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], i32 8 |
| // CHECK_CONSTRAINED-NEXT: br label [[END_BLOCK]] |
| // |
| __device__ int foo2() { |
| __shared__ int shval; |
| int lcval; |
| return printf("%s %p %p %p\n", dstr, dstr, &shval, &lcval); |
| } |
| |
| __device__ unsigned short g = 30; |
| __device__ unsigned long n = 30; |
| |
| __device__ float f1 = 3.14f; |
| __device__ double f2 = 2.71828; |
| __device__ _Float16 f3 = 2.71; |
| __device__ __bf16 f4 = 3.142; |
| __device__ _BitInt(55) Int55 = 31; |
| __device__ _BitInt(44) Int44 = 312; |
| __device__ _BitInt(128) Int128 = 45637; |
| |
| // CHECK-LABEL: define dso_local noundef i32 @_Z4foo3v |
| // CHECK-NEXT: entry: |
| // CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5) |
| // CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr |
| // CHECK-NEXT: store i32 25, ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), align 4 |
| // CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), align 4 |
| // CHECK-NEXT: [[TMP1:%.*]] = load i16, ptr addrspacecast (ptr addrspace(1) @g to ptr), align 2 |
| // CHECK-NEXT: [[CONV:%.*]] = zext i16 [[TMP1]] to i32 |
| // CHECK-NEXT: [[TMP2:%.*]] = load i64, ptr addrspacecast (ptr addrspace(1) @n to ptr), align 8 |
| // CHECK-NEXT: [[TMP3:%.*]] = load float, ptr addrspacecast (ptr addrspace(1) @f1 to ptr), align 4 |
| // CHECK-NEXT: [[CONV1:%.*]] = fpext float [[TMP3]] to double |
| // CHECK-NEXT: [[TMP4:%.*]] = load double, ptr addrspacecast (ptr addrspace(1) @f2 to ptr), align 8 |
| // CHECK-NEXT: [[TMP5:%.*]] = load half, ptr addrspacecast (ptr addrspace(1) @f3 to ptr), align 2 |
| // CHECK-NEXT: [[TMP6:%.*]] = load bfloat, ptr addrspacecast (ptr addrspace(1) @f4 to ptr), align 2 |
| // CHECK-NEXT: [[TMP7:%.*]] = load i55, ptr addrspacecast (ptr addrspace(1) @Int55 to ptr), align 8 |
| // CHECK-NEXT: [[TMP8:%.*]] = load i44, ptr addrspacecast (ptr addrspace(1) @Int44 to ptr), align 8 |
| // CHECK-NEXT: [[TMP9:%.*]] = load i128, ptr addrspacecast (ptr addrspace(1) @Int128 to ptr), align 8 |
| // CHECK-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 108) |
| // CHECK-NEXT: [[TMP10:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null |
| // CHECK-NEXT: br i1 [[TMP10]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]] |
| // CHECK: end.block: |
| // CHECK-NEXT: [[TMP11:%.*]] = xor i1 [[TMP10]], true |
| // CHECK-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP11]] to i32 |
| // CHECK-NEXT: ret i32 [[PRINTF_RESULT]] |
| // CHECK: argpush.block: |
| // CHECK-NEXT: store i32 434, ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4 |
| // CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4 |
| // CHECK-NEXT: store i64 7271852820361268873, ptr addrspace(1) [[TMP12]], align 8 |
| // CHECK-NEXT: [[TMP13:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP12]], i32 8 |
| // CHECK-NEXT: [[TMP14:%.*]] = zext i32 [[TMP0]] to i64 |
| // CHECK-NEXT: store i64 [[TMP14]], ptr addrspace(1) [[TMP13]], align 8 |
| // CHECK-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP13]], i32 8 |
| // CHECK-NEXT: store ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8 |
| // CHECK-NEXT: [[PRINTBUFFNEXTPTR2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8 |
| // CHECK-NEXT: [[TMP15:%.*]] = zext i32 [[CONV]] to i64 |
| // CHECK-NEXT: store i64 [[TMP15]], ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], align 8 |
| // CHECK-NEXT: [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], i32 8 |
| // CHECK-NEXT: store i64 [[TMP2]], ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], align 8 |
| // CHECK-NEXT: [[PRINTBUFFNEXTPTR4:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], i32 8 |
| // CHECK-NEXT: store double [[CONV1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], align 8 |
| // CHECK-NEXT: [[PRINTBUFFNEXTPTR5:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], i32 8 |
| // CHECK-NEXT: store double [[TMP4]], ptr addrspace(1) [[PRINTBUFFNEXTPTR5]], align 8 |
| // CHECK-NEXT: [[PRINTBUFFNEXTPTR6:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR5]], i32 8 |
| // CHECK-NEXT: [[TMP16:%.*]] = fpext half [[TMP5]] to double |
| // CHECK-NEXT: store double [[TMP16]], ptr addrspace(1) [[PRINTBUFFNEXTPTR6]], align 8 |
| // CHECK-NEXT: [[PRINTBUFFNEXTPTR7:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR6]], i32 8 |
| // CHECK-NEXT: [[TMP17:%.*]] = fpext bfloat [[TMP6]] to double |
| // CHECK-NEXT: store double [[TMP17]], ptr addrspace(1) [[PRINTBUFFNEXTPTR7]], align 8 |
| // CHECK-NEXT: [[PRINTBUFFNEXTPTR8:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR7]], i32 8 |
| // CHECK-NEXT: [[TMP18:%.*]] = zext i55 [[TMP7]] to i64 |
| // CHECK-NEXT: store i64 [[TMP18]], ptr addrspace(1) [[PRINTBUFFNEXTPTR8]], align 8 |
| // CHECK-NEXT: [[PRINTBUFFNEXTPTR9:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR8]], i32 8 |
| // CHECK-NEXT: [[TMP19:%.*]] = zext i44 [[TMP8]] to i64 |
| // CHECK-NEXT: store i64 [[TMP19]], ptr addrspace(1) [[PRINTBUFFNEXTPTR9]], align 8 |
| // CHECK-NEXT: [[PRINTBUFFNEXTPTR10:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR9]], i32 8 |
| // CHECK-NEXT: store i128 [[TMP9]], ptr addrspace(1) [[PRINTBUFFNEXTPTR10]], align 8 |
| // CHECK-NEXT: [[PRINTBUFFNEXTPTR11:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR10]], i32 16 |
| // CHECK-NEXT: br label [[END_BLOCK]] |
| // |
| // CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z4foo3v |
| // CHECK_CONSTRAINED-NEXT: entry: |
| // CHECK_CONSTRAINED-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5) |
| // CHECK_CONSTRAINED-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr |
| // CHECK_CONSTRAINED-NEXT: store i32 25, ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), align 4 |
| // CHECK_CONSTRAINED-NEXT: [[TMP0:%.*]] = load i32, ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), align 4 |
| // CHECK_CONSTRAINED-NEXT: [[TMP1:%.*]] = load i16, ptr addrspacecast (ptr addrspace(1) @g to ptr), align 2 |
| // CHECK_CONSTRAINED-NEXT: [[CONV:%.*]] = zext i16 [[TMP1]] to i32 |
| // CHECK_CONSTRAINED-NEXT: [[TMP2:%.*]] = load i64, ptr addrspacecast (ptr addrspace(1) @n to ptr), align 8 |
| // CHECK_CONSTRAINED-NEXT: [[TMP3:%.*]] = load float, ptr addrspacecast (ptr addrspace(1) @f1 to ptr), align 4 |
| // CHECK_CONSTRAINED-NEXT: [[CONV1:%.*]] = fpext float [[TMP3]] to double |
| // CHECK_CONSTRAINED-NEXT: [[TMP4:%.*]] = load double, ptr addrspacecast (ptr addrspace(1) @f2 to ptr), align 8 |
| // CHECK_CONSTRAINED-NEXT: [[TMP5:%.*]] = load half, ptr addrspacecast (ptr addrspace(1) @f3 to ptr), align 2 |
| // CHECK_CONSTRAINED-NEXT: [[TMP6:%.*]] = load bfloat, ptr addrspacecast (ptr addrspace(1) @f4 to ptr), align 2 |
| // CHECK_CONSTRAINED-NEXT: [[TMP7:%.*]] = load i55, ptr addrspacecast (ptr addrspace(1) @Int55 to ptr), align 8 |
| // CHECK_CONSTRAINED-NEXT: [[TMP8:%.*]] = load i44, ptr addrspacecast (ptr addrspace(1) @Int44 to ptr), align 8 |
| // CHECK_CONSTRAINED-NEXT: [[TMP9:%.*]] = load i128, ptr addrspacecast (ptr addrspace(1) @Int128 to ptr), align 8 |
| // CHECK_CONSTRAINED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 108) |
| // CHECK_CONSTRAINED-NEXT: [[TMP10:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null |
| // CHECK_CONSTRAINED-NEXT: br i1 [[TMP10]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]] |
| // CHECK_CONSTRAINED: end.block: |
| // CHECK_CONSTRAINED-NEXT: [[TMP11:%.*]] = xor i1 [[TMP10]], true |
| // CHECK_CONSTRAINED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP11]] to i32 |
| // CHECK_CONSTRAINED-NEXT: ret i32 [[PRINTF_RESULT]] |
| // CHECK_CONSTRAINED: argpush.block: |
| // CHECK_CONSTRAINED-NEXT: store i32 434, ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4 |
| // CHECK_CONSTRAINED-NEXT: [[TMP12:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4 |
| // CHECK_CONSTRAINED-NEXT: store i64 7271852820361268873, ptr addrspace(1) [[TMP12]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[TMP13:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP12]], i32 8 |
| // CHECK_CONSTRAINED-NEXT: [[TMP14:%.*]] = zext i32 [[TMP0]] to i64 |
| // CHECK_CONSTRAINED-NEXT: store i64 [[TMP14]], ptr addrspace(1) [[TMP13]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP13]], i32 8 |
| // CHECK_CONSTRAINED-NEXT: store ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8 |
| // CHECK_CONSTRAINED-NEXT: [[TMP15:%.*]] = zext i32 [[CONV]] to i64 |
| // CHECK_CONSTRAINED-NEXT: store i64 [[TMP15]], ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], i32 8 |
| // CHECK_CONSTRAINED-NEXT: store i64 [[TMP2]], ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR4:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], i32 8 |
| // CHECK_CONSTRAINED-NEXT: store double [[CONV1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR5:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], i32 8 |
| // CHECK_CONSTRAINED-NEXT: store double [[TMP4]], ptr addrspace(1) [[PRINTBUFFNEXTPTR5]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR6:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR5]], i32 8 |
| // CHECK_CONSTRAINED-NEXT: [[TMP16:%.*]] = fpext half [[TMP5]] to double |
| // CHECK_CONSTRAINED-NEXT: store double [[TMP16]], ptr addrspace(1) [[PRINTBUFFNEXTPTR6]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR7:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR6]], i32 8 |
| // CHECK_CONSTRAINED-NEXT: [[TMP17:%.*]] = fpext bfloat [[TMP6]] to double |
| // CHECK_CONSTRAINED-NEXT: store double [[TMP17]], ptr addrspace(1) [[PRINTBUFFNEXTPTR7]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR8:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR7]], i32 8 |
| // CHECK_CONSTRAINED-NEXT: [[TMP18:%.*]] = zext i55 [[TMP7]] to i64 |
| // CHECK_CONSTRAINED-NEXT: store i64 [[TMP18]], ptr addrspace(1) [[PRINTBUFFNEXTPTR8]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR9:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR8]], i32 8 |
| // CHECK_CONSTRAINED-NEXT: [[TMP19:%.*]] = zext i44 [[TMP8]] to i64 |
| // CHECK_CONSTRAINED-NEXT: store i64 [[TMP19]], ptr addrspace(1) [[PRINTBUFFNEXTPTR9]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR10:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR9]], i32 8 |
| // CHECK_CONSTRAINED-NEXT: store i128 [[TMP9]], ptr addrspace(1) [[PRINTBUFFNEXTPTR10]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR11:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR10]], i32 16 |
| // CHECK_CONSTRAINED-NEXT: br label [[END_BLOCK]] |
| // |
| __device__ int foo3() { |
| __shared__ int s; |
| s = 25; |
| return printf("Random values: %d,%p,%hd,%ld,%f,%f,%f,%f,%d,%d,%d\n",s, &s, g, n, f1, f2, f3, f4, Int55, Int44, Int128); |
| } |
| |
| //A non trivial case, |
| // CHECK-LABEL: define dso_local noundef i32 @_Z4foo4v |
| // CHECK-NEXT: entry: |
| // CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5) |
| // CHECK-NEXT: [[S:%.*]] = alloca ptr, align 8, addrspace(5) |
| // CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr |
| // CHECK-NEXT: [[S_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S]] to ptr |
| // CHECK-NEXT: store ptr addrspacecast (ptr addrspace(4) @.str.4 to ptr), ptr [[S_ASCAST]], align 8 |
| // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[S_ASCAST]], align 8 |
| // CHECK-NEXT: [[TMP1:%.*]] = icmp eq ptr [[TMP0]], null |
| // CHECK-NEXT: br i1 [[TMP1]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]] |
| // CHECK: strlen.while: |
| // CHECK-NEXT: [[TMP2:%.*]] = phi ptr [ [[TMP0]], [[ENTRY:%.*]] ], [ [[TMP3:%.*]], [[STRLEN_WHILE]] ] |
| // CHECK-NEXT: [[TMP3]] = getelementptr i8, ptr [[TMP2]], i64 1 |
| // CHECK-NEXT: [[TMP4:%.*]] = load i8, ptr [[TMP2]], align 1 |
| // CHECK-NEXT: [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0 |
| // CHECK-NEXT: br i1 [[TMP5]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]] |
| // CHECK: strlen.while.done: |
| // CHECK-NEXT: [[TMP6:%.*]] = ptrtoint ptr [[TMP0]] to i64 |
| // CHECK-NEXT: [[TMP7:%.*]] = ptrtoint ptr [[TMP2]] to i64 |
| // CHECK-NEXT: [[TMP8:%.*]] = sub i64 [[TMP7]], [[TMP6]] |
| // CHECK-NEXT: [[TMP9:%.*]] = add i64 [[TMP8]], 1 |
| // CHECK-NEXT: br label [[STRLEN_JOIN]] |
| // CHECK: strlen.join: |
| // CHECK-NEXT: [[TMP10:%.*]] = phi i64 [ [[TMP9]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ] |
| // CHECK-NEXT: [[TMP11:%.*]] = add i64 [[TMP10]], 7 |
| // CHECK-NEXT: [[TMP12:%.*]] = and i64 [[TMP11]], 4294967288 |
| // CHECK-NEXT: [[TMP13:%.*]] = add i64 [[TMP12]], 12 |
| // CHECK-NEXT: [[TMP14:%.*]] = trunc i64 [[TMP13]] to i32 |
| // CHECK-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP14]]) |
| // CHECK-NEXT: [[TMP15:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null |
| // CHECK-NEXT: br i1 [[TMP15]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]] |
| // CHECK: end.block: |
| // CHECK-NEXT: [[TMP16:%.*]] = xor i1 [[TMP15]], true |
| // CHECK-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP16]] to i32 |
| // CHECK-NEXT: ret i32 [[PRINTF_RESULT]] |
| // CHECK: argpush.block: |
| // CHECK-NEXT: [[TMP17:%.*]] = shl i32 [[TMP14]], 2 |
| // CHECK-NEXT: store i32 [[TMP17]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4 |
| // CHECK-NEXT: [[TMP18:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4 |
| // CHECK-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[TMP18]], ptr align 1 [[TMP0]], i64 [[TMP10]], i1 false) |
| // CHECK-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP18]], i64 [[TMP12]] |
| // CHECK-NEXT: store i64 10, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8 |
| // CHECK-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8 |
| // CHECK-NEXT: br label [[END_BLOCK]] |
| // |
| // CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z4foo4v |
| // CHECK_CONSTRAINED-NEXT: entry: |
| // CHECK_CONSTRAINED-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5) |
| // CHECK_CONSTRAINED-NEXT: [[S:%.*]] = alloca ptr, align 8, addrspace(5) |
| // CHECK_CONSTRAINED-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr |
| // CHECK_CONSTRAINED-NEXT: [[S_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S]] to ptr |
| // CHECK_CONSTRAINED-NEXT: store ptr addrspacecast (ptr addrspace(4) @.str.4 to ptr), ptr [[S_ASCAST]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[TMP0:%.*]] = load ptr, ptr [[S_ASCAST]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[TMP1:%.*]] = icmp eq ptr [[TMP0]], null |
| // CHECK_CONSTRAINED-NEXT: br i1 [[TMP1]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]] |
| // CHECK_CONSTRAINED: strlen.while: |
| // CHECK_CONSTRAINED-NEXT: [[TMP2:%.*]] = phi ptr [ [[TMP0]], [[ENTRY:%.*]] ], [ [[TMP3:%.*]], [[STRLEN_WHILE]] ] |
| // CHECK_CONSTRAINED-NEXT: [[TMP3]] = getelementptr i8, ptr [[TMP2]], i64 1 |
| // CHECK_CONSTRAINED-NEXT: [[TMP4:%.*]] = load i8, ptr [[TMP2]], align 1 |
| // CHECK_CONSTRAINED-NEXT: [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0 |
| // CHECK_CONSTRAINED-NEXT: br i1 [[TMP5]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]] |
| // CHECK_CONSTRAINED: strlen.while.done: |
| // CHECK_CONSTRAINED-NEXT: [[TMP6:%.*]] = ptrtoint ptr [[TMP0]] to i64 |
| // CHECK_CONSTRAINED-NEXT: [[TMP7:%.*]] = ptrtoint ptr [[TMP2]] to i64 |
| // CHECK_CONSTRAINED-NEXT: [[TMP8:%.*]] = sub i64 [[TMP7]], [[TMP6]] |
| // CHECK_CONSTRAINED-NEXT: [[TMP9:%.*]] = add i64 [[TMP8]], 1 |
| // CHECK_CONSTRAINED-NEXT: br label [[STRLEN_JOIN]] |
| // CHECK_CONSTRAINED: strlen.join: |
| // CHECK_CONSTRAINED-NEXT: [[TMP10:%.*]] = phi i64 [ [[TMP9]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ] |
| // CHECK_CONSTRAINED-NEXT: [[TMP11:%.*]] = add i64 [[TMP10]], 7 |
| // CHECK_CONSTRAINED-NEXT: [[TMP12:%.*]] = and i64 [[TMP11]], 4294967288 |
| // CHECK_CONSTRAINED-NEXT: [[TMP13:%.*]] = add i64 [[TMP12]], 12 |
| // CHECK_CONSTRAINED-NEXT: [[TMP14:%.*]] = trunc i64 [[TMP13]] to i32 |
| // CHECK_CONSTRAINED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP14]]) |
| // CHECK_CONSTRAINED-NEXT: [[TMP15:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null |
| // CHECK_CONSTRAINED-NEXT: br i1 [[TMP15]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]] |
| // CHECK_CONSTRAINED: end.block: |
| // CHECK_CONSTRAINED-NEXT: [[TMP16:%.*]] = xor i1 [[TMP15]], true |
| // CHECK_CONSTRAINED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP16]] to i32 |
| // CHECK_CONSTRAINED-NEXT: ret i32 [[PRINTF_RESULT]] |
| // CHECK_CONSTRAINED: argpush.block: |
| // CHECK_CONSTRAINED-NEXT: [[TMP17:%.*]] = shl i32 [[TMP14]], 2 |
| // CHECK_CONSTRAINED-NEXT: store i32 [[TMP17]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4 |
| // CHECK_CONSTRAINED-NEXT: [[TMP18:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4 |
| // CHECK_CONSTRAINED-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[TMP18]], ptr align 1 [[TMP0]], i64 [[TMP10]], i1 false) |
| // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP18]], i64 [[TMP12]] |
| // CHECK_CONSTRAINED-NEXT: store i64 10, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8 |
| // CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8 |
| // CHECK_CONSTRAINED-NEXT: br label [[END_BLOCK]] |
| // |
| __device__ int foo4() { |
| const char* s = "format str%d"; |
| return printf(s, 10); |
| } |
