@@ -1075,7 +1075,6 @@ extern "C" __device__ double test_cospi(double x) {
10751075 return cospi (x);
10761076}
10771077
1078- //
10791078// DEFAULT-LABEL: @test_cyl_bessel_i0f(
10801079// DEFAULT-NEXT: entry:
10811080// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_i0_f32(float noundef [[X:%.*]]) #[[ATTR14]]
@@ -1748,7 +1747,6 @@ extern "C" __device__ double test_fmax(double x, double y) {
17481747 return fmax (x, y);
17491748}
17501749
1751- //
17521750// DEFAULT-LABEL: @test_fminf(
17531751// DEFAULT-NEXT: entry:
17541752// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
@@ -3086,10 +3084,9 @@ extern "C" __device__ long int test_lround(double x) {
30863084// AMDGCNSPIRV-LABEL: @test_modff(
30873085// AMDGCNSPIRV-NEXT: entry:
30883086// AMDGCNSPIRV-NEXT: [[__TMP_I:%.*]] = alloca float, align 4
3089- // AMDGCNSPIRV-NEXT: [[__TMP_ASCAST_I:%.*]] = addrspacecast ptr [[__TMP_I]] to ptr addrspace(4)
30903087// AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.lifetime.start.p0(i64 4, ptr nonnull [[__TMP_I]]) #[[ATTR15:[0-9]+]]
30913088// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = call contract spir_func noundef addrspace(4) float @__ocml_modf_f32(float noundef [[X:%.*]], ptr noundef nonnull [[__TMP_I]]) #[[ATTR14]]
3092- // AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(4) [[__TMP_ASCAST_I ]], align 4, !tbaa [[TBAA17:![0-9]+]]
3089+ // AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load float, ptr [[__TMP_I ]], align 4, !tbaa [[TBAA17:![0-9]+]]
30933090// AMDGCNSPIRV-NEXT: store float [[TMP0]], ptr addrspace(4) [[Y:%.*]], align 4, !tbaa [[TBAA17]]
30943091// AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.lifetime.end.p0(i64 4, ptr nonnull [[__TMP_I]]) #[[ATTR15]]
30953092// AMDGCNSPIRV-NEXT: ret float [[CALL_I]]
@@ -3131,10 +3128,9 @@ extern "C" __device__ float test_modff(float x, float* y) {
31313128// AMDGCNSPIRV-LABEL: @test_modf(
31323129// AMDGCNSPIRV-NEXT: entry:
31333130// AMDGCNSPIRV-NEXT: [[__TMP_I:%.*]] = alloca double, align 8
3134- // AMDGCNSPIRV-NEXT: [[__TMP_ASCAST_I:%.*]] = addrspacecast ptr [[__TMP_I]] to ptr addrspace(4)
31353131// AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.lifetime.start.p0(i64 8, ptr nonnull [[__TMP_I]]) #[[ATTR15]]
31363132// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = call contract spir_func noundef addrspace(4) double @__ocml_modf_f64(double noundef [[X:%.*]], ptr noundef nonnull [[__TMP_I]]) #[[ATTR14]]
3137- // AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(4) [[__TMP_ASCAST_I ]], align 8, !tbaa [[TBAA19:![0-9]+]]
3133+ // AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load double, ptr [[__TMP_I ]], align 8, !tbaa [[TBAA19:![0-9]+]]
31383134// AMDGCNSPIRV-NEXT: store double [[TMP0]], ptr addrspace(4) [[Y:%.*]], align 8, !tbaa [[TBAA19]]
31393135// AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.lifetime.end.p0(i64 8, ptr nonnull [[__TMP_I]]) #[[ATTR15]]
31403136// AMDGCNSPIRV-NEXT: ret double [[CALL_I]]
@@ -4471,10 +4467,9 @@ extern "C" __device__ double test_remainder(double x, double y) {
44714467// AMDGCNSPIRV-LABEL: @test_remquof(
44724468// AMDGCNSPIRV-NEXT: entry:
44734469// AMDGCNSPIRV-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4
4474- // AMDGCNSPIRV-NEXT: [[__TMP_ASCAST_I:%.*]] = addrspacecast ptr [[__TMP_I]] to ptr addrspace(4)
44754470// AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.lifetime.start.p0(i64 4, ptr nonnull [[__TMP_I]]) #[[ATTR15]]
44764471// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = call contract spir_func noundef addrspace(4) float @__ocml_remquo_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], ptr noundef nonnull [[__TMP_I]]) #[[ATTR14]]
4477- // AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(4) [[__TMP_ASCAST_I ]], align 4, !tbaa [[TBAA13]]
4472+ // AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load i32, ptr [[__TMP_I ]], align 4, !tbaa [[TBAA13]]
44784473// AMDGCNSPIRV-NEXT: store i32 [[TMP0]], ptr addrspace(4) [[Z:%.*]], align 4, !tbaa [[TBAA13]]
44794474// AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.lifetime.end.p0(i64 4, ptr nonnull [[__TMP_I]]) #[[ATTR15]]
44804475// AMDGCNSPIRV-NEXT: ret float [[CALL_I]]
@@ -4516,10 +4511,9 @@ extern "C" __device__ float test_remquof(float x, float y, int* z) {
45164511// AMDGCNSPIRV-LABEL: @test_remquo(
45174512// AMDGCNSPIRV-NEXT: entry:
45184513// AMDGCNSPIRV-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4
4519- // AMDGCNSPIRV-NEXT: [[__TMP_ASCAST_I:%.*]] = addrspacecast ptr [[__TMP_I]] to ptr addrspace(4)
45204514// AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.lifetime.start.p0(i64 4, ptr nonnull [[__TMP_I]]) #[[ATTR15]]
45214515// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = call contract spir_func noundef addrspace(4) double @__ocml_remquo_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], ptr noundef nonnull [[__TMP_I]]) #[[ATTR14]]
4522- // AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(4) [[__TMP_ASCAST_I ]], align 4, !tbaa [[TBAA13]]
4516+ // AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load i32, ptr [[__TMP_I ]], align 4, !tbaa [[TBAA13]]
45234517// AMDGCNSPIRV-NEXT: store i32 [[TMP0]], ptr addrspace(4) [[Z:%.*]], align 4, !tbaa [[TBAA13]]
45244518// AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.lifetime.end.p0(i64 4, ptr nonnull [[__TMP_I]]) #[[ATTR15]]
45254519// AMDGCNSPIRV-NEXT: ret double [[CALL_I]]
@@ -5230,11 +5224,10 @@ extern "C" __device__ BOOL_TYPE test___signbit(double x) {
52305224// AMDGCNSPIRV-LABEL: @test_sincosf(
52315225// AMDGCNSPIRV-NEXT: entry:
52325226// AMDGCNSPIRV-NEXT: [[__TMP_I:%.*]] = alloca float, align 4
5233- // AMDGCNSPIRV-NEXT: [[__TMP_ASCAST_I:%.*]] = addrspacecast ptr [[__TMP_I]] to ptr addrspace(4)
52345227// AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.lifetime.start.p0(i64 4, ptr nonnull [[__TMP_I]]) #[[ATTR15]]
52355228// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = call contract spir_func addrspace(4) float @__ocml_sincos_f32(float noundef [[X:%.*]], ptr noundef nonnull [[__TMP_I]]) #[[ATTR14]]
52365229// AMDGCNSPIRV-NEXT: store float [[CALL_I]], ptr addrspace(4) [[Y:%.*]], align 4, !tbaa [[TBAA17]]
5237- // AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(4) [[__TMP_ASCAST_I ]], align 4, !tbaa [[TBAA17]]
5230+ // AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load float, ptr [[__TMP_I ]], align 4, !tbaa [[TBAA17]]
52385231// AMDGCNSPIRV-NEXT: store float [[TMP0]], ptr addrspace(4) [[Z:%.*]], align 4, !tbaa [[TBAA17]]
52395232// AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.lifetime.end.p0(i64 4, ptr nonnull [[__TMP_I]]) #[[ATTR15]]
52405233// AMDGCNSPIRV-NEXT: ret void
@@ -5279,11 +5272,10 @@ extern "C" __device__ void test_sincosf(float x, float *y, float *z) {
52795272// AMDGCNSPIRV-LABEL: @test_sincos(
52805273// AMDGCNSPIRV-NEXT: entry:
52815274// AMDGCNSPIRV-NEXT: [[__TMP_I:%.*]] = alloca double, align 8
5282- // AMDGCNSPIRV-NEXT: [[__TMP_ASCAST_I:%.*]] = addrspacecast ptr [[__TMP_I]] to ptr addrspace(4)
52835275// AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.lifetime.start.p0(i64 8, ptr nonnull [[__TMP_I]]) #[[ATTR15]]
52845276// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = call contract spir_func addrspace(4) double @__ocml_sincos_f64(double noundef [[X:%.*]], ptr noundef nonnull [[__TMP_I]]) #[[ATTR14]]
52855277// AMDGCNSPIRV-NEXT: store double [[CALL_I]], ptr addrspace(4) [[Y:%.*]], align 8, !tbaa [[TBAA19]]
5286- // AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(4) [[__TMP_ASCAST_I ]], align 8, !tbaa [[TBAA19]]
5278+ // AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load double, ptr [[__TMP_I ]], align 8, !tbaa [[TBAA19]]
52875279// AMDGCNSPIRV-NEXT: store double [[TMP0]], ptr addrspace(4) [[Z:%.*]], align 8, !tbaa [[TBAA19]]
52885280// AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.lifetime.end.p0(i64 8, ptr nonnull [[__TMP_I]]) #[[ATTR15]]
52895281// AMDGCNSPIRV-NEXT: ret void
@@ -5328,11 +5320,10 @@ extern "C" __device__ void test_sincos(double x, double *y, double *z) {
53285320// AMDGCNSPIRV-LABEL: @test_sincospif(
53295321// AMDGCNSPIRV-NEXT: entry:
53305322// AMDGCNSPIRV-NEXT: [[__TMP_I:%.*]] = alloca float, align 4
5331- // AMDGCNSPIRV-NEXT: [[__TMP_ASCAST_I:%.*]] = addrspacecast ptr [[__TMP_I]] to ptr addrspace(4)
53325323// AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.lifetime.start.p0(i64 4, ptr nonnull [[__TMP_I]]) #[[ATTR15]]
53335324// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = call contract spir_func addrspace(4) float @__ocml_sincospi_f32(float noundef [[X:%.*]], ptr noundef nonnull [[__TMP_I]]) #[[ATTR14]]
53345325// AMDGCNSPIRV-NEXT: store float [[CALL_I]], ptr addrspace(4) [[Y:%.*]], align 4, !tbaa [[TBAA17]]
5335- // AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(4) [[__TMP_ASCAST_I ]], align 4, !tbaa [[TBAA17]]
5326+ // AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load float, ptr [[__TMP_I ]], align 4, !tbaa [[TBAA17]]
53365327// AMDGCNSPIRV-NEXT: store float [[TMP0]], ptr addrspace(4) [[Z:%.*]], align 4, !tbaa [[TBAA17]]
53375328// AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.lifetime.end.p0(i64 4, ptr nonnull [[__TMP_I]]) #[[ATTR15]]
53385329// AMDGCNSPIRV-NEXT: ret void
@@ -5377,11 +5368,10 @@ extern "C" __device__ void test_sincospif(float x, float *y, float *z) {
53775368// AMDGCNSPIRV-LABEL: @test_sincospi(
53785369// AMDGCNSPIRV-NEXT: entry:
53795370// AMDGCNSPIRV-NEXT: [[__TMP_I:%.*]] = alloca double, align 8
5380- // AMDGCNSPIRV-NEXT: [[__TMP_ASCAST_I:%.*]] = addrspacecast ptr [[__TMP_I]] to ptr addrspace(4)
53815371// AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.lifetime.start.p0(i64 8, ptr nonnull [[__TMP_I]]) #[[ATTR15]]
53825372// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = call contract spir_func addrspace(4) double @__ocml_sincospi_f64(double noundef [[X:%.*]], ptr noundef nonnull [[__TMP_I]]) #[[ATTR14]]
53835373// AMDGCNSPIRV-NEXT: store double [[CALL_I]], ptr addrspace(4) [[Y:%.*]], align 8, !tbaa [[TBAA19]]
5384- // AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(4) [[__TMP_ASCAST_I ]], align 8, !tbaa [[TBAA19]]
5374+ // AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load double, ptr [[__TMP_I ]], align 8, !tbaa [[TBAA19]]
53855375// AMDGCNSPIRV-NEXT: store double [[TMP0]], ptr addrspace(4) [[Z:%.*]], align 8, !tbaa [[TBAA19]]
53865376// AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.lifetime.end.p0(i64 8, ptr nonnull [[__TMP_I]]) #[[ATTR15]]
53875377// AMDGCNSPIRV-NEXT: ret void
0 commit comments