@@ -3676,7 +3676,6 @@ extern "C" __device__ long int test_lround(double x) {
36763676// AMDGCNSPIRV-LABEL: @test_modff(
36773677// AMDGCNSPIRV-NEXT: entry:
36783678// AMDGCNSPIRV-NEXT: [[__TMP_I:%.*]] = alloca float, align 4
3679- // AMDGCNSPIRV-NEXT: [[__TMP_ASCAST_I:%.*]] = addrspacecast ptr [[__TMP_I]] to ptr addrspace(4) ; TODO: remove??
36803679// AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.lifetime.start.p0(ptr nonnull [[__TMP_I]]) #[[ATTR15:[0-9]+]]
36813680// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = call contract spir_func noundef addrspace(4) float @__ocml_modf_f32(float noundef [[X:%.*]], ptr noundef nonnull [[__TMP_I]]) #[[ATTR14]]
36823681// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load float, ptr [[__TMP_I]], align 4, !tbaa [[TBAA17:![0-9]+]]
@@ -3731,7 +3730,6 @@ extern "C" __device__ float test_modff(float x, float* y) {
37313730// AMDGCNSPIRV-LABEL: @test_modf(
37323731// AMDGCNSPIRV-NEXT: entry:
37333732// AMDGCNSPIRV-NEXT: [[__TMP_I:%.*]] = alloca double, align 8
3734- // AMDGCNSPIRV-NEXT: [[__TMP_ASCAST_I:%.*]] = addrspacecast ptr [[__TMP_I]] to ptr addrspace(4) ; TODO: remove??
37353733// AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.lifetime.start.p0(ptr nonnull [[__TMP_I]]) #[[ATTR15]]
37363734// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = call contract spir_func noundef addrspace(4) double @__ocml_modf_f64(double noundef [[X:%.*]], ptr noundef nonnull [[__TMP_I]]) #[[ATTR14]]
37373735// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load double, ptr [[__TMP_I]], align 8, !tbaa [[TBAA19:![0-9]+]]
@@ -5402,7 +5400,6 @@ extern "C" __device__ double test_remainder(double x, double y) {
54025400// AMDGCNSPIRV-LABEL: @test_remquof(
54035401// AMDGCNSPIRV-NEXT: entry:
54045402// AMDGCNSPIRV-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4
5405- // AMDGCNSPIRV-NEXT: [[__TMP_ASCAST_I:%.*]] = addrspacecast ptr [[__TMP_I]] to ptr addrspace(4) ; TODO: remove??
54065403// AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.lifetime.start.p0(ptr nonnull [[__TMP_I]]) #[[ATTR15]]
54075404// 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]]
54085405// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load i32, ptr [[__TMP_I]], align 4, !tbaa [[TBAA13]]
@@ -5457,7 +5454,6 @@ extern "C" __device__ float test_remquof(float x, float y, int* z) {
54575454// AMDGCNSPIRV-LABEL: @test_remquo(
54585455// AMDGCNSPIRV-NEXT: entry:
54595456// AMDGCNSPIRV-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4
5460- // AMDGCNSPIRV-NEXT: [[__TMP_ASCAST_I:%.*]] = addrspacecast ptr [[__TMP_I]] to ptr addrspace(4) ; TODO: remove??
54615457// AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.lifetime.start.p0(ptr nonnull [[__TMP_I]]) #[[ATTR15]]
54625458// 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]]
54635459// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load i32, ptr [[__TMP_I]], align 4, !tbaa [[TBAA13]]
@@ -6242,7 +6238,6 @@ extern "C" __device__ BOOL_TYPE test___signbit(double x) {
62426238// AMDGCNSPIRV-LABEL: @test_sincosf(
62436239// AMDGCNSPIRV-NEXT: entry:
62446240// AMDGCNSPIRV-NEXT: [[__TMP_I:%.*]] = alloca float, align 4
6245- // AMDGCNSPIRV-NEXT: [[__TMP_ASCAST_I:%.*]] = addrspacecast ptr [[__TMP_I]] to ptr addrspace(4) ; TODO: remove??
62466241// AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.lifetime.start.p0(ptr nonnull [[__TMP_I]]) #[[ATTR15]]
62476242// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = call contract spir_func addrspace(4) float @__ocml_sincos_f32(float noundef [[X:%.*]], ptr noundef nonnull [[__TMP_I]]) #[[ATTR14]]
62486243// AMDGCNSPIRV-NEXT: store float [[CALL_I]], ptr addrspace(4) [[Y:%.*]], align 4, !tbaa [[TBAA17]]
@@ -6302,7 +6297,6 @@ extern "C" __device__ void test_sincosf(float x, float *y, float *z) {
63026297// AMDGCNSPIRV-LABEL: @test_sincos(
63036298// AMDGCNSPIRV-NEXT: entry:
63046299// AMDGCNSPIRV-NEXT: [[__TMP_I:%.*]] = alloca double, align 8
6305- // AMDGCNSPIRV-NEXT: [[__TMP_ASCAST_I:%.*]] = addrspacecast ptr [[__TMP_I]] to ptr addrspace(4) ; TODO: remove??
63066300// AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.lifetime.start.p0(ptr nonnull [[__TMP_I]]) #[[ATTR15]]
63076301// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = call contract spir_func addrspace(4) double @__ocml_sincos_f64(double noundef [[X:%.*]], ptr noundef nonnull [[__TMP_I]]) #[[ATTR14]]
63086302// AMDGCNSPIRV-NEXT: store double [[CALL_I]], ptr addrspace(4) [[Y:%.*]], align 8, !tbaa [[TBAA19]]
@@ -6362,7 +6356,6 @@ extern "C" __device__ void test_sincos(double x, double *y, double *z) {
63626356// AMDGCNSPIRV-LABEL: @test_sincospif(
63636357// AMDGCNSPIRV-NEXT: entry:
63646358// AMDGCNSPIRV-NEXT: [[__TMP_I:%.*]] = alloca float, align 4
6365- // AMDGCNSPIRV-NEXT: [[__TMP_ASCAST_I:%.*]] = addrspacecast ptr [[__TMP_I]] to ptr addrspace(4) ; TODO: remove??
63666359// AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.lifetime.start.p0(ptr nonnull [[__TMP_I]]) #[[ATTR15]]
63676360// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = call contract spir_func addrspace(4) float @__ocml_sincospi_f32(float noundef [[X:%.*]], ptr noundef nonnull [[__TMP_I]]) #[[ATTR14]]
63686361// AMDGCNSPIRV-NEXT: store float [[CALL_I]], ptr addrspace(4) [[Y:%.*]], align 4, !tbaa [[TBAA17]]
@@ -6422,7 +6415,6 @@ extern "C" __device__ void test_sincospif(float x, float *y, float *z) {
64226415// AMDGCNSPIRV-LABEL: @test_sincospi(
64236416// AMDGCNSPIRV-NEXT: entry:
64246417// AMDGCNSPIRV-NEXT: [[__TMP_I:%.*]] = alloca double, align 8
6425- // AMDGCNSPIRV-NEXT: [[__TMP_ASCAST_I:%.*]] = addrspacecast ptr [[__TMP_I]] to ptr addrspace(4) ; TODO: remove??
64266418// AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.lifetime.start.p0(ptr nonnull [[__TMP_I]]) #[[ATTR15]]
64276419// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = call contract spir_func addrspace(4) double @__ocml_sincospi_f64(double noundef [[X:%.*]], ptr noundef nonnull [[__TMP_I]]) #[[ATTR14]]
64286420// AMDGCNSPIRV-NEXT: store double [[CALL_I]], ptr addrspace(4) [[Y:%.*]], align 8, !tbaa [[TBAA19]]
0 commit comments