479 lines · cpp
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 22// REQUIRES: amdgpu-registered-target3// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -emit-llvm -mprintf-kind=buffered -fcuda-is-device \4// RUN: -o - %s | FileCheck --enable-var-scope %s5// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -emit-llvm -ffp-exception-behavior=strict -mprintf-kind=buffered -fcuda-is-device \6// RUN: -o - %s | FileCheck --enable-var-scope --check-prefix=CHECK_CONSTRAINED %s7 8#define __device__ __attribute__((device))9#define __shared__ __attribute__((shared))10#define __constant__ __attribute__((constant))11 12extern "C" __device__ int printf(const char *format, ...);13 14// CHECK-LABEL: define dso_local noundef i32 @_Z4foo1v15// CHECK-NEXT: entry:16// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)17// CHECK-NEXT: [[S:%.*]] = alloca ptr, align 8, addrspace(5)18// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr19// CHECK-NEXT: [[S_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S]] to ptr20// CHECK-NEXT: store ptr addrspacecast (ptr addrspace(4) @.str to ptr), ptr [[S_ASCAST]], align 821// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[S_ASCAST]], align 822// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[S_ASCAST]], align 823// CHECK-NEXT: [[TMP2:%.*]] = icmp eq ptr [[TMP0]], null24// CHECK-NEXT: br i1 [[TMP2]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]25// CHECK: strlen.while:26// CHECK-NEXT: [[TMP3:%.*]] = phi ptr [ [[TMP0]], [[ENTRY:%.*]] ], [ [[TMP4:%.*]], [[STRLEN_WHILE]] ]27// CHECK-NEXT: [[TMP4]] = getelementptr i8, ptr [[TMP3]], i64 128// CHECK-NEXT: [[TMP5:%.*]] = load i8, ptr [[TMP3]], align 129// CHECK-NEXT: [[TMP6:%.*]] = icmp eq i8 [[TMP5]], 030// CHECK-NEXT: br i1 [[TMP6]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]31// CHECK: strlen.while.done:32// CHECK-NEXT: [[TMP7:%.*]] = ptrtoint ptr [[TMP0]] to i6433// CHECK-NEXT: [[TMP8:%.*]] = ptrtoint ptr [[TMP3]] to i6434// CHECK-NEXT: [[TMP9:%.*]] = sub i64 [[TMP8]], [[TMP7]]35// CHECK-NEXT: [[TMP10:%.*]] = add i64 [[TMP9]], 136// CHECK-NEXT: br label [[STRLEN_JOIN]]37// CHECK: strlen.join:38// CHECK-NEXT: [[TMP11:%.*]] = phi i64 [ [[TMP10]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]39// CHECK-NEXT: [[TMP12:%.*]] = add i64 [[TMP11]], 740// CHECK-NEXT: [[TMP13:%.*]] = and i64 [[TMP12]], 429496728841// CHECK-NEXT: [[TMP14:%.*]] = add i64 [[TMP13]], 5242// CHECK-NEXT: [[TMP15:%.*]] = trunc i64 [[TMP14]] to i3243// CHECK-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP15]])44// CHECK-NEXT: [[TMP16:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null45// CHECK-NEXT: br i1 [[TMP16]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]46// CHECK: end.block:47// CHECK-NEXT: [[TMP17:%.*]] = xor i1 [[TMP16]], true48// CHECK-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP17]] to i3249// CHECK-NEXT: ret i32 [[PRINTF_RESULT]]50// CHECK: argpush.block:51// CHECK-NEXT: [[TMP18:%.*]] = shl i32 [[TMP15]], 252// CHECK-NEXT: [[TMP19:%.*]] = or i32 [[TMP18]], 253// CHECK-NEXT: store i32 [[TMP19]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 454// CHECK-NEXT: [[TMP20:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 455// CHECK-NEXT: store i64 1107004088646384690, ptr addrspace(1) [[TMP20]], align 856// CHECK-NEXT: [[TMP21:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP20]], i32 857// CHECK-NEXT: store i64 8, ptr addrspace(1) [[TMP21]], align 858// CHECK-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP21]], i32 859// CHECK-NEXT: store double 3.141590e+00, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 860// CHECK-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 861// CHECK-NEXT: store i64 8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], align 862// CHECK-NEXT: [[PRINTBUFFNEXTPTR2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], i32 863// CHECK-NEXT: store i64 4, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], align 864// CHECK-NEXT: [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], i32 865// CHECK-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[PRINTBUFFNEXTPTR3]], ptr align 1 [[TMP0]], i64 [[TMP11]], i1 false)66// CHECK-NEXT: [[PRINTBUFFNEXTPTR4:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], i64 [[TMP13]]67// CHECK-NEXT: store ptr [[TMP1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], align 868// CHECK-NEXT: [[PRINTBUFFNEXTPTR5:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], i32 869// CHECK-NEXT: br label [[END_BLOCK]]70//71// CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z4foo1v72// CHECK_CONSTRAINED-NEXT: entry:73// CHECK_CONSTRAINED-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)74// CHECK_CONSTRAINED-NEXT: [[S:%.*]] = alloca ptr, align 8, addrspace(5)75// CHECK_CONSTRAINED-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr76// CHECK_CONSTRAINED-NEXT: [[S_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S]] to ptr77// CHECK_CONSTRAINED-NEXT: store ptr addrspacecast (ptr addrspace(4) @.str to ptr), ptr [[S_ASCAST]], align 878// CHECK_CONSTRAINED-NEXT: [[TMP0:%.*]] = load ptr, ptr [[S_ASCAST]], align 879// CHECK_CONSTRAINED-NEXT: [[TMP1:%.*]] = load ptr, ptr [[S_ASCAST]], align 880// CHECK_CONSTRAINED-NEXT: [[TMP2:%.*]] = icmp eq ptr [[TMP0]], null81// CHECK_CONSTRAINED-NEXT: br i1 [[TMP2]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]82// CHECK_CONSTRAINED: strlen.while:83// CHECK_CONSTRAINED-NEXT: [[TMP3:%.*]] = phi ptr [ [[TMP0]], [[ENTRY:%.*]] ], [ [[TMP4:%.*]], [[STRLEN_WHILE]] ]84// CHECK_CONSTRAINED-NEXT: [[TMP4]] = getelementptr i8, ptr [[TMP3]], i64 185// CHECK_CONSTRAINED-NEXT: [[TMP5:%.*]] = load i8, ptr [[TMP3]], align 186// CHECK_CONSTRAINED-NEXT: [[TMP6:%.*]] = icmp eq i8 [[TMP5]], 087// CHECK_CONSTRAINED-NEXT: br i1 [[TMP6]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]88// CHECK_CONSTRAINED: strlen.while.done:89// CHECK_CONSTRAINED-NEXT: [[TMP7:%.*]] = ptrtoint ptr [[TMP0]] to i6490// CHECK_CONSTRAINED-NEXT: [[TMP8:%.*]] = ptrtoint ptr [[TMP3]] to i6491// CHECK_CONSTRAINED-NEXT: [[TMP9:%.*]] = sub i64 [[TMP8]], [[TMP7]]92// CHECK_CONSTRAINED-NEXT: [[TMP10:%.*]] = add i64 [[TMP9]], 193// CHECK_CONSTRAINED-NEXT: br label [[STRLEN_JOIN]]94// CHECK_CONSTRAINED: strlen.join:95// CHECK_CONSTRAINED-NEXT: [[TMP11:%.*]] = phi i64 [ [[TMP10]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]96// CHECK_CONSTRAINED-NEXT: [[TMP12:%.*]] = add i64 [[TMP11]], 797// CHECK_CONSTRAINED-NEXT: [[TMP13:%.*]] = and i64 [[TMP12]], 429496728898// CHECK_CONSTRAINED-NEXT: [[TMP14:%.*]] = add i64 [[TMP13]], 5299// CHECK_CONSTRAINED-NEXT: [[TMP15:%.*]] = trunc i64 [[TMP14]] to i32100// CHECK_CONSTRAINED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP15]])101// CHECK_CONSTRAINED-NEXT: [[TMP16:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null102// CHECK_CONSTRAINED-NEXT: br i1 [[TMP16]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]103// CHECK_CONSTRAINED: end.block:104// CHECK_CONSTRAINED-NEXT: [[TMP17:%.*]] = xor i1 [[TMP16]], true105// CHECK_CONSTRAINED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP17]] to i32106// CHECK_CONSTRAINED-NEXT: ret i32 [[PRINTF_RESULT]]107// CHECK_CONSTRAINED: argpush.block:108// CHECK_CONSTRAINED-NEXT: [[TMP18:%.*]] = shl i32 [[TMP15]], 2109// CHECK_CONSTRAINED-NEXT: [[TMP19:%.*]] = or i32 [[TMP18]], 2110// CHECK_CONSTRAINED-NEXT: store i32 [[TMP19]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4111// CHECK_CONSTRAINED-NEXT: [[TMP20:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4112// CHECK_CONSTRAINED-NEXT: store i64 1107004088646384690, ptr addrspace(1) [[TMP20]], align 8113// CHECK_CONSTRAINED-NEXT: [[TMP21:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP20]], i32 8114// CHECK_CONSTRAINED-NEXT: store i64 8, ptr addrspace(1) [[TMP21]], align 8115// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP21]], i32 8116// CHECK_CONSTRAINED-NEXT: store double 3.141590e+00, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8117// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8118// CHECK_CONSTRAINED-NEXT: store i64 8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], align 8119// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], i32 8120// CHECK_CONSTRAINED-NEXT: store i64 4, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], align 8121// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], i32 8122// CHECK_CONSTRAINED-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[PRINTBUFFNEXTPTR3]], ptr align 1 [[TMP0]], i64 [[TMP11]], i1 false)123// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR4:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], i64 [[TMP13]]124// CHECK_CONSTRAINED-NEXT: store ptr [[TMP1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], align 8125// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR5:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], i32 8126// CHECK_CONSTRAINED-NEXT: br label [[END_BLOCK]]127//128__device__ int foo1() {129 const char *s = "hello world";130 return printf("%.*f %*.*s %p\n", 8, 3.14159, 8, 4, s, s);131}132 133__device__ char *dstr;134__device__ const135// CHECK-LABEL: define dso_local noundef i32 @_Z4foo2v136// CHECK-NEXT: entry:137// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)138// CHECK-NEXT: [[LCVAL:%.*]] = alloca i32, align 4, addrspace(5)139// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr140// CHECK-NEXT: [[LCVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[LCVAL]] to ptr141// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspacecast (ptr addrspace(1) @dstr to ptr), align 8142// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr addrspacecast (ptr addrspace(1) @dstr to ptr), align 8143// CHECK-NEXT: [[TMP2:%.*]] = icmp eq ptr [[TMP0]], null144// CHECK-NEXT: br i1 [[TMP2]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]145// CHECK: strlen.while:146// CHECK-NEXT: [[TMP3:%.*]] = phi ptr [ [[TMP0]], [[ENTRY:%.*]] ], [ [[TMP4:%.*]], [[STRLEN_WHILE]] ]147// CHECK-NEXT: [[TMP4]] = getelementptr i8, ptr [[TMP3]], i64 1148// CHECK-NEXT: [[TMP5:%.*]] = load i8, ptr [[TMP3]], align 1149// CHECK-NEXT: [[TMP6:%.*]] = icmp eq i8 [[TMP5]], 0150// CHECK-NEXT: br i1 [[TMP6]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]151// CHECK: strlen.while.done:152// CHECK-NEXT: [[TMP7:%.*]] = ptrtoint ptr [[TMP0]] to i64153// CHECK-NEXT: [[TMP8:%.*]] = ptrtoint ptr [[TMP3]] to i64154// CHECK-NEXT: [[TMP9:%.*]] = sub i64 [[TMP8]], [[TMP7]]155// CHECK-NEXT: [[TMP10:%.*]] = add i64 [[TMP9]], 1156// CHECK-NEXT: br label [[STRLEN_JOIN]]157// CHECK: strlen.join:158// CHECK-NEXT: [[TMP11:%.*]] = phi i64 [ [[TMP10]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]159// CHECK-NEXT: [[TMP12:%.*]] = add i64 [[TMP11]], 7160// CHECK-NEXT: [[TMP13:%.*]] = and i64 [[TMP12]], 4294967288161// CHECK-NEXT: [[TMP14:%.*]] = add i64 [[TMP13]], 36162// CHECK-NEXT: [[TMP15:%.*]] = trunc i64 [[TMP14]] to i32163// CHECK-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP15]])164// CHECK-NEXT: [[TMP16:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null165// CHECK-NEXT: br i1 [[TMP16]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]166// CHECK: end.block:167// CHECK-NEXT: [[TMP17:%.*]] = xor i1 [[TMP16]], true168// CHECK-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP17]] to i32169// CHECK-NEXT: ret i32 [[PRINTF_RESULT]]170// CHECK: argpush.block:171// CHECK-NEXT: [[TMP18:%.*]] = shl i32 [[TMP15]], 2172// CHECK-NEXT: [[TMP19:%.*]] = or i32 [[TMP18]], 2173// CHECK-NEXT: store i32 [[TMP19]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4174// CHECK-NEXT: [[TMP20:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4175// CHECK-NEXT: store i64 7257695813269076350, ptr addrspace(1) [[TMP20]], align 8176// CHECK-NEXT: [[TMP21:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP20]], i32 8177// CHECK-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[TMP21]], ptr align 1 [[TMP0]], i64 [[TMP11]], i1 false)178// CHECK-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP21]], i64 [[TMP13]]179// CHECK-NEXT: store ptr [[TMP1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8180// CHECK-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8181// CHECK-NEXT: store ptr addrspacecast (ptr addrspace(3) @_ZZ4foo2vE5shval to ptr), ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], align 8182// CHECK-NEXT: [[PRINTBUFFNEXTPTR2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], i32 8183// CHECK-NEXT: store ptr [[LCVAL_ASCAST]], ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], align 8184// CHECK-NEXT: [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], i32 8185// CHECK-NEXT: br label [[END_BLOCK]]186//187// CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z4foo2v188// CHECK_CONSTRAINED-NEXT: entry:189// CHECK_CONSTRAINED-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)190// CHECK_CONSTRAINED-NEXT: [[LCVAL:%.*]] = alloca i32, align 4, addrspace(5)191// CHECK_CONSTRAINED-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr192// CHECK_CONSTRAINED-NEXT: [[LCVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[LCVAL]] to ptr193// CHECK_CONSTRAINED-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspacecast (ptr addrspace(1) @dstr to ptr), align 8194// CHECK_CONSTRAINED-NEXT: [[TMP1:%.*]] = load ptr, ptr addrspacecast (ptr addrspace(1) @dstr to ptr), align 8195// CHECK_CONSTRAINED-NEXT: [[TMP2:%.*]] = icmp eq ptr [[TMP0]], null196// CHECK_CONSTRAINED-NEXT: br i1 [[TMP2]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]197// CHECK_CONSTRAINED: strlen.while:198// CHECK_CONSTRAINED-NEXT: [[TMP3:%.*]] = phi ptr [ [[TMP0]], [[ENTRY:%.*]] ], [ [[TMP4:%.*]], [[STRLEN_WHILE]] ]199// CHECK_CONSTRAINED-NEXT: [[TMP4]] = getelementptr i8, ptr [[TMP3]], i64 1200// CHECK_CONSTRAINED-NEXT: [[TMP5:%.*]] = load i8, ptr [[TMP3]], align 1201// CHECK_CONSTRAINED-NEXT: [[TMP6:%.*]] = icmp eq i8 [[TMP5]], 0202// CHECK_CONSTRAINED-NEXT: br i1 [[TMP6]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]203// CHECK_CONSTRAINED: strlen.while.done:204// CHECK_CONSTRAINED-NEXT: [[TMP7:%.*]] = ptrtoint ptr [[TMP0]] to i64205// CHECK_CONSTRAINED-NEXT: [[TMP8:%.*]] = ptrtoint ptr [[TMP3]] to i64206// CHECK_CONSTRAINED-NEXT: [[TMP9:%.*]] = sub i64 [[TMP8]], [[TMP7]]207// CHECK_CONSTRAINED-NEXT: [[TMP10:%.*]] = add i64 [[TMP9]], 1208// CHECK_CONSTRAINED-NEXT: br label [[STRLEN_JOIN]]209// CHECK_CONSTRAINED: strlen.join:210// CHECK_CONSTRAINED-NEXT: [[TMP11:%.*]] = phi i64 [ [[TMP10]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]211// CHECK_CONSTRAINED-NEXT: [[TMP12:%.*]] = add i64 [[TMP11]], 7212// CHECK_CONSTRAINED-NEXT: [[TMP13:%.*]] = and i64 [[TMP12]], 4294967288213// CHECK_CONSTRAINED-NEXT: [[TMP14:%.*]] = add i64 [[TMP13]], 36214// CHECK_CONSTRAINED-NEXT: [[TMP15:%.*]] = trunc i64 [[TMP14]] to i32215// CHECK_CONSTRAINED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP15]])216// CHECK_CONSTRAINED-NEXT: [[TMP16:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null217// CHECK_CONSTRAINED-NEXT: br i1 [[TMP16]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]218// CHECK_CONSTRAINED: end.block:219// CHECK_CONSTRAINED-NEXT: [[TMP17:%.*]] = xor i1 [[TMP16]], true220// CHECK_CONSTRAINED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP17]] to i32221// CHECK_CONSTRAINED-NEXT: ret i32 [[PRINTF_RESULT]]222// CHECK_CONSTRAINED: argpush.block:223// CHECK_CONSTRAINED-NEXT: [[TMP18:%.*]] = shl i32 [[TMP15]], 2224// CHECK_CONSTRAINED-NEXT: [[TMP19:%.*]] = or i32 [[TMP18]], 2225// CHECK_CONSTRAINED-NEXT: store i32 [[TMP19]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4226// CHECK_CONSTRAINED-NEXT: [[TMP20:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4227// CHECK_CONSTRAINED-NEXT: store i64 7257695813269076350, ptr addrspace(1) [[TMP20]], align 8228// CHECK_CONSTRAINED-NEXT: [[TMP21:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP20]], i32 8229// CHECK_CONSTRAINED-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[TMP21]], ptr align 1 [[TMP0]], i64 [[TMP11]], i1 false)230// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP21]], i64 [[TMP13]]231// CHECK_CONSTRAINED-NEXT: store ptr [[TMP1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8232// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8233// CHECK_CONSTRAINED-NEXT: store ptr addrspacecast (ptr addrspace(3) @_ZZ4foo2vE5shval to ptr), ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], align 8234// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], i32 8235// CHECK_CONSTRAINED-NEXT: store ptr [[LCVAL_ASCAST]], ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], align 8236// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], i32 8237// CHECK_CONSTRAINED-NEXT: br label [[END_BLOCK]]238//239__device__ int foo2() {240 __shared__ int shval;241 int lcval;242 return printf("%s %p %p %p\n", dstr, dstr, &shval, &lcval);243}244 245__device__ unsigned short g = 30;246__device__ unsigned long n = 30;247 248__device__ float f1 = 3.14f;249__device__ double f2 = 2.71828;250__device__ _Float16 f3 = 2.71;251__device__ __bf16 f4 = 3.142;252__device__ _BitInt(55) Int55 = 31;253__device__ _BitInt(44) Int44 = 312;254__device__ _BitInt(128) Int128 = 45637;255 256// CHECK-LABEL: define dso_local noundef i32 @_Z4foo3v257// CHECK-NEXT: entry:258// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)259// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr260// CHECK-NEXT: store i32 25, ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), align 4261// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), align 4262// CHECK-NEXT: [[TMP1:%.*]] = load i16, ptr addrspacecast (ptr addrspace(1) @g to ptr), align 2263// CHECK-NEXT: [[CONV:%.*]] = zext i16 [[TMP1]] to i32264// CHECK-NEXT: [[TMP2:%.*]] = load i64, ptr addrspacecast (ptr addrspace(1) @n to ptr), align 8265// CHECK-NEXT: [[TMP3:%.*]] = load float, ptr addrspacecast (ptr addrspace(1) @f1 to ptr), align 4266// CHECK-NEXT: [[CONV1:%.*]] = fpext contract float [[TMP3]] to double267// CHECK-NEXT: [[TMP4:%.*]] = load double, ptr addrspacecast (ptr addrspace(1) @f2 to ptr), align 8268// CHECK-NEXT: [[TMP5:%.*]] = load half, ptr addrspacecast (ptr addrspace(1) @f3 to ptr), align 2269// CHECK-NEXT: [[TMP6:%.*]] = load bfloat, ptr addrspacecast (ptr addrspace(1) @f4 to ptr), align 2270// CHECK-NEXT: [[TMP7:%.*]] = load i64, ptr addrspacecast (ptr addrspace(1) @Int55 to ptr), align 8271// CHECK-NEXT: [[LOADEDV:%.*]] = trunc i64 [[TMP7]] to i55272// CHECK-NEXT: [[TMP8:%.*]] = load i64, ptr addrspacecast (ptr addrspace(1) @Int44 to ptr), align 8273// CHECK-NEXT: [[LOADEDV2:%.*]] = trunc i64 [[TMP8]] to i44274// CHECK-NEXT: [[TMP9:%.*]] = load i128, ptr addrspacecast (ptr addrspace(1) @Int128 to ptr), align 8275// CHECK-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 108)276// CHECK-NEXT: [[TMP10:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null277// CHECK-NEXT: br i1 [[TMP10]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]278// CHECK: end.block:279// CHECK-NEXT: [[TMP11:%.*]] = xor i1 [[TMP10]], true280// CHECK-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP11]] to i32281// CHECK-NEXT: ret i32 [[PRINTF_RESULT]]282// CHECK: argpush.block:283// CHECK-NEXT: store i32 434, ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4284// CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4285// CHECK-NEXT: store i64 7271852820361268873, ptr addrspace(1) [[TMP12]], align 8286// CHECK-NEXT: [[TMP13:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP12]], i32 8287// CHECK-NEXT: [[TMP14:%.*]] = zext i32 [[TMP0]] to i64288// CHECK-NEXT: store i64 [[TMP14]], ptr addrspace(1) [[TMP13]], align 8289// CHECK-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP13]], i32 8290// CHECK-NEXT: store ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8291// CHECK-NEXT: [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8292// CHECK-NEXT: [[TMP15:%.*]] = zext i32 [[CONV]] to i64293// CHECK-NEXT: store i64 [[TMP15]], ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], align 8294// CHECK-NEXT: [[PRINTBUFFNEXTPTR4:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], i32 8295// CHECK-NEXT: store i64 [[TMP2]], ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], align 8296// CHECK-NEXT: [[PRINTBUFFNEXTPTR5:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], i32 8297// CHECK-NEXT: store double [[CONV1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR5]], align 8298// CHECK-NEXT: [[PRINTBUFFNEXTPTR6:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR5]], i32 8299// CHECK-NEXT: store double [[TMP4]], ptr addrspace(1) [[PRINTBUFFNEXTPTR6]], align 8300// CHECK-NEXT: [[PRINTBUFFNEXTPTR7:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR6]], i32 8301// CHECK-NEXT: [[TMP16:%.*]] = fpext half [[TMP5]] to double302// CHECK-NEXT: store double [[TMP16]], ptr addrspace(1) [[PRINTBUFFNEXTPTR7]], align 8303// CHECK-NEXT: [[PRINTBUFFNEXTPTR8:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR7]], i32 8304// CHECK-NEXT: [[TMP17:%.*]] = fpext bfloat [[TMP6]] to double305// CHECK-NEXT: store double [[TMP17]], ptr addrspace(1) [[PRINTBUFFNEXTPTR8]], align 8306// CHECK-NEXT: [[PRINTBUFFNEXTPTR9:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR8]], i32 8307// CHECK-NEXT: [[TMP18:%.*]] = zext i55 [[LOADEDV]] to i64308// CHECK-NEXT: store i64 [[TMP18]], ptr addrspace(1) [[PRINTBUFFNEXTPTR9]], align 8309// CHECK-NEXT: [[PRINTBUFFNEXTPTR10:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR9]], i32 8310// CHECK-NEXT: [[TMP19:%.*]] = zext i44 [[LOADEDV2]] to i64311// CHECK-NEXT: store i64 [[TMP19]], ptr addrspace(1) [[PRINTBUFFNEXTPTR10]], align 8312// CHECK-NEXT: [[PRINTBUFFNEXTPTR11:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR10]], i32 8313// CHECK-NEXT: store i128 [[TMP9]], ptr addrspace(1) [[PRINTBUFFNEXTPTR11]], align 8314// CHECK-NEXT: [[PRINTBUFFNEXTPTR12:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR11]], i32 16315// CHECK-NEXT: br label [[END_BLOCK]]316//317// CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z4foo3v318// CHECK_CONSTRAINED-NEXT: entry:319// CHECK_CONSTRAINED-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)320// CHECK_CONSTRAINED-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr321// CHECK_CONSTRAINED-NEXT: store i32 25, ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), align 4322// CHECK_CONSTRAINED-NEXT: [[TMP0:%.*]] = load i32, ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), align 4323// CHECK_CONSTRAINED-NEXT: [[TMP1:%.*]] = load i16, ptr addrspacecast (ptr addrspace(1) @g to ptr), align 2324// CHECK_CONSTRAINED-NEXT: [[CONV:%.*]] = zext i16 [[TMP1]] to i32325// CHECK_CONSTRAINED-NEXT: [[TMP2:%.*]] = load i64, ptr addrspacecast (ptr addrspace(1) @n to ptr), align 8326// CHECK_CONSTRAINED-NEXT: [[TMP3:%.*]] = load float, ptr addrspacecast (ptr addrspace(1) @f1 to ptr), align 4327// CHECK_CONSTRAINED-NEXT: [[CONV1:%.*]] = fpext contract float [[TMP3]] to double328// CHECK_CONSTRAINED-NEXT: [[TMP4:%.*]] = load double, ptr addrspacecast (ptr addrspace(1) @f2 to ptr), align 8329// CHECK_CONSTRAINED-NEXT: [[TMP5:%.*]] = load half, ptr addrspacecast (ptr addrspace(1) @f3 to ptr), align 2330// CHECK_CONSTRAINED-NEXT: [[TMP6:%.*]] = load bfloat, ptr addrspacecast (ptr addrspace(1) @f4 to ptr), align 2331// CHECK_CONSTRAINED-NEXT: [[TMP7:%.*]] = load i64, ptr addrspacecast (ptr addrspace(1) @Int55 to ptr), align 8332// CHECK_CONSTRAINED-NEXT: [[LOADEDV:%.*]] = trunc i64 [[TMP7]] to i55333// CHECK_CONSTRAINED-NEXT: [[TMP8:%.*]] = load i64, ptr addrspacecast (ptr addrspace(1) @Int44 to ptr), align 8334// CHECK_CONSTRAINED-NEXT: [[LOADEDV2:%.*]] = trunc i64 [[TMP8]] to i44335// CHECK_CONSTRAINED-NEXT: [[TMP9:%.*]] = load i128, ptr addrspacecast (ptr addrspace(1) @Int128 to ptr), align 8336// CHECK_CONSTRAINED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 108)337// CHECK_CONSTRAINED-NEXT: [[TMP10:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null338// CHECK_CONSTRAINED-NEXT: br i1 [[TMP10]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]339// CHECK_CONSTRAINED: end.block:340// CHECK_CONSTRAINED-NEXT: [[TMP11:%.*]] = xor i1 [[TMP10]], true341// CHECK_CONSTRAINED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP11]] to i32342// CHECK_CONSTRAINED-NEXT: ret i32 [[PRINTF_RESULT]]343// CHECK_CONSTRAINED: argpush.block:344// CHECK_CONSTRAINED-NEXT: store i32 434, ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4345// CHECK_CONSTRAINED-NEXT: [[TMP12:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4346// CHECK_CONSTRAINED-NEXT: store i64 7271852820361268873, ptr addrspace(1) [[TMP12]], align 8347// CHECK_CONSTRAINED-NEXT: [[TMP13:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP12]], i32 8348// CHECK_CONSTRAINED-NEXT: [[TMP14:%.*]] = zext i32 [[TMP0]] to i64349// CHECK_CONSTRAINED-NEXT: store i64 [[TMP14]], ptr addrspace(1) [[TMP13]], align 8350// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP13]], i32 8351// CHECK_CONSTRAINED-NEXT: store ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8352// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8353// CHECK_CONSTRAINED-NEXT: [[TMP15:%.*]] = zext i32 [[CONV]] to i64354// CHECK_CONSTRAINED-NEXT: store i64 [[TMP15]], ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], align 8355// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR4:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], i32 8356// CHECK_CONSTRAINED-NEXT: store i64 [[TMP2]], ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], align 8357// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR5:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], i32 8358// CHECK_CONSTRAINED-NEXT: store double [[CONV1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR5]], align 8359// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR6:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR5]], i32 8360// CHECK_CONSTRAINED-NEXT: store double [[TMP4]], ptr addrspace(1) [[PRINTBUFFNEXTPTR6]], align 8361// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR7:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR6]], i32 8362// CHECK_CONSTRAINED-NEXT: [[TMP16:%.*]] = fpext half [[TMP5]] to double363// CHECK_CONSTRAINED-NEXT: store double [[TMP16]], ptr addrspace(1) [[PRINTBUFFNEXTPTR7]], align 8364// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR8:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR7]], i32 8365// CHECK_CONSTRAINED-NEXT: [[TMP17:%.*]] = fpext bfloat [[TMP6]] to double366// CHECK_CONSTRAINED-NEXT: store double [[TMP17]], ptr addrspace(1) [[PRINTBUFFNEXTPTR8]], align 8367// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR9:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR8]], i32 8368// CHECK_CONSTRAINED-NEXT: [[TMP18:%.*]] = zext i55 [[LOADEDV]] to i64369// CHECK_CONSTRAINED-NEXT: store i64 [[TMP18]], ptr addrspace(1) [[PRINTBUFFNEXTPTR9]], align 8370// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR10:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR9]], i32 8371// CHECK_CONSTRAINED-NEXT: [[TMP19:%.*]] = zext i44 [[LOADEDV2]] to i64372// CHECK_CONSTRAINED-NEXT: store i64 [[TMP19]], ptr addrspace(1) [[PRINTBUFFNEXTPTR10]], align 8373// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR11:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR10]], i32 8374// CHECK_CONSTRAINED-NEXT: store i128 [[TMP9]], ptr addrspace(1) [[PRINTBUFFNEXTPTR11]], align 8375// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR12:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR11]], i32 16376// CHECK_CONSTRAINED-NEXT: br label [[END_BLOCK]]377//378__device__ int foo3() {379 __shared__ int s;380 s = 25;381 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);382}383 384//A non trivial case,385// CHECK-LABEL: define dso_local noundef i32 @_Z4foo4v386// CHECK-NEXT: entry:387// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)388// CHECK-NEXT: [[S:%.*]] = alloca ptr, align 8, addrspace(5)389// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr390// CHECK-NEXT: [[S_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S]] to ptr391// CHECK-NEXT: store ptr addrspacecast (ptr addrspace(4) @.str.4 to ptr), ptr [[S_ASCAST]], align 8392// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[S_ASCAST]], align 8393// CHECK-NEXT: [[TMP1:%.*]] = icmp eq ptr [[TMP0]], null394// CHECK-NEXT: br i1 [[TMP1]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]395// CHECK: strlen.while:396// CHECK-NEXT: [[TMP2:%.*]] = phi ptr [ [[TMP0]], [[ENTRY:%.*]] ], [ [[TMP3:%.*]], [[STRLEN_WHILE]] ]397// CHECK-NEXT: [[TMP3]] = getelementptr i8, ptr [[TMP2]], i64 1398// CHECK-NEXT: [[TMP4:%.*]] = load i8, ptr [[TMP2]], align 1399// CHECK-NEXT: [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0400// CHECK-NEXT: br i1 [[TMP5]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]401// CHECK: strlen.while.done:402// CHECK-NEXT: [[TMP6:%.*]] = ptrtoint ptr [[TMP0]] to i64403// CHECK-NEXT: [[TMP7:%.*]] = ptrtoint ptr [[TMP2]] to i64404// CHECK-NEXT: [[TMP8:%.*]] = sub i64 [[TMP7]], [[TMP6]]405// CHECK-NEXT: [[TMP9:%.*]] = add i64 [[TMP8]], 1406// CHECK-NEXT: br label [[STRLEN_JOIN]]407// CHECK: strlen.join:408// CHECK-NEXT: [[TMP10:%.*]] = phi i64 [ [[TMP9]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]409// CHECK-NEXT: [[TMP11:%.*]] = add i64 [[TMP10]], 7410// CHECK-NEXT: [[TMP12:%.*]] = and i64 [[TMP11]], 4294967288411// CHECK-NEXT: [[TMP13:%.*]] = add i64 [[TMP12]], 12412// CHECK-NEXT: [[TMP14:%.*]] = trunc i64 [[TMP13]] to i32413// CHECK-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP14]])414// CHECK-NEXT: [[TMP15:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null415// CHECK-NEXT: br i1 [[TMP15]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]416// CHECK: end.block:417// CHECK-NEXT: [[TMP16:%.*]] = xor i1 [[TMP15]], true418// CHECK-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP16]] to i32419// CHECK-NEXT: ret i32 [[PRINTF_RESULT]]420// CHECK: argpush.block:421// CHECK-NEXT: [[TMP17:%.*]] = shl i32 [[TMP14]], 2422// CHECK-NEXT: store i32 [[TMP17]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4423// CHECK-NEXT: [[TMP18:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4424// CHECK-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[TMP18]], ptr align 1 [[TMP0]], i64 [[TMP10]], i1 false)425// CHECK-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP18]], i64 [[TMP12]]426// CHECK-NEXT: store i64 10, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8427// CHECK-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8428// CHECK-NEXT: br label [[END_BLOCK]]429//430// CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z4foo4v431// CHECK_CONSTRAINED-NEXT: entry:432// CHECK_CONSTRAINED-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)433// CHECK_CONSTRAINED-NEXT: [[S:%.*]] = alloca ptr, align 8, addrspace(5)434// CHECK_CONSTRAINED-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr435// CHECK_CONSTRAINED-NEXT: [[S_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S]] to ptr436// CHECK_CONSTRAINED-NEXT: store ptr addrspacecast (ptr addrspace(4) @.str.4 to ptr), ptr [[S_ASCAST]], align 8437// CHECK_CONSTRAINED-NEXT: [[TMP0:%.*]] = load ptr, ptr [[S_ASCAST]], align 8438// CHECK_CONSTRAINED-NEXT: [[TMP1:%.*]] = icmp eq ptr [[TMP0]], null439// CHECK_CONSTRAINED-NEXT: br i1 [[TMP1]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]440// CHECK_CONSTRAINED: strlen.while:441// CHECK_CONSTRAINED-NEXT: [[TMP2:%.*]] = phi ptr [ [[TMP0]], [[ENTRY:%.*]] ], [ [[TMP3:%.*]], [[STRLEN_WHILE]] ]442// CHECK_CONSTRAINED-NEXT: [[TMP3]] = getelementptr i8, ptr [[TMP2]], i64 1443// CHECK_CONSTRAINED-NEXT: [[TMP4:%.*]] = load i8, ptr [[TMP2]], align 1444// CHECK_CONSTRAINED-NEXT: [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0445// CHECK_CONSTRAINED-NEXT: br i1 [[TMP5]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]446// CHECK_CONSTRAINED: strlen.while.done:447// CHECK_CONSTRAINED-NEXT: [[TMP6:%.*]] = ptrtoint ptr [[TMP0]] to i64448// CHECK_CONSTRAINED-NEXT: [[TMP7:%.*]] = ptrtoint ptr [[TMP2]] to i64449// CHECK_CONSTRAINED-NEXT: [[TMP8:%.*]] = sub i64 [[TMP7]], [[TMP6]]450// CHECK_CONSTRAINED-NEXT: [[TMP9:%.*]] = add i64 [[TMP8]], 1451// CHECK_CONSTRAINED-NEXT: br label [[STRLEN_JOIN]]452// CHECK_CONSTRAINED: strlen.join:453// CHECK_CONSTRAINED-NEXT: [[TMP10:%.*]] = phi i64 [ [[TMP9]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]454// CHECK_CONSTRAINED-NEXT: [[TMP11:%.*]] = add i64 [[TMP10]], 7455// CHECK_CONSTRAINED-NEXT: [[TMP12:%.*]] = and i64 [[TMP11]], 4294967288456// CHECK_CONSTRAINED-NEXT: [[TMP13:%.*]] = add i64 [[TMP12]], 12457// CHECK_CONSTRAINED-NEXT: [[TMP14:%.*]] = trunc i64 [[TMP13]] to i32458// CHECK_CONSTRAINED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP14]])459// CHECK_CONSTRAINED-NEXT: [[TMP15:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null460// CHECK_CONSTRAINED-NEXT: br i1 [[TMP15]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]461// CHECK_CONSTRAINED: end.block:462// CHECK_CONSTRAINED-NEXT: [[TMP16:%.*]] = xor i1 [[TMP15]], true463// CHECK_CONSTRAINED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP16]] to i32464// CHECK_CONSTRAINED-NEXT: ret i32 [[PRINTF_RESULT]]465// CHECK_CONSTRAINED: argpush.block:466// CHECK_CONSTRAINED-NEXT: [[TMP17:%.*]] = shl i32 [[TMP14]], 2467// CHECK_CONSTRAINED-NEXT: store i32 [[TMP17]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4468// CHECK_CONSTRAINED-NEXT: [[TMP18:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4469// CHECK_CONSTRAINED-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[TMP18]], ptr align 1 [[TMP0]], i64 [[TMP10]], i1 false)470// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP18]], i64 [[TMP12]]471// CHECK_CONSTRAINED-NEXT: store i64 10, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8472// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8473// CHECK_CONSTRAINED-NEXT: br label [[END_BLOCK]]474//475__device__ int foo4() {476 const char* s = "format str%d";477 return printf(s, 10);478}479