PR clean up

Author: YixingZhang007

test/extensions/INTEL/SPV_INTEL_joint_matrix/joint_matrix_load_store_offset.ll

@@ -11,29 +11,28 @@
 ; CHECK-SPIRV-DAG: Capability CooperativeMatrixKHR
 ; CHECK-SPIRV-DAG: Capability CooperativeMatrixOffsetInstructionsINTEL
 ; CHECK-SPIRV-DAG: Extension "SPV_INTEL_joint_matrix"
-; CHECK-SPIRV-DAG: TypeInt [[#Int16Ty:]] 16 0  // 18
-; CHECK-SPIRV-DAG: TypeInt [[#Int32Ty:]] 32 0  // 106
-; CHECK-SPIRV-DAG: Constant [[#Int32Ty]] [[#Const16:]] 16 // 112
-; CHECK-SPIRV-DAG: Constant [[#Int32Ty]] [[#Const3:]] 3  // 113
-; CHECK-SPIRV-DAG: Constant [[#Int32Ty]] [[#Const2:]] 2 // 114
-; CHECK-SPIRV-DAG: Constant [[#Int32Ty]] [[#Const0:]] 0 // 116
-; CHECK-SPIRV-DAG: Constant [[#Int32Ty]] [[#Const1:]] 1 // 133
-; CHECK-SPIRV-DAG: TypeFloat [[#Float32Ty:]] 32  // 8
-
-; CHECK-SPIRV-DAG: TypeJointMatrixINTEL [[#MatTy1:]] [[#Float32Ty]] [[#Const16]] [[#Const16]] [[#Const3]] [[#Const3]] [[#Const2]] //TypeJointMatrixINTEL 115 8 112 112 113 113 114  
-; CHECK-SPIRV-DAG: TypeJointMatrixINTEL [[#MatTy2:]] [[#Int16Ty]] [[#Const16]] [[#Const16]] [[#Const0]] [[#Const3]] [[#Const0]] [[#Const1]] // TypeJointMatrixINTEL 134 18 112 112 116 113 116 133 
-; CHECK-SPIRV-DAG: TypeJointMatrixINTEL [[#MatTy3:]] [[#Int16Ty]] [[#Const16]] [[#Const16]] [[#Const2:]] [[#Const3:]] [[#Const1:]] [[#Const1:]] // TypeJointMatrixINTEL 137 18 112 112 114 113 133 133
-; CHECK-SPIRV: CooperativeMatrixLoadOffsetINTEL [[#MatTy1]] // CooperativeMatrixLoadOffsetINTEL 115 118 104 109 111 116 117 116 
-; CHECK-SPIRV: CooperativeMatrixLoadOffsetINTEL [[#MatTy2]] // CooperativeMatrixLoadOffsetINTEL 134 135 121 109 132 116 117 116
-; CHECK-SPIRV: CooperativeMatrixLoadOffsetINTEL [[#MatTy3]] // CooperativeMatrixLoadOffsetINTEL 137 139 124 136 126 114 138 116
-; CHECK-SPIRV: JointMatrixMadINTEL [[#MatTy1]] //JointMatrixMadINTEL 115 127 135 139 128 113 
-; CHECK-SPIRV: CooperativeMatrixStoreOffsetINTEL // CooperativeMatrixStoreOffsetINTEL 104 109 111 128 116 117 116 
+; CHECK-SPIRV-DAG: TypeInt [[#Int16Ty:]] 16 0
+; CHECK-SPIRV-DAG: TypeInt [[#Int32Ty:]] 32 0
+; CHECK-SPIRV-DAG: Constant [[#Int32Ty]] [[#Const16:]] 16
+; CHECK-SPIRV-DAG: Constant [[#Int32Ty]] [[#Const3:]] 3
+; CHECK-SPIRV-DAG: Constant [[#Int32Ty]] [[#Const2:]] 2
+; CHECK-SPIRV-DAG: Constant [[#Int32Ty]] [[#Const0:]] 0
+; CHECK-SPIRV-DAG: Constant [[#Int32Ty]] [[#Const1:]] 1
+; CHECK-SPIRV-DAG: TypeFloat [[#Float32Ty:]] 32
+; CHECK-SPIRV-DAG: TypeJointMatrixINTEL [[#MatTy1:]] [[#Float32Ty]] [[#Const16]] [[#Const16]] [[#Const3]] [[#Const3]] [[#Const2]]
+; CHECK-SPIRV-DAG: TypeJointMatrixINTEL [[#MatTy2:]] [[#Int16Ty]] [[#Const16]] [[#Const16]] [[#Const0]] [[#Const3]] [[#Const0]] [[#Const1]]
+; CHECK-SPIRV-DAG: TypeJointMatrixINTEL [[#MatTy3:]] [[#Int16Ty]] [[#Const16]] [[#Const16]] [[#Const2:]] [[#Const3:]] [[#Const1:]] [[#Const1:]]
+; CHECK-SPIRV: CooperativeMatrixLoadOffsetINTEL [[#MatTy1]]
+; CHECK-SPIRV: CooperativeMatrixLoadOffsetINTEL [[#MatTy2]]
+; CHECK-SPIRV: CooperativeMatrixLoadOffsetINTEL [[#MatTy3]]
+; CHECK-SPIRV: JointMatrixMadINTEL [[#MatTy1]]
+; CHECK-SPIRV: CooperativeMatrixStoreOffsetINTEL
 
 ; CHECK-LLVM: call spir_func target("spirv.JointMatrixINTEL", float, 16, 16, 3, 3, 2) @_Z93__spirv_CooperativeMatrixLoadOffsetINTEL_RPU3AS143__spirv_JointMatrixINTEL__float_16_16_3_3_2PU3AS1fiiili(ptr addrspace(1) %add.ptr.i.i, i32 %conv.i, i32 %conv2.i, i32 0, i64 32, i32 0)
 ; CHECK-LLVM: call spir_func target("spirv.JointMatrixINTEL", i16, 16, 16, 0, 3, 0, 1) @"_Z95__spirv_CooperativeMatrixLoadOffsetINTEL_RPU3AS145__spirv_JointMatrixINTEL__short_16_16_0_3_0_1PU3AS138class.sycl::_V1::ext::oneapi::bfloat16iiili"(ptr addrspace(1) %add.ptr.i.i120, i32 %conv.i, i32 %22, i32 0, i64 32, i32 0)
 ; CHECK-LLVM: call spir_func target("spirv.JointMatrixINTEL", i16, 16, 16, 2, 3, 1, 1) @"_Z95__spirv_CooperativeMatrixLoadOffsetINTEL_RPU3AS145__spirv_JointMatrixINTEL__short_16_16_2_3_1_1PU3AS138class.sycl::_V1::ext::oneapi::bfloat16iiili"(ptr addrspace(1) %add.ptr.i.i128, i32 %23, i32 %conv2.i60, i32 2, i64 64, i32 0)
 ; CHECK-LLVM: call spir_func target("spirv.JointMatrixINTEL", float, 16, 16, 3, 3, 2) @_Z27__spirv_JointMatrixMadINTELPU3AS145__spirv_JointMatrixINTEL__short_16_16_0_3_0_1PU3AS145__spirv_JointMatrixINTEL__short_16_16_2_3_1_1PU3AS143__spirv_JointMatrixINTEL__float_16_16_3_3_2i(target("spirv.JointMatrixINTEL", i16, 16, 16, 0, 3, 0, 1) %call3.i50, target("spirv.JointMatrixINTEL", i16, 16, 16, 2, 3, 1, 1) %call3.i61, target("spirv.JointMatrixINTEL", float, 16, 16, 3, 3, 2) %sub_c.i.sroa.0.0, i32 3)
-
+; CHECK-LLVM: call spir_func void @_Z41__spirv_CooperativeMatrixStoreOffsetINTELPU3AS1fiiPU3AS143__spirv_JointMatrixINTEL__float_16_16_3_3_2ili(ptr addrspace(1) %add.ptr.i.i, i32 %conv.i, i32 %conv2.i, target("spirv.JointMatrixINTEL", float, 16, 16, 3, 3, 2) %sub_c.i.sroa.0.0, i32 0, i64 32, i32 0)
 
 ; ModuleID = 'joint_matrix_bfloat16.cpp'
 source_filename = "joint_matrix_bfloat16.cpp"
@@ -147,9 +146,7 @@ declare dso_local spir_func noundef target("spirv.JointMatrixINTEL", float, 16,
 ; Function Attrs: convergent nounwind
 declare dso_local spir_func void @_Z41__spirv_CooperativeMatrixStoreOffsetINTELIU3AS1ffLm16ELm16ELN5__spv9MatrixUseE2ELNS1_12MatrixLayoutE3ELNS1_5Scope4FlagE3EEvPT_iiPNS1_24__spirv_JointMatrixINTELIT0_XT1_EXT2_EXT4_EXT5_EXT3_EEES3_mi(ptr addrspace(1) noundef, i32 noundef, i32 noundef, target("spirv.JointMatrixINTEL", float, 16, 16, 3, 3, 2) noundef, i32 noundef, i64 noundef, i32 noundef) local_unnamed_addr #2
 
-declare dso_local spir_func i32 @_Z18__spirv_ocl_printfPU3AS2Kcz(ptr addrspace(2), ...)
-
-attributes #0 = { convergent mustprogress norecurse nounwind "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="../joint_matrix_bfloat16.cpp" "sycl-optlevel"="2" "uniform-work-group-size"="true" }
+attributes #0 = { convergent mustprogress norecurse nounwind "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="joint_matrix_bfloat16.cpp" "sycl-optlevel"="2" "uniform-work-group-size"="true" }
 attributes #1 = { mustprogress nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: write) }
 attributes #2 = { convergent nounwind "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
 attributes #3 = { convergent nounwind }