Skip to content

Commit 48f4e6d

Browse files
committed
[IR] Allow fast math flags on fptosi, fptoui and sitofp
1 parent add906f commit 48f4e6d

31 files changed

+439
-363
lines changed

clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -164,7 +164,7 @@ __device__ double ffp4(double *p, float f) {
164164

165165
__device__ double ffp5(double *p, int i) {
166166
// FUN-LABEL: @_Z4ffp5Pdi
167-
// CHECK: sitofp i32 {{.*}} to double
167+
// CHECK: sitofp contract i32 {{.*}} to double
168168
// SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, [[DEFMD]]
169169
// UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, [[DEFMD]]
170170
__atomic_fetch_sub(p, i, memory_order_relaxed);

clang/test/CodeGenCUDA/managed-var.cu

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -145,7 +145,7 @@ float load3() {
145145
// HOST: %4 = ptrtoint ptr %3 to i64
146146
// HOST: %5 = sub i64 %4, %1
147147
// HOST: %sub.ptr.div = sdiv exact i64 %5, 4
148-
// HOST: %conv = sitofp i64 %sub.ptr.div to float
148+
// HOST: %conv = sitofp contract i64 %sub.ptr.div to float
149149
// HOST: ret float %conv
150150
float addr_taken2() {
151151
return (float)reinterpret_cast<long>(&(v2[1].y)-&(v[1].x));

clang/test/CodeGenHLSL/BasicFeatures/AggregateSplatCast.hlsl

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -38,7 +38,7 @@ export void call8() {
3838
// CHECK-NEXT: store <1 x float> splat (float 1.000000e+00), ptr [[B]], align 4
3939
// CHECK-NEXT: [[L:%.*]] = load <1 x float>, ptr [[B]], align 4
4040
// CHECK-NEXT: [[VL:%.*]] = extractelement <1 x float> [[L]], i32 0
41-
// CHECK-NEXT: [[C:%.*]] = fptosi float [[VL]] to i32
41+
// CHECK-NEXT: [[C:%.*]] = fptosi reassoc nnan ninf nsz arcp afn float [[VL]] to i32
4242
// CHECK-NEXT: [[SI:%.*]] = insertelement <4 x i32> poison, i32 [[C]], i64 0
4343
// CHECK-NEXT: [[S:%.*]] = shufflevector <4 x i32> [[SI]], <4 x i32> poison, <4 x i32> zeroinitializer
4444
// CHECK-NEXT: store <4 x i32> [[S]], ptr [[A]], align 16
@@ -62,7 +62,7 @@ struct S {
6262
// CHECK-NEXT: [[G1:%.*]] = getelementptr inbounds %struct.S, ptr [[s]], i32 0, i32 0
6363
// CHECK-NEXT: [[G2:%.*]] = getelementptr inbounds %struct.S, ptr [[s]], i32 0, i32 1
6464
// CHECK-NEXT: store i32 [[VL]], ptr [[G1]], align 4
65-
// CHECK-NEXT: [[C:%.*]] = sitofp i32 [[VL]] to float
65+
// CHECK-NEXT: [[C:%.*]] = sitofp reassoc nnan ninf nsz arcp afn i32 [[VL]] to float
6666
// CHECK-NEXT: store float [[C]], ptr [[G2]], align 4
6767
export void call3() {
6868
int1 A = {1};
@@ -79,7 +79,7 @@ export void call3() {
7979
// CHECK-NEXT: [[G1:%.*]] = getelementptr inbounds %struct.S, ptr [[s]], i32 0, i32 0
8080
// CHECK-NEXT: [[G2:%.*]] = getelementptr inbounds %struct.S, ptr [[s]], i32 0, i32 1
8181
// CHECK-NEXT: store i32 [[VL]], ptr [[G1]], align 4
82-
// CHECK-NEXT: [[C:%.*]] = sitofp i32 [[VL]] to float
82+
// CHECK-NEXT: [[C:%.*]] = sitofp reassoc nnan ninf nsz arcp afn i32 [[VL]] to float
8383
// CHECK-NEXT: store float [[C]], ptr [[G2]], align 4
8484
export void call5() {
8585
int1 A = {1};

clang/test/CodeGenHLSL/BasicFeatures/ArrayElementwiseCast.hlsl

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -46,7 +46,7 @@ export void call1() {
4646
// CHECK-NEXT: [[G1:%.*]] = getelementptr inbounds [1 x float], ptr [[B]], i32 0, i32 0
4747
// CHECK-NEXT: [[G2:%.*]] = getelementptr inbounds [1 x i32], ptr [[Tmp]], i32 0, i32 0
4848
// CHECK-NEXT: [[L:%.*]] = load i32, ptr [[G2]], align 4
49-
// CHECK-NEXT: [[C:%.*]] = sitofp i32 [[L]] to float
49+
// CHECK-NEXT: [[C:%.*]] = sitofp reassoc nnan ninf nsz arcp afn i32 [[L]] to float
5050
// CHECK-NEXT: store float [[C]], ptr [[G1]], align 4
5151
export void call2() {
5252
int A[1] = {0};
@@ -63,7 +63,7 @@ export void call2() {
6363
// CHECK-NEXT: [[C:%.*]] = load <1 x float>, ptr [[A]], align 4
6464
// CHECK-NEXT: [[G1:%.*]] = getelementptr inbounds [1 x i32], ptr [[B]], i32 0, i32 0
6565
// CHECK-NEXT: [[V:%.*]] = extractelement <1 x float> [[C]], i64 0
66-
// CHECK-NEXT: [[C:%.*]] = fptosi float [[V]] to i32
66+
// CHECK-NEXT: [[C:%.*]] = fptosi reassoc nnan ninf nsz arcp afn float [[V]] to i32
6767
// CHECK-NEXT: store i32 [[C]], ptr [[G1]], align 4
6868
export void call3() {
6969
float1 A = {1.2};
@@ -84,11 +84,11 @@ export void call3() {
8484
// CHECK-NEXT: [[VG:%.*]] = getelementptr inbounds [1 x <2 x float>], ptr [[Tmp]], i32 0, i32 0
8585
// CHECK-NEXT: [[L:%.*]] = load <2 x float>, ptr [[VG]], align 8
8686
// CHECK-NEXT: [[VL:%.*]] = extractelement <2 x float> [[L]], i32 0
87-
// CHECK-NEXT: [[C:%.*]] = fptosi float [[VL]] to i32
87+
// CHECK-NEXT: [[C:%.*]] = fptosi reassoc nnan ninf nsz arcp afn float [[VL]] to i32
8888
// CHECK-NEXT: store i32 [[C]], ptr [[G1]], align 4
8989
// CHECK-NEXT: [[L4:%.*]] = load <2 x float>, ptr [[VG]], align 8
9090
// CHECK-NEXT: [[VL5:%.*]] = extractelement <2 x float> [[L4]], i32 1
91-
// CHECK-NEXT: [[C6:%.*]] = fptosi float [[VL5]] to i32
91+
// CHECK-NEXT: [[C6:%.*]] = fptosi reassoc nnan ninf nsz arcp afn float [[VL5]] to i32
9292
// CHECK-NEXT: store i32 [[C6]], ptr [[G2]], align 4
9393
export void call5() {
9494
float2 A[1] = {{1.2,3.4}};

clang/test/CodeGenHLSL/BasicFeatures/InitLists.hlsl

Lines changed: 12 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -98,7 +98,7 @@ TwoFloats case2() {
9898
// CHECK-NEXT: store i32 [[VAL]], ptr [[VAL_ADDR]], align 4
9999
// CHECK-NEXT: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_TWOFLOATS]], ptr [[AGG_RESULT]], i32 0, i32 0
100100
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[VAL_ADDR]], align 4
101-
// CHECK-NEXT: [[CONV:%.*]] = sitofp i32 [[TMP0]] to float
101+
// CHECK-NEXT: [[CONV:%.*]] = sitofp reassoc nnan ninf nsz arcp afn i32 [[TMP0]] to float
102102
// CHECK-NEXT: store float [[CONV]], ptr [[X]], align 1
103103
// CHECK-NEXT: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_TWOFLOATS]], ptr [[AGG_RESULT]], i32 0, i32 1
104104
// CHECK-NEXT: store float 2.000000e+00, ptr [[Y]], align 1
@@ -119,12 +119,12 @@ TwoFloats case3(int Val) {
119119
// CHECK-NEXT: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_TWOFLOATS]], ptr [[AGG_RESULT]], i32 0, i32 0
120120
// CHECK-NEXT: [[TMP0:%.*]] = load <2 x i32>, ptr [[TWOVALS_ADDR]], align 8
121121
// CHECK-NEXT: [[VECEXT:%.*]] = extractelement <2 x i32> [[TMP0]], i64 0
122-
// CHECK-NEXT: [[CONV:%.*]] = sitofp i32 [[VECEXT]] to float
122+
// CHECK-NEXT: [[CONV:%.*]] = sitofp reassoc nnan ninf nsz arcp afn i32 [[VECEXT]] to float
123123
// CHECK-NEXT: store float [[CONV]], ptr [[X]], align 1
124124
// CHECK-NEXT: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_TWOFLOATS]], ptr [[AGG_RESULT]], i32 0, i32 1
125125
// CHECK-NEXT: [[TMP1:%.*]] = load <2 x i32>, ptr [[TWOVALS_ADDR]], align 8
126126
// CHECK-NEXT: [[VECEXT1:%.*]] = extractelement <2 x i32> [[TMP1]], i64 1
127-
// CHECK-NEXT: [[CONV2:%.*]] = sitofp i32 [[VECEXT1]] to float
127+
// CHECK-NEXT: [[CONV2:%.*]] = sitofp reassoc nnan ninf nsz arcp afn i32 [[VECEXT1]] to float
128128
// CHECK-NEXT: store float [[CONV2]], ptr [[Y]], align 1
129129
// CHECK-NEXT: ret void
130130
//
@@ -162,12 +162,12 @@ TwoInts case5(int2 TwoVals) {
162162
// CHECK-NEXT: [[Z:%.*]] = getelementptr inbounds nuw [[STRUCT_TWOINTS]], ptr [[AGG_RESULT]], i32 0, i32 0
163163
// CHECK-NEXT: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_TWOFLOATS]], ptr [[TF4]], i32 0, i32 0
164164
// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[X]], align 1
165-
// CHECK-NEXT: [[CONV:%.*]] = fptosi float [[TMP0]] to i32
165+
// CHECK-NEXT: [[CONV:%.*]] = fptosi reassoc nnan ninf nsz arcp afn float [[TMP0]] to i32
166166
// CHECK-NEXT: store i32 [[CONV]], ptr [[Z]], align 1
167167
// CHECK-NEXT: [[W:%.*]] = getelementptr inbounds nuw [[STRUCT_TWOINTS]], ptr [[AGG_RESULT]], i32 0, i32 1
168168
// CHECK-NEXT: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_TWOFLOATS]], ptr [[TF4]], i32 0, i32 1
169169
// CHECK-NEXT: [[TMP1:%.*]] = load float, ptr [[Y]], align 1
170-
// CHECK-NEXT: [[CONV1:%.*]] = fptosi float [[TMP1]] to i32
170+
// CHECK-NEXT: [[CONV1:%.*]] = fptosi reassoc nnan ninf nsz arcp afn float [[TMP1]] to i32
171171
// CHECK-NEXT: store i32 [[CONV1]], ptr [[W]], align 1
172172
// CHECK-NEXT: ret void
173173
//
@@ -202,7 +202,7 @@ TwoInts case6(TwoFloats TF4) {
202202
// CHECK-NEXT: store i32 [[TMP4]], ptr [[TAILSTATE]], align 1
203203
// CHECK-NEXT: [[HAIRCOUNT:%.*]] = getelementptr inbounds nuw [[STRUCT_DOGGO]], ptr [[AGG_RESULT]], i32 0, i32 2
204204
// CHECK-NEXT: [[TMP5:%.*]] = load i32, ptr [[VAL_ADDR]], align 4
205-
// CHECK-NEXT: [[CONV:%.*]] = sitofp i32 [[TMP5]] to float
205+
// CHECK-NEXT: [[CONV:%.*]] = sitofp reassoc nnan ninf nsz arcp afn i32 [[TMP5]] to float
206206
// CHECK-NEXT: store float [[CONV]], ptr [[HAIRCOUNT]], align 1
207207
// CHECK-NEXT: [[EARDIRECTION:%.*]] = getelementptr inbounds nuw [[STRUCT_DOGGO]], ptr [[AGG_RESULT]], i32 0, i32 3
208208
// CHECK-NEXT: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_TWOFLOATS]], ptr [[TF1]], i32 0, i32 0
@@ -272,7 +272,7 @@ Doggo case7(TwoInts TI1, TwoInts TI2, int Val, TwoFloats TF1, TwoFloats TF2,
272272
// CHECK-NEXT: [[COUNTER:%.*]] = getelementptr inbounds nuw [[STRUCT_ANIMALBITS]], ptr [[AGG_RESULT]], i32 0, i32 2
273273
// CHECK-NEXT: [[HAIRCOUNT:%.*]] = getelementptr inbounds nuw [[STRUCT_DOGGO]], ptr [[D1]], i32 0, i32 2
274274
// CHECK-NEXT: [[TMP5:%.*]] = load float, ptr [[HAIRCOUNT]], align 1
275-
// CHECK-NEXT: [[CONV:%.*]] = fptosi float [[TMP5]] to i64
275+
// CHECK-NEXT: [[CONV:%.*]] = fptosi reassoc nnan ninf nsz arcp afn float [[TMP5]] to i64
276276
// CHECK-NEXT: store i64 [[CONV]], ptr [[COUNTER]], align 1
277277
// CHECK-NEXT: [[LEFTDIR:%.*]] = getelementptr inbounds nuw [[STRUCT_ANIMALBITS]], ptr [[AGG_RESULT]], i32 0, i32 3
278278
// CHECK-NEXT: [[EARDIRECTION:%.*]] = getelementptr inbounds nuw [[STRUCT_DOGGO]], ptr [[D1]], i32 0, i32 3
@@ -428,7 +428,7 @@ AnimalBits case8(Doggo D1) {
428428
// CHECK-NEXT: [[HAIRCOUNT58:%.*]] = getelementptr inbounds nuw [[STRUCT_DOGGO]], ptr [[ARRAYINIT_ELEMENT44]], i32 0, i32 2
429429
// CHECK-NEXT: [[COUNTER:%.*]] = getelementptr inbounds nuw [[STRUCT_ANIMALBITS]], ptr [[A1]], i32 0, i32 2
430430
// CHECK-NEXT: [[TMP19:%.*]] = load i64, ptr [[COUNTER]], align 1
431-
// CHECK-NEXT: [[CONV:%.*]] = sitofp i64 [[TMP19]] to float
431+
// CHECK-NEXT: [[CONV:%.*]] = sitofp reassoc nnan ninf nsz arcp afn i64 [[TMP19]] to float
432432
// CHECK-NEXT: store float [[CONV]], ptr [[HAIRCOUNT58]], align 1
433433
// CHECK-NEXT: [[EARDIRECTION59:%.*]] = getelementptr inbounds nuw [[STRUCT_DOGGO]], ptr [[ARRAYINIT_ELEMENT44]], i32 0, i32 3
434434
// CHECK-NEXT: [[LEFTDIR:%.*]] = getelementptr inbounds nuw [[STRUCT_ANIMALBITS]], ptr [[A1]], i32 0, i32 3
@@ -563,7 +563,7 @@ AnimalBits case8(Doggo D1) {
563563
// CHECK-NEXT: [[HAIRCOUNT149:%.*]] = getelementptr inbounds nuw [[STRUCT_KITTEH]], ptr [[ARRAYINIT_ELEMENT133]], i32 0, i32 2
564564
// CHECK-NEXT: [[COUNTER150:%.*]] = getelementptr inbounds nuw [[STRUCT_ANIMALBITS]], ptr [[A1]], i32 0, i32 2
565565
// CHECK-NEXT: [[TMP47:%.*]] = load i64, ptr [[COUNTER150]], align 1
566-
// CHECK-NEXT: [[CONV151:%.*]] = sitofp i64 [[TMP47]] to float
566+
// CHECK-NEXT: [[CONV151:%.*]] = sitofp reassoc nnan ninf nsz arcp afn i64 [[TMP47]] to float
567567
// CHECK-NEXT: store float [[CONV151]], ptr [[HAIRCOUNT149]], align 1
568568
// CHECK-NEXT: [[CLAWS152:%.*]] = getelementptr inbounds nuw [[STRUCT_KITTEH]], ptr [[ARRAYINIT_ELEMENT133]], i32 0, i32 3
569569
// CHECK-NEXT: [[LEFTDIR153:%.*]] = getelementptr inbounds nuw [[STRUCT_ANIMALBITS]], ptr [[A1]], i32 0, i32 3
@@ -698,7 +698,7 @@ AnimalBits case8(Doggo D1) {
698698
// CHECK-NEXT: [[HAIRCOUNT246:%.*]] = getelementptr inbounds nuw [[STRUCT_KITTEH]], ptr [[ARRAYINIT_ELEMENT230]], i32 0, i32 2
699699
// CHECK-NEXT: [[COUNTER247:%.*]] = getelementptr inbounds nuw [[STRUCT_ANIMALBITS]], ptr [[A1]], i32 0, i32 2
700700
// CHECK-NEXT: [[TMP75:%.*]] = load i64, ptr [[COUNTER247]], align 1
701-
// CHECK-NEXT: [[CONV248:%.*]] = sitofp i64 [[TMP75]] to float
701+
// CHECK-NEXT: [[CONV248:%.*]] = sitofp reassoc nnan ninf nsz arcp afn i64 [[TMP75]] to float
702702
// CHECK-NEXT: store float [[CONV248]], ptr [[HAIRCOUNT246]], align 1
703703
// CHECK-NEXT: [[CLAWS249:%.*]] = getelementptr inbounds nuw [[STRUCT_KITTEH]], ptr [[ARRAYINIT_ELEMENT230]], i32 0, i32 3
704704
// CHECK-NEXT: [[LEFTDIR250:%.*]] = getelementptr inbounds nuw [[STRUCT_ANIMALBITS]], ptr [[A1]], i32 0, i32 3
@@ -888,13 +888,13 @@ TwoInts case14(SlicyBits SB) {
888888
// CHECK-NEXT: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_TWOFLOATS]], ptr [[AGG_RESULT]], i32 0, i32 0
889889
// CHECK-NEXT: [[BF_LOAD:%.*]] = load i8, ptr [[SB]], align 1
890890
// CHECK-NEXT: [[BF_CAST:%.*]] = sext i8 [[BF_LOAD]] to i32
891-
// CHECK-NEXT: [[CONV:%.*]] = sitofp i32 [[BF_CAST]] to float
891+
// CHECK-NEXT: [[CONV:%.*]] = sitofp reassoc nnan ninf nsz arcp afn i32 [[BF_CAST]] to float
892892
// CHECK-NEXT: store float [[CONV]], ptr [[X]], align 1
893893
// CHECK-NEXT: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_TWOFLOATS]], ptr [[AGG_RESULT]], i32 0, i32 1
894894
// CHECK-NEXT: [[W:%.*]] = getelementptr inbounds nuw [[STRUCT_SLICYBITS]], ptr [[SB]], i32 0, i32 1
895895
// CHECK-NEXT: [[BF_LOAD1:%.*]] = load i8, ptr [[W]], align 1
896896
// CHECK-NEXT: [[BF_CAST2:%.*]] = sext i8 [[BF_LOAD1]] to i32
897-
// CHECK-NEXT: [[CONV3:%.*]] = sitofp i32 [[BF_CAST2]] to float
897+
// CHECK-NEXT: [[CONV3:%.*]] = sitofp reassoc nnan ninf nsz arcp afn i32 [[BF_CAST2]] to float
898898
// CHECK-NEXT: store float [[CONV3]], ptr [[Y]], align 1
899899
// CHECK-NEXT: ret void
900900
//

clang/test/CodeGenHLSL/BasicFeatures/OutputArguments.hlsl

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -13,14 +13,14 @@ void trunc_Param(inout int X) {}
1313
// CHECK: [[F:%.*]] = alloca float
1414
// CHECK: [[ArgTmp:%.*]] = alloca i32
1515
// CHECK: [[FVal:%.*]] = load float, ptr {{.*}}
16-
// CHECK: [[IVal:%.*]] = fptosi float [[FVal]] to i32
16+
// CHECK: [[IVal:%.*]] = fptosi reassoc nnan ninf nsz arcp afn float [[FVal]] to i32
1717
// CHECK: store i32 [[IVal]], ptr [[ArgTmp]]
1818
// CHECK: call void {{.*}}trunc_Param{{.*}}(ptr noalias noundef nonnull align 4 dereferenceable(4) [[ArgTmp]])
1919
// CHECK: [[IRet:%.*]] = load i32, ptr [[ArgTmp]]
20-
// CHECK: [[FRet:%.*]] = sitofp i32 [[IRet]] to float
20+
// CHECK: [[FRet:%.*]] = sitofp reassoc nnan ninf nsz arcp afn i32 [[IRet]] to float
2121
// CHECK: store float [[FRet]], ptr [[F]]
22-
// OPT: [[IVal:%.*]] = fptosi float {{.*}} to i32
23-
// OPT: [[FVal:%.*]] = sitofp i32 [[IVal]] to float
22+
// OPT: [[IVal:%.*]] = fptosi reassoc nnan ninf nsz arcp afn float {{.*}} to i32
23+
// OPT: [[FVal:%.*]] = sitofp reassoc nnan ninf nsz arcp afn i32 [[IVal]] to float
2424
// OPT: ret float [[FVal]]
2525
export float case1(float F) {
2626
trunc_Param(F);
@@ -202,15 +202,15 @@ void trunc_vec(inout int3 V) {}
202202
// CHECK: [[V:%.*]] = alloca <3 x float>
203203
// CHECK: [[Tmp:%.*]] = alloca <3 x i32>
204204
// CHECK: [[FVal:%.*]] = load <3 x float>, ptr [[V]]
205-
// CHECK: [[IVal:%.*]] = fptosi <3 x float> [[FVal]] to <3 x i32>
205+
// CHECK: [[IVal:%.*]] = fptosi reassoc nnan ninf nsz arcp afn <3 x float> [[FVal]] to <3 x i32>
206206
// CHECK: store <3 x i32> [[IVal]], ptr [[Tmp]]
207207
// CHECK: call void {{.*}}trunc_vec{{.*}}(ptr noalias noundef nonnull align 16 dereferenceable(16) [[Tmp]])
208208
// CHECK: [[IRet:%.*]] = load <3 x i32>, ptr [[Tmp]]
209-
// CHECK: [[FRet:%.*]] = sitofp <3 x i32> [[IRet]] to <3 x float>
209+
// CHECK: [[FRet:%.*]] = sitofp reassoc nnan ninf nsz arcp afn <3 x i32> [[IRet]] to <3 x float>
210210
// CHECK: store <3 x float> [[FRet]], ptr [[V]]
211211

212-
// OPT: [[IVal:%.*]] = fptosi <3 x float> {{.*}} to <3 x i32>
213-
// OPT: [[FVal:%.*]] = sitofp <3 x i32> [[IVal]] to <3 x float>
212+
// OPT: [[IVal:%.*]] = fptosi reassoc nnan ninf nsz arcp afn <3 x float> {{.*}} to <3 x i32>
213+
// OPT: [[FVal:%.*]] = sitofp reassoc nnan ninf nsz arcp afn <3 x i32> [[IVal]] to <3 x float>
214214
// OPT: ret <3 x float> [[FVal]]
215215

216216
export float3 case8(float3 V) {

clang/test/CodeGenHLSL/BasicFeatures/StructElementwiseCast.hlsl

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -32,7 +32,7 @@ export void call0() {
3232
// CHECK-NEXT: [[VL:%.*]] = extractelement <2 x i32> [[L]], i64 0
3333
// CHECK-NEXT: store i32 [[VL]], ptr [[G1]], align 4
3434
// CHECK-NEXT: [[VL2:%.*]] = extractelement <2 x i32> [[L]], i64 1
35-
// CHECK-NEXT: [[C:%.*]] = sitofp i32 [[VL2]] to float
35+
// CHECK-NEXT: [[C:%.*]] = sitofp reassoc nnan ninf nsz arcp afn i32 [[VL2]] to float
3636
// CHECK-NEXT: store float [[C]], ptr [[G2]], align 4
3737
export void call1() {
3838
int2 A = {1,2};
@@ -54,7 +54,7 @@ export void call1() {
5454
// CHECK-NEXT: [[L:%.*]] = load i32, ptr [[G3]], align 4
5555
// CHECK-NEXT: store i32 [[L]], ptr [[G1]], align 4
5656
// CHECK-NEXT: [[L4:%.*]] = load i32, ptr [[G4]], align 4
57-
// CHECK-NEXT: [[C:%.*]] = sitofp i32 [[L4]] to float
57+
// CHECK-NEXT: [[C:%.*]] = sitofp reassoc nnan ninf nsz arcp afn i32 [[L4]] to float
5858
// CHECK-NEXT: store float [[C]], ptr [[G2]], align 4
5959
export void call2() {
6060
int A[2] = {1,2};
@@ -104,7 +104,7 @@ export void call6() {
104104
// CHECK-NEXT: [[L:%.*]] = load i32, ptr [[G3]], align 4
105105
// CHECK-NEXT: store i32 [[L]], ptr [[G1]], align 4
106106
// CHECK-NEXT: [[L4:%.*]] = load i32, ptr [[G4]], align 4
107-
// CHECK-NEXT: [[C:%.*]] = sitofp i32 [[L4]] to float
107+
// CHECK-NEXT: [[C:%.*]] = sitofp reassoc nnan ninf nsz arcp afn i32 [[L4]] to float
108108
// CHECK-NEXT: store float [[C]], ptr [[G2]], align 4
109109
export void call7() {
110110
int A[2] = {1,2};
@@ -132,7 +132,7 @@ struct T {
132132
// CHECK-NEXT: %load = load i32, ptr %gep2, align 4
133133
// CHECK-NEXT: store i32 %load, ptr %gep, align 4
134134
// CHECK-NEXT: %load5 = load i32, ptr %gep3, align 4
135-
// CHECK-NEXT: %conv = sitofp i32 %load5 to float
135+
// CHECK-NEXT: %conv = sitofp reassoc nnan ninf nsz arcp afn i32 %load5 to float
136136
// CHECK-NEXT: store float %conv, ptr %gep1, align 4
137137
export void call8() {
138138
T t = {1,2,3};

clang/test/CodeGenHLSL/BasicFeatures/VectorElementwiseCast.hlsl

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -40,7 +40,7 @@ struct S {
4040
// CHECK-NEXT: [[L:%.*]] = load i32, ptr [[G1]], align 4
4141
// CHECK-NEXT: [[C:%.*]] = insertelement <2 x i32> [[B]], i32 [[L]], i64 0
4242
// CHECK-NEXT: [[L2:%.*]] = load float, ptr [[G2]], align 4
43-
// CHECK-NEXT: [[D:%.*]] = fptosi float [[L2]] to i32
43+
// CHECK-NEXT: [[D:%.*]] = fptosi reassoc nnan ninf nsz arcp afn float [[L2]] to i32
4444
// CHECK-NEXT: [[E:%.*]] = insertelement <2 x i32> [[C]], i32 [[D]], i64 1
4545
// CHECK-NEXT: store <2 x i32> [[E]], ptr [[A]], align 8
4646
export void call3() {

clang/test/CodeGenHLSL/BasicFeatures/standard_conversion_sequences.hlsl

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -43,7 +43,7 @@ void d4_to_f2() {
4343
// CHECK: [[i2:%.*]] = alloca <2 x i32>
4444
// CHECK: store <2 x float> splat (float 4.000000e+00), ptr [[f2]]
4545
// CHECK: [[vecf2:%.*]] = load <2 x float>, ptr [[f2]]
46-
// CHECK: [[veci2:%.*]] = fptosi <2 x float> [[vecf2]] to <2 x i32>
46+
// CHECK: [[veci2:%.*]] = fptosi reassoc nnan ninf nsz arcp afn <2 x float> [[vecf2]] to <2 x i32>
4747
// CHECK: store <2 x i32> [[veci2]], ptr [[i2]]
4848
void f2_to_i2() {
4949
vector<float,2> f2 = 4.0;
@@ -55,7 +55,7 @@ void f2_to_i2() {
5555
// CHECK: [[i2:%.*]] = alloca <2 x i32>
5656
// CHECK: store <4 x double> splat (double 5.000000e+00), ptr [[d4]]
5757
// CHECK: [[vecd4:%.*]] = load <4 x double>, ptr [[d4]]
58-
// CHECK: [[veci4:%.*]] = fptosi <4 x double> [[vecd4]] to <4 x i32>
58+
// CHECK: [[veci4:%.*]] = fptosi reassoc nnan ninf nsz arcp afn <4 x double> [[vecd4]] to <4 x i32>
5959
// CHECK: [[veci2:%.*]] = shufflevector <4 x i32> [[veci4]], <4 x i32> poison, <2 x i32> <i32 0, i32 1>
6060
// CHECK: store <2 x i32> [[veci2]], ptr [[i2]]
6161
void d4_to_i2() {
@@ -68,7 +68,7 @@ void d4_to_i2() {
6868
// CHECK: [[l4:%.*]] = alloca <4 x i64>
6969
// CHECK: store <4 x double> splat (double 6.000000e+00), ptr [[d4]]
7070
// CHECK: [[vecd4:%.*]] = load <4 x double>, ptr [[d4]]
71-
// CHECK: [[vecl4:%.*]] = fptosi <4 x double> [[vecd4]] to <4 x i64>
71+
// CHECK: [[vecl4:%.*]] = fptosi reassoc nnan ninf nsz arcp afn <4 x double> [[vecd4]] to <4 x i64>
7272
// CHECK: store <4 x i64> [[vecl4]], ptr [[l4]]
7373
void d4_to_l4() {
7474
vector<double,4> d4 = 6.0;

0 commit comments

Comments
 (0)