@@ -4983,7 +4983,7 @@ extern "C" __device__ double test_normcdfinv(double x) {
4983
4983
// DEFAULT-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
4984
4984
// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT_LOOPEXIT:%.*]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]]
4985
4985
// DEFAULT: _ZL5normfiPKf.exit.loopexit:
4986
- // DEFAULT-NEXT: [[TMP1:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[ADD_I]])
4986
+ // DEFAULT-NEXT: [[TMP1:%.*]] = tail call contract float @llvm.sqrt.f32(float [[ADD_I]])
4987
4987
// DEFAULT-NEXT: br label [[_ZL5NORMFIPKF_EXIT]]
4988
4988
// DEFAULT: _ZL5normfiPKf.exit:
4989
4989
// DEFAULT-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[TMP1]], [[_ZL5NORMFIPKF_EXIT_LOOPEXIT]] ]
@@ -5005,7 +5005,7 @@ extern "C" __device__ double test_normcdfinv(double x) {
5005
5005
// FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
5006
5006
// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT_LOOPEXIT:%.*]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]]
5007
5007
// FINITEONLY: _ZL5normfiPKf.exit.loopexit:
5008
- // FINITEONLY-NEXT: [[TMP1:%.*]] = tail call nnan ninf contract noundef float @llvm.sqrt.f32(float [[ADD_I]])
5008
+ // FINITEONLY-NEXT: [[TMP1:%.*]] = tail call nnan ninf contract float @llvm.sqrt.f32(float [[ADD_I]])
5009
5009
// FINITEONLY-NEXT: br label [[_ZL5NORMFIPKF_EXIT]]
5010
5010
// FINITEONLY: _ZL5normfiPKf.exit:
5011
5011
// FINITEONLY-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[TMP1]], [[_ZL5NORMFIPKF_EXIT_LOOPEXIT]] ]
@@ -5027,7 +5027,7 @@ extern "C" __device__ double test_normcdfinv(double x) {
5027
5027
// APPROX-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
5028
5028
// APPROX-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT_LOOPEXIT:%.*]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]]
5029
5029
// APPROX: _ZL5normfiPKf.exit.loopexit:
5030
- // APPROX-NEXT: [[TMP1:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[ADD_I]])
5030
+ // APPROX-NEXT: [[TMP1:%.*]] = tail call contract float @llvm.sqrt.f32(float [[ADD_I]])
5031
5031
// APPROX-NEXT: br label [[_ZL5NORMFIPKF_EXIT]]
5032
5032
// APPROX: _ZL5normfiPKf.exit:
5033
5033
// APPROX-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[TMP1]], [[_ZL5NORMFIPKF_EXIT_LOOPEXIT]] ]
@@ -5049,7 +5049,7 @@ extern "C" __device__ double test_normcdfinv(double x) {
5049
5049
// NCRDIV-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
5050
5050
// NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT_LOOPEXIT:%.*]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
5051
5051
// NCRDIV: _ZL5normfiPKf.exit.loopexit:
5052
- // NCRDIV-NEXT: [[TMP1:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[ADD_I]]), !fpmath [[META22:![0-9]+]]
5052
+ // NCRDIV-NEXT: [[TMP1:%.*]] = tail call contract float @llvm.sqrt.f32(float [[ADD_I]])
5053
5053
// NCRDIV-NEXT: br label [[_ZL5NORMFIPKF_EXIT]]
5054
5054
// NCRDIV: _ZL5normfiPKf.exit:
5055
5055
// NCRDIV-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[TMP1]], [[_ZL5NORMFIPKF_EXIT_LOOPEXIT]] ]
@@ -5071,7 +5071,7 @@ extern "C" __device__ double test_normcdfinv(double x) {
5071
5071
// AMDGCNSPIRV-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
5072
5072
// AMDGCNSPIRV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT_LOOPEXIT:%.*]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
5073
5073
// AMDGCNSPIRV: _ZL5normfiPKf.exit.loopexit:
5074
- // AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = tail call contract noundef addrspace(4) float @llvm.sqrt.f32(float [[ADD_I]])
5074
+ // AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = tail call contract addrspace(4) float @llvm.sqrt.f32(float [[ADD_I]])
5075
5075
// AMDGCNSPIRV-NEXT: br label [[_ZL5NORMFIPKF_EXIT]]
5076
5076
// AMDGCNSPIRV: _ZL5normfiPKf.exit:
5077
5077
// AMDGCNSPIRV-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[TMP1]], [[_ZL5NORMFIPKF_EXIT_LOOPEXIT]] ]
@@ -5097,7 +5097,7 @@ extern "C" __device__ float test_normf(int x, const float *y) {
5097
5097
// DEFAULT-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
5098
5098
// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT_LOOPEXIT:%.*]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
5099
5099
// DEFAULT: _ZL4normiPKd.exit.loopexit:
5100
- // DEFAULT-NEXT: [[TMP1:%.*]] = tail call contract noundef double @llvm.sqrt.f64(double [[ADD_I]])
5100
+ // DEFAULT-NEXT: [[TMP1:%.*]] = tail call contract double @llvm.sqrt.f64(double [[ADD_I]])
5101
5101
// DEFAULT-NEXT: br label [[_ZL4NORMIPKD_EXIT]]
5102
5102
// DEFAULT: _ZL4normiPKd.exit:
5103
5103
// DEFAULT-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[TMP1]], [[_ZL4NORMIPKD_EXIT_LOOPEXIT]] ]
@@ -5119,7 +5119,7 @@ extern "C" __device__ float test_normf(int x, const float *y) {
5119
5119
// FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
5120
5120
// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT_LOOPEXIT:%.*]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
5121
5121
// FINITEONLY: _ZL4normiPKd.exit.loopexit:
5122
- // FINITEONLY-NEXT: [[TMP1:%.*]] = tail call nnan ninf contract noundef double @llvm.sqrt.f64(double [[ADD_I]])
5122
+ // FINITEONLY-NEXT: [[TMP1:%.*]] = tail call nnan ninf contract double @llvm.sqrt.f64(double [[ADD_I]])
5123
5123
// FINITEONLY-NEXT: br label [[_ZL4NORMIPKD_EXIT]]
5124
5124
// FINITEONLY: _ZL4normiPKd.exit:
5125
5125
// FINITEONLY-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[TMP1]], [[_ZL4NORMIPKD_EXIT_LOOPEXIT]] ]
@@ -5141,7 +5141,7 @@ extern "C" __device__ float test_normf(int x, const float *y) {
5141
5141
// APPROX-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
5142
5142
// APPROX-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT_LOOPEXIT:%.*]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
5143
5143
// APPROX: _ZL4normiPKd.exit.loopexit:
5144
- // APPROX-NEXT: [[TMP1:%.*]] = tail call contract noundef double @llvm.sqrt.f64(double [[ADD_I]])
5144
+ // APPROX-NEXT: [[TMP1:%.*]] = tail call contract double @llvm.sqrt.f64(double [[ADD_I]])
5145
5145
// APPROX-NEXT: br label [[_ZL4NORMIPKD_EXIT]]
5146
5146
// APPROX: _ZL4normiPKd.exit:
5147
5147
// APPROX-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[TMP1]], [[_ZL4NORMIPKD_EXIT_LOOPEXIT]] ]
@@ -5161,9 +5161,9 @@ extern "C" __device__ float test_normf(int x, const float *y) {
5161
5161
// NCRDIV-NEXT: [[ADD_I]] = fadd contract double [[__R_0_I4]], [[MUL_I]]
5162
5162
// NCRDIV-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr [[__A_ADDR_0_I3]], i64 8
5163
5163
// NCRDIV-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
5164
- // NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT_LOOPEXIT:%.*]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP23 :![0-9]+]]
5164
+ // NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT_LOOPEXIT:%.*]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22 :![0-9]+]]
5165
5165
// NCRDIV: _ZL4normiPKd.exit.loopexit:
5166
- // NCRDIV-NEXT: [[TMP1:%.*]] = tail call contract noundef double @llvm.sqrt.f64(double [[ADD_I]])
5166
+ // NCRDIV-NEXT: [[TMP1:%.*]] = tail call contract double @llvm.sqrt.f64(double [[ADD_I]])
5167
5167
// NCRDIV-NEXT: br label [[_ZL4NORMIPKD_EXIT]]
5168
5168
// NCRDIV: _ZL4normiPKd.exit:
5169
5169
// NCRDIV-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[TMP1]], [[_ZL4NORMIPKD_EXIT_LOOPEXIT]] ]
@@ -5185,7 +5185,7 @@ extern "C" __device__ float test_normf(int x, const float *y) {
5185
5185
// AMDGCNSPIRV-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
5186
5186
// AMDGCNSPIRV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT_LOOPEXIT:%.*]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]]
5187
5187
// AMDGCNSPIRV: _ZL4normiPKd.exit.loopexit:
5188
- // AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = tail call contract noundef addrspace(4) double @llvm.sqrt.f64(double [[ADD_I]])
5188
+ // AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = tail call contract addrspace(4) double @llvm.sqrt.f64(double [[ADD_I]])
5189
5189
// AMDGCNSPIRV-NEXT: br label [[_ZL4NORMIPKD_EXIT]]
5190
5190
// AMDGCNSPIRV: _ZL4normiPKd.exit:
5191
5191
// AMDGCNSPIRV-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[TMP1]], [[_ZL4NORMIPKD_EXIT_LOOPEXIT]] ]
@@ -5727,7 +5727,7 @@ extern "C" __device__ double test_rint(double x) {
5727
5727
// NCRDIV-NEXT: [[ADD_I]] = fadd contract float [[__R_0_I4]], [[MUL_I]]
5728
5728
// NCRDIV-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr [[__A_ADDR_0_I3]], i64 4
5729
5729
// NCRDIV-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
5730
- // NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP24 :![0-9]+]]
5730
+ // NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP23 :![0-9]+]]
5731
5731
// NCRDIV: _ZL6rnormfiPKf.exit:
5732
5732
// NCRDIV-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
5733
5733
// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_rsqrt_f32(float noundef [[__R_0_I_LCSSA]]) #[[ATTR15]]
@@ -5831,7 +5831,7 @@ extern "C" __device__ float test_rnormf(int x, const float* y) {
5831
5831
// NCRDIV-NEXT: [[ADD_I]] = fadd contract double [[__R_0_I4]], [[MUL_I]]
5832
5832
// NCRDIV-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr [[__A_ADDR_0_I3]], i64 8
5833
5833
// NCRDIV-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
5834
- // NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP25 :![0-9]+]]
5834
+ // NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP24 :![0-9]+]]
5835
5835
// NCRDIV: _ZL5rnormiPKd.exit:
5836
5836
// NCRDIV-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
5837
5837
// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_rsqrt_f64(double noundef [[__R_0_I_LCSSA]]) #[[ATTR15]]
@@ -6636,7 +6636,7 @@ extern "C" __device__ double test_sinpi(double x) {
6636
6636
//
6637
6637
// NCRDIV-LABEL: @test_sqrtf(
6638
6638
// NCRDIV-NEXT: entry:
6639
- // NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[X:%.*]]), !fpmath [[META22 ]]
6639
+ // NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[X:%.*]]), !fpmath [[META25:![0-9]+ ]]
6640
6640
// NCRDIV-NEXT: ret float [[TMP0]]
6641
6641
//
6642
6642
// AMDGCNSPIRV-LABEL: @test_sqrtf(
0 commit comments