[FuncAttrs] Deduce noundef attributes for return values (#76553)

This patch deduces `noundef` attributes for return values. IIUC, a function returns `noundef` values iff all of its return values are guaranteed not to be `undef` or `poison`. Definition of `noundef` from LangRef: ``` noundef This attribute applies to parameters and return values. If the value representation contains any undefined or poison bits, the behavior is undefined. Note that this does not refer to padding introduced by the type’s storage representation. ``` Alive2: https://alive2.llvm.org/ce/z/g8Eis6 Compile-time impact: http://llvm-compile-time-tracker.com/compare.php?from=30dcc33c4ea3ab50397a7adbe85fe977d4a400bd&to=c5e8738d4bfbf1e97e3f455fded90b791f223d74&stat=instructions:u |stage1-O3|stage1-ReleaseThinLTO|stage1-ReleaseLTO-g|stage1-O0-g|stage2-O3|stage2-O0-g|stage2-clang| |--|--|--|--|--|--|--| |+0.01%|+0.01%|-0.01%|+0.01%|+0.03%|-0.04%|+0.01%| The motivation of this patch is to reduce the number of `freeze` insts and enable more optimizations.
llvm · Dec 31, 2023 · 1228bec · 1228bec
1 parent c7c912c
commit 1228bec
Show file tree

Hide file tree

Showing 47 changed files with 317 additions and 128 deletions.
diff --git a/clang/test/CodeGen/X86/ms-x86-intrinsics.c b/clang/test/CodeGen/X86/ms-x86-intrinsics.c
@@ -145,7 +145,7 @@ unsigned __int64 test__shiftleft128(unsigned __int64 l, unsigned __int64 h,
                                     unsigned char d) {
   return __shiftleft128(l, h, d);
 }
-// CHECK-X64-LABEL: define dso_local i64 @test__shiftleft128(i64 noundef %l, i64 noundef %h, i8 noundef %d)
+// CHECK-X64-LABEL: define dso_local noundef i64 @test__shiftleft128(i64 noundef %l, i64 noundef %h, i8 noundef %d)
 // CHECK-X64: = zext i8 %{{.*}} to i64
 // CHECK-X64: = tail call i64 @llvm.fshl.i64(i64 %h, i64 %l, i64 %{{.*}})
 // CHECK-X64:  ret i64 %
@@ -154,7 +154,7 @@ unsigned __int64 test__shiftright128(unsigned __int64 l, unsigned __int64 h,
                                      unsigned char d) {
   return __shiftright128(l, h, d);
 }
-// CHECK-X64-LABEL: define dso_local i64 @test__shiftright128(i64 noundef %l, i64 noundef %h, i8 noundef %d)
+// CHECK-X64-LABEL: define dso_local noundef i64 @test__shiftright128(i64 noundef %l, i64 noundef %h, i8 noundef %d)
 // CHECK-X64: = zext i8 %{{.*}} to i64
 // CHECK-X64: = tail call i64 @llvm.fshr.i64(i64 %h, i64 %l, i64 %{{.*}})
 // CHECK-X64:  ret i64 %

diff --git a/clang/test/CodeGen/arm-bf16-params-returns.c b/clang/test/CodeGen/arm-bf16-params-returns.c
@@ -11,7 +11,7 @@
 __bf16 test_ret_bf16(__bf16 v) {
   return v;
 }
-// CHECK32-HARD: define{{.*}} arm_aapcs_vfpcc bfloat @test_ret_bf16(bfloat noundef returned %v) {{.*}} {
+// CHECK32-HARD: define{{.*}} arm_aapcs_vfpcc noundef bfloat @test_ret_bf16(bfloat noundef returned %v) {{.*}} {
 // CHECK32-HARD: ret bfloat %v
 // CHECK32-SOFTFP: define{{.*}} bfloat @test_ret_bf16(bfloat noundef returned %v) {{.*}} {
 // CHECK32-SOFTFP: ret bfloat %v
@@ -23,11 +23,11 @@ __bf16 test_ret_bf16(__bf16 v) {
 bfloat16x4_t test_ret_bf16x4_t(bfloat16x4_t v) {
   return v;
 }
-// CHECK32-HARD: define{{.*}} arm_aapcs_vfpcc <4 x bfloat> @test_ret_bf16x4_t(<4 x bfloat> noundef returned %v) {{.*}} {
+// CHECK32-HARD: define{{.*}} arm_aapcs_vfpcc noundef <4 x bfloat> @test_ret_bf16x4_t(<4 x bfloat> noundef returned %v) {{.*}} {
 // CHECK32-HARD: ret <4 x bfloat> %v
-// CHECK32-SOFTFP: define{{.*}} <2 x i32> @test_ret_bf16x4_t(<2 x i32> [[V0:.*]]) {{.*}} {
+// CHECK32-SOFTFP: define{{.*}} noundef <2 x i32> @test_ret_bf16x4_t(<2 x i32> [[V0:.*]]) {{.*}} {
 // CHECK32-SOFTFP: ret <2 x i32> %v
-// CHECK64NEON: define{{.*}} <4 x bfloat> @test_ret_bf16x4_t(<4 x bfloat> noundef returned %v) {{.*}} {
+// CHECK64NEON: define{{.*}} noundef <4 x bfloat> @test_ret_bf16x4_t(<4 x bfloat> noundef returned %v) {{.*}} {
 // CHECK64NEON: ret <4 x bfloat> %v
 
 #endif
diff --git a/clang/test/CodeGen/arm-vector_type-params-returns.c b/clang/test/CodeGen/arm-vector_type-params-returns.c
@@ -27,7 +27,7 @@
 #endif
 
 // function return types
-// CHECK-LABEL: define dso_local <8 x half> @test_ret_v8f16(
+// CHECK-LABEL: define dso_local noundef <8 x half> @test_ret_v8f16(
 // CHECK-SAME: <8 x half> noundef returned [[V:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    ret <8 x half> [[V]]
@@ -36,7 +36,7 @@ float16x8_t test_ret_v8f16(float16x8_t v) {
   return v;
 }
 
-// CHECK-LABEL: define dso_local <4 x float> @test_ret_v4f32(
+// CHECK-LABEL: define dso_local noundef <4 x float> @test_ret_v4f32(
 // CHECK-SAME: <4 x float> noundef returned [[V:%.*]]) local_unnamed_addr #[[ATTR0]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    ret <4 x float> [[V]]
@@ -45,7 +45,7 @@ float32x4_t test_ret_v4f32(float32x4_t v) {
   return v;
 }
 
-// CHECK-LABEL: define dso_local <2 x double> @test_ret_v2f64(
+// CHECK-LABEL: define dso_local noundef <2 x double> @test_ret_v2f64(
 // CHECK-SAME: <2 x double> noundef returned [[V:%.*]]) local_unnamed_addr #[[ATTR0]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    ret <2 x double> [[V]]
@@ -54,7 +54,7 @@ float64x2_t test_ret_v2f64(float64x2_t v) {
   return v;
 }
 
-// CHECK-LABEL: define dso_local <8 x bfloat> @test_ret_v8bf16(
+// CHECK-LABEL: define dso_local noundef <8 x bfloat> @test_ret_v8bf16(
 // CHECK-SAME: <8 x bfloat> noundef returned [[V:%.*]]) local_unnamed_addr #[[ATTR0]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    ret <8 x bfloat> [[V]]
@@ -63,7 +63,7 @@ bfloat16x8_t test_ret_v8bf16(bfloat16x8_t v) {
   return v;
 }
 
-// CHECK-LABEL: define dso_local <16 x i8> @test_ret_v16s8(
+// CHECK-LABEL: define dso_local noundef <16 x i8> @test_ret_v16s8(
 // CHECK-SAME: <16 x i8> noundef returned [[V:%.*]]) local_unnamed_addr #[[ATTR0]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    ret <16 x i8> [[V]]
@@ -72,7 +72,7 @@ int8x16_t test_ret_v16s8(int8x16_t v) {
   return v;
 }
 
-// CHECK-LABEL: define dso_local <8 x i16> @test_ret_v8s16(
+// CHECK-LABEL: define dso_local noundef <8 x i16> @test_ret_v8s16(
 // CHECK-SAME: <8 x i16> noundef returned [[V:%.*]]) local_unnamed_addr #[[ATTR0]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    ret <8 x i16> [[V]]
@@ -81,7 +81,7 @@ int16x8_t test_ret_v8s16(int16x8_t v) {
   return v;
 }
 
-// CHECK-LABEL: define dso_local <4 x i32> @test_ret_v32s4(
+// CHECK-LABEL: define dso_local noundef <4 x i32> @test_ret_v32s4(
 // CHECK-SAME: <4 x i32> noundef returned [[V:%.*]]) local_unnamed_addr #[[ATTR0]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    ret <4 x i32> [[V]]
@@ -90,7 +90,7 @@ int32x4_t test_ret_v32s4(int32x4_t v) {
   return v;
 }
 
-// CHECK-LABEL: define dso_local <2 x i64> @test_ret_v64s2(
+// CHECK-LABEL: define dso_local noundef <2 x i64> @test_ret_v64s2(
 // CHECK-SAME: <2 x i64> noundef returned [[V:%.*]]) local_unnamed_addr #[[ATTR0]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    ret <2 x i64> [[V]]
@@ -99,7 +99,7 @@ int64x2_t test_ret_v64s2(int64x2_t v) {
   return v;
 }
 
-// CHECK-LABEL: define dso_local <16 x i8> @test_ret_v16u8(
+// CHECK-LABEL: define dso_local noundef <16 x i8> @test_ret_v16u8(
 // CHECK-SAME: <16 x i8> noundef returned [[V:%.*]]) local_unnamed_addr #[[ATTR0]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    ret <16 x i8> [[V]]
@@ -108,7 +108,7 @@ uint8x16_t test_ret_v16u8(uint8x16_t v) {
   return v;
 }
 
-// CHECK-LABEL: define dso_local <8 x i16> @test_ret_v8u16(
+// CHECK-LABEL: define dso_local noundef <8 x i16> @test_ret_v8u16(
 // CHECK-SAME: <8 x i16> noundef returned [[V:%.*]]) local_unnamed_addr #[[ATTR0]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    ret <8 x i16> [[V]]
@@ -117,7 +117,7 @@ uint16x8_t test_ret_v8u16(uint16x8_t v) {
   return v;
 }
 
-// CHECK-LABEL: define dso_local <4 x i32> @test_ret_v32u4(
+// CHECK-LABEL: define dso_local noundef <4 x i32> @test_ret_v32u4(
 // CHECK-SAME: <4 x i32> noundef returned [[V:%.*]]) local_unnamed_addr #[[ATTR0]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    ret <4 x i32> [[V]]
@@ -126,7 +126,7 @@ uint32x4_t test_ret_v32u4(uint32x4_t v) {
   return v;
 }
 
-// CHECK-LABEL: define dso_local <2 x i64> @test_ret_v64u2(
+// CHECK-LABEL: define dso_local noundef <2 x i64> @test_ret_v64u2(
 // CHECK-SAME: <2 x i64> noundef returned [[V:%.*]]) local_unnamed_addr #[[ATTR0]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    ret <2 x i64> [[V]]

diff --git a/clang/test/CodeGen/ifunc.c b/clang/test/CodeGen/ifunc.c
@@ -51,11 +51,11 @@ void* goo_ifunc(void) {
 // CHECK: call i32 @foo(i32
 // CHECK: call void @goo()
 
-// SAN: define internal nonnull ptr @foo_ifunc() #[[#FOO_IFUNC:]] {
-// MACSAN: define internal nonnull ptr @foo_ifunc() #[[#FOO_IFUNC:]] {
+// SAN: define internal nonnull {{(noundef )?}}ptr @foo_ifunc() #[[#FOO_IFUNC:]] {
+// MACSAN: define internal nonnull {{(noundef )?}}ptr @foo_ifunc() #[[#FOO_IFUNC:]] {
 
-// SAN: define dso_local noalias ptr @goo_ifunc() #[[#GOO_IFUNC:]] {
-// MACSAN: define noalias ptr @goo_ifunc() #[[#GOO_IFUNC:]] {
+// SAN: define dso_local noalias {{(noundef )?}}ptr @goo_ifunc() #[[#GOO_IFUNC:]] {
+// MACSAN: define noalias {{(noundef )?}}ptr @goo_ifunc() #[[#GOO_IFUNC:]] {
 
 // SAN-DAG: attributes #[[#FOO_IFUNC]] = {{{.*}} disable_sanitizer_instrumentation {{.*}}
 // MACSAN-DAG: attributes #[[#FOO_IFUNC]] = {{{.*}} disable_sanitizer_instrumentation {{.*}}

diff --git a/clang/test/CodeGen/isfpclass.c b/clang/test/CodeGen/isfpclass.c
@@ -1,7 +1,7 @@
 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 2
 // RUN: %clang_cc1 -triple aarch64-linux-gnu -S -O1 -emit-llvm %s -o - | FileCheck %s
 
-// CHECK-LABEL: define dso_local i1 @check_isfpclass_finite
+// CHECK-LABEL: define dso_local noundef i1 @check_isfpclass_finite
 // CHECK-SAME: (float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[TMP0:%.*]] = tail call float @llvm.fabs.f32(float [[X]])
@@ -12,7 +12,7 @@ _Bool check_isfpclass_finite(float x) {
   return __builtin_isfpclass(x, 504 /*Finite*/);
 }
 
-// CHECK-LABEL: define dso_local i1 @check_isfpclass_finite_strict
+// CHECK-LABEL: define dso_local noundef i1 @check_isfpclass_finite_strict
 // CHECK-SAME: (float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR2:[0-9]+]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[TMP0:%.*]] = tail call i1 @llvm.is.fpclass.f32(float [[X]], i32 504) #[[ATTR6:[0-9]+]]
@@ -23,7 +23,7 @@ _Bool check_isfpclass_finite_strict(float x) {
   return __builtin_isfpclass(x, 504 /*Finite*/);
 }
 
-// CHECK-LABEL: define dso_local i1 @check_isfpclass_nan_f32
+// CHECK-LABEL: define dso_local noundef i1 @check_isfpclass_nan_f32
 // CHECK-SAME: (float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3:[0-9]+]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[TMP0:%.*]] = fcmp uno float [[X]], 0.000000e+00
@@ -33,7 +33,7 @@ _Bool check_isfpclass_nan_f32(float x) {
   return __builtin_isfpclass(x, 3 /*NaN*/);
 }
 
-// CHECK-LABEL: define dso_local i1 @check_isfpclass_nan_f32_strict
+// CHECK-LABEL: define dso_local noundef i1 @check_isfpclass_nan_f32_strict
 // CHECK-SAME: (float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR2]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[TMP0:%.*]] = tail call i1 @llvm.is.fpclass.f32(float [[X]], i32 3) #[[ATTR6]]
@@ -44,7 +44,7 @@ _Bool check_isfpclass_nan_f32_strict(float x) {
   return __builtin_isfpclass(x, 3 /*NaN*/);
 }
 
-// CHECK-LABEL: define dso_local i1 @check_isfpclass_snan_f64
+// CHECK-LABEL: define dso_local noundef i1 @check_isfpclass_snan_f64
 // CHECK-SAME: (double noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[TMP0:%.*]] = tail call i1 @llvm.is.fpclass.f64(double [[X]], i32 1)
@@ -54,7 +54,7 @@ _Bool check_isfpclass_snan_f64(double x) {
   return __builtin_isfpclass(x, 1 /*SNaN*/);
 }
 
-// CHECK-LABEL: define dso_local i1 @check_isfpclass_snan_f64_strict
+// CHECK-LABEL: define dso_local noundef i1 @check_isfpclass_snan_f64_strict
 // CHECK-SAME: (double noundef [[X:%.*]]) local_unnamed_addr #[[ATTR2]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[TMP0:%.*]] = tail call i1 @llvm.is.fpclass.f64(double [[X]], i32 1) #[[ATTR6]]
@@ -65,7 +65,7 @@ _Bool check_isfpclass_snan_f64_strict(double x) {
   return __builtin_isfpclass(x, 1 /*NaN*/);
 }
 
-// CHECK-LABEL: define dso_local i1 @check_isfpclass_zero_f16
+// CHECK-LABEL: define dso_local noundef i1 @check_isfpclass_zero_f16
 // CHECK-SAME: (half noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[TMP0:%.*]] = fcmp oeq half [[X]], 0xH0000
@@ -75,7 +75,7 @@ _Bool check_isfpclass_zero_f16(_Float16 x) {
   return __builtin_isfpclass(x, 96 /*Zero*/);
 }
 
-// CHECK-LABEL: define dso_local i1 @check_isfpclass_zero_f16_strict
+// CHECK-LABEL: define dso_local noundef i1 @check_isfpclass_zero_f16_strict
 // CHECK-SAME: (half noundef [[X:%.*]]) local_unnamed_addr #[[ATTR2]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[TMP0:%.*]] = tail call i1 @llvm.is.fpclass.f16(half [[X]], i32 96) #[[ATTR6]]
@@ -86,7 +86,7 @@ _Bool check_isfpclass_zero_f16_strict(_Float16 x) {
   return __builtin_isfpclass(x, 96 /*Zero*/);
 }
 
-// CHECK-LABEL: define dso_local i1 @check_isnan
+// CHECK-LABEL: define dso_local noundef i1 @check_isnan
 // CHECK-SAME: (float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR2]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[TMP0:%.*]] = tail call i1 @llvm.is.fpclass.f32(float [[X]], i32 3) #[[ATTR6]]
@@ -97,7 +97,7 @@ _Bool check_isnan(float x) {
   return __builtin_isnan(x);
 }
 
-// CHECK-LABEL: define dso_local i1 @check_isinf
+// CHECK-LABEL: define dso_local noundef i1 @check_isinf
 // CHECK-SAME: (float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR2]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[TMP0:%.*]] = tail call i1 @llvm.is.fpclass.f32(float [[X]], i32 516) #[[ATTR6]]
@@ -108,7 +108,7 @@ _Bool check_isinf(float x) {
   return __builtin_isinf(x);
 }
 
-// CHECK-LABEL: define dso_local i1 @check_isfinite
+// CHECK-LABEL: define dso_local noundef i1 @check_isfinite
 // CHECK-SAME: (float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR2]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[TMP0:%.*]] = tail call i1 @llvm.is.fpclass.f32(float [[X]], i32 504) #[[ATTR6]]
@@ -119,7 +119,7 @@ _Bool check_isfinite(float x) {
   return __builtin_isfinite(x);
 }
 
-// CHECK-LABEL: define dso_local i1 @check_isnormal
+// CHECK-LABEL: define dso_local noundef i1 @check_isnormal
 // CHECK-SAME: (float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR2]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[TMP0:%.*]] = tail call i1 @llvm.is.fpclass.f32(float [[X]], i32 264) #[[ATTR6]]
@@ -136,7 +136,7 @@ typedef double __attribute__((ext_vector_type(4))) double4;
 typedef int __attribute__((ext_vector_type(4))) int4;
 typedef long __attribute__((ext_vector_type(4))) long4;
 
-// CHECK-LABEL: define dso_local <4 x i32> @check_isfpclass_nan_v4f32
+// CHECK-LABEL: define dso_local noundef <4 x i32> @check_isfpclass_nan_v4f32
 // CHECK-SAME: (<4 x float> noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[TMP0:%.*]] = fcmp uno <4 x float> [[X]], zeroinitializer
@@ -147,7 +147,7 @@ int4 check_isfpclass_nan_v4f32(float4 x) {
   return __builtin_isfpclass(x, 3 /*NaN*/);
 }
 
-// CHECK-LABEL: define dso_local <4 x i32> @check_isfpclass_nan_strict_v4f32
+// CHECK-LABEL: define dso_local noundef <4 x i32> @check_isfpclass_nan_strict_v4f32
 // CHECK-SAME: (<4 x float> noundef [[X:%.*]]) local_unnamed_addr #[[ATTR2]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[TMP0:%.*]] = tail call <4 x i1> @llvm.is.fpclass.v4f32(<4 x float> [[X]], i32 3) #[[ATTR6]]

diff --git a/clang/test/CodeGen/ms-mixed-ptr-sizes.c b/clang/test/CodeGen/ms-mixed-ptr-sizes.c
@@ -49,7 +49,7 @@ void test_other(struct Foo *f, __attribute__((address_space(10))) int *i) {
 }
 
 int test_compare1(int *__ptr32 __uptr i, int *__ptr64 j) {
-  // ALL-LABEL: define dso_local i32 @test_compare1
+  // ALL-LABEL: define dso_local noundef i32 @test_compare1
   // X64: %{{.+}} = addrspacecast ptr %j to ptr addrspace(271)
   // X64: %cmp = icmp eq ptr addrspace(271) %{{.+}}, %i
   // X86: %{{.+}} = addrspacecast ptr addrspace(272) %j to ptr addrspace(271)
@@ -58,7 +58,7 @@ int test_compare1(int *__ptr32 __uptr i, int *__ptr64 j) {
 }
 
 int test_compare2(int *__ptr32 __sptr i, int *__ptr64 j) {
-  // ALL-LABEL: define dso_local i32 @test_compare2
+  // ALL-LABEL: define dso_local noundef i32 @test_compare2
   // X64: %{{.+}} = addrspacecast ptr %j to ptr addrspace(270)
   // X64: %cmp = icmp eq ptr addrspace(270) %{{.+}}, %i
   // X86: %{{.+}} = addrspacecast ptr addrspace(272) %j to ptr
@@ -67,7 +67,7 @@ int test_compare2(int *__ptr32 __sptr i, int *__ptr64 j) {
 }
 
 int test_compare3(int *__ptr32 __uptr i, int *__ptr64 j) {
-  // ALL-LABEL: define dso_local i32 @test_compare3
+  // ALL-LABEL: define dso_local noundef i32 @test_compare3
   // X64: %{{.+}} = addrspacecast ptr addrspace(271) %i to ptr
   // X64: %cmp = icmp eq ptr %{{.+}}, %j
   // X86: %{{.+}} = addrspacecast ptr addrspace(271) %i to ptr addrspace(272)
@@ -76,7 +76,7 @@ int test_compare3(int *__ptr32 __uptr i, int *__ptr64 j) {
 }
 
 int test_compare4(int *__ptr32 __sptr i, int *__ptr64 j) {
-  // ALL-LABEL: define dso_local i32 @test_compare4
+  // ALL-LABEL: define dso_local noundef i32 @test_compare4
   // X64: %{{.+}} = addrspacecast ptr addrspace(270) %i to ptr
   // X64: %cmp = icmp eq ptr %{{.+}}, %j
   // X86: %{{.+}} = addrspacecast ptr %i to ptr addrspace(272)

diff --git a/clang/test/CodeGenCUDA/link-builtin-bitcode-denormal-fp-mode.cu b/clang/test/CodeGenCUDA/link-builtin-bitcode-denormal-fp-mode.cu
@@ -111,16 +111,16 @@ __global__ void kernel_f64(double* out, double* a, double* b, double* c) {
 }
 }
 
-// INTERNALIZE: define internal half @do_f16_stuff({{.*}}) #[[$FUNCATTR:[0-9]+]]
-// INTERNALIZE: define internal float @do_f32_stuff({{.*}}) #[[$FUNCATTR]]
-// INTERNALIZE: define internal double @do_f64_stuff({{.*}}) #[[$FUNCATTR]]
-// INTERNALIZE: define internal float @weak_do_f32_stuff({{.*}}) #[[$WEAK_FUNCATTR:[0-9]+]]
+// INTERNALIZE: define internal {{(noundef )?}}half @do_f16_stuff({{.*}}) #[[$FUNCATTR:[0-9]+]]
+// INTERNALIZE: define internal {{(noundef )?}}float @do_f32_stuff({{.*}}) #[[$FUNCATTR]]
+// INTERNALIZE: define internal {{(noundef )?}}double @do_f64_stuff({{.*}}) #[[$FUNCATTR]]
+// INTERNALIZE: define internal {{(noundef )?}}float @weak_do_f32_stuff({{.*}}) #[[$WEAK_FUNCATTR:[0-9]+]]
 
 
-// NOINTERNALIZE: define dso_local half @do_f16_stuff({{.*}}) #[[$FUNCATTR:[0-9]+]]
-// NOINTERNALIZE: define dso_local float @do_f32_stuff({{.*}}) #[[$FUNCATTR]]
-// NOINTERNALIZE: define dso_local double @do_f64_stuff({{.*}}) #[[$FUNCATTR]]
-// NOINTERNALIZE: define weak float @weak_do_f32_stuff({{.*}}) #[[$WEAK_FUNCATTR:[0-9]+]]
+// NOINTERNALIZE: define dso_local {{(noundef )?}}half @do_f16_stuff({{.*}}) #[[$FUNCATTR:[0-9]+]]
+// NOINTERNALIZE: define dso_local {{(noundef )?}}float @do_f32_stuff({{.*}}) #[[$FUNCATTR]]
+// NOINTERNALIZE: define dso_local {{(noundef )?}}double @do_f64_stuff({{.*}}) #[[$FUNCATTR]]
+// NOINTERNALIZE: define weak {{(noundef )?}}float @weak_do_f32_stuff({{.*}}) #[[$WEAK_FUNCATTR:[0-9]+]]