From 0b08a2cccf016503a630a6473058f98f9fd1b045 Mon Sep 17 00:00:00 2001 From: Feng Zou Date: Mon, 27 Oct 2025 22:11:49 +0800 Subject: [PATCH] Update LIT tests after commit 1ffff05 https://github.com/llvm/llvm-project/commit/1ffff05: [clang][SPIR][SPIRV] Materialize non-generic null pointers via addrspacecast (#161773) --- sycl/test/check_device_code/group_load.cpp | 161 +++++++++--------- .../group_load_store_alignment.cpp | 24 +-- .../group_load_store_native_key.cpp | 44 ++--- sycl/test/check_device_code/group_store.cpp | 116 ++++++------- 4 files changed, 173 insertions(+), 172 deletions(-) diff --git a/sycl/test/check_device_code/group_load.cpp b/sycl/test/check_device_code/group_load.cpp index 3f03dfcd65135..401479e1c78e8 100644 --- a/sycl/test/check_device_code/group_load.cpp +++ b/sycl/test/check_device_code/group_load.cpp @@ -60,8 +60,8 @@ namespace blocked { // CHECK-GLOBAL-NEXT: [[CALL_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR8:[0-9]+]] // CHECK-GLOBAL-NEXT: [[IDXPROM_I_I:%.*]] = sext i32 [[CALL_I_I_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[P:%.*]], i64 [[IDXPROM_I_I]] -// CHECK-GLOBAL-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(1) [[ARRAYIDX_I_I]], align 4, !tbaa [[TBAA7:![0-9]+]] -// CHECK-GLOBAL-NEXT: store i32 [[TMP0]], ptr addrspace(4) [[OUT:%.*]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(1) [[ARRAYIDX_I_I]], align 4, !tbaa [[TBAA6:![0-9]+]] +// CHECK-GLOBAL-NEXT: store i32 [[TMP0]], ptr addrspace(4) [[OUT:%.*]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I]], !llvm.loop [[LOOP11:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS1iiNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS9_9naive_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT2_EEvE4typeESP_SN_RSO_SQ_.exit: // CHECK-GLOBAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR7]] @@ -78,8 +78,8 @@ namespace blocked { // CHECK-LOCAL-NEXT: [[CALL_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR7:[0-9]+]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I:%.*]] = sext i32 [[CALL_I_I_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P:%.*]], i64 [[IDXPROM_I_I]] -// CHECK-LOCAL-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I]], align 4, !tbaa [[TBAA7:![0-9]+]] -// CHECK-LOCAL-NEXT: store i32 [[TMP0]], ptr addrspace(4) [[OUT:%.*]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I]], align 4, !tbaa [[TBAA6:![0-9]+]] +// CHECK-LOCAL-NEXT: store i32 [[TMP0]], ptr addrspace(4) [[OUT:%.*]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I]], !llvm.loop [[LOOP11:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS3iiNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS9_9naive_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT2_EEvE4typeESP_SN_RSO_SQ_.exit: // CHECK-LOCAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR6]] @@ -90,17 +90,18 @@ SYCL_EXTERNAL void test_naive(sycl::sub_group &sg, plain_ptr p, int &out) { group_load(sg, p, out, naive_blocked{}); } +// // CHECK-GLOBAL-LABEL: @_ZN7blocked14test_optimizedERN4sycl3_V19sub_groupEPU3AS1iRi( // CHECK-GLOBAL-NEXT: entry: -// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], null +// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(1)) // CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) -// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef i32 @_Z30__spirv_SubgroupBlockReadINTELIjET_PU3AS1Kj(ptr addrspace(1) noundef nonnull [[P]]) #[[ATTR7]] +// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef i32 @_Z30__spirv_SubgroupBlockReadINTELIjET_PU3AS1Kj(ptr addrspace(1) noundef [[P]]) #[[ATTR7]] // CHECK-GLOBAL-NEXT: store i32 [[CALL6_I_I]], ptr addrspace(4) [[OUT:%.*]], align 4 // CHECK-GLOBAL-NEXT: ret void // // CHECK-LOCAL-LABEL: @_ZN7blocked14test_optimizedERN4sycl3_V19sub_groupEPU3AS3iRi( // CHECK-LOCAL-NEXT: entry: -// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], null +// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(3)) // CHECK-LOCAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-LOCAL-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(3) [[P]] to i64 // CHECK-LOCAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP0]], 15 @@ -116,14 +117,14 @@ SYCL_EXTERNAL void test_naive(sycl::sub_group &sg, plain_ptr p, int &out) { // CHECK-LOCAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR7]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[CALL_I_I_I_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P]], i64 [[IDXPROM_I_I_I]] -// CHECK-LOCAL-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] -// CHECK-LOCAL-NEXT: store i32 [[TMP1]], ptr addrspace(4) [[OUT:%.*]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] +// CHECK-LOCAL-NEXT: store i32 [[TMP1]], ptr addrspace(4) [[OUT:%.*]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP13:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental6detail15group_load_implINS0_9sub_groupEPU3AS3iiLm1ENS3_10propertiesINS4_20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS4_9naive_keyEJEEENSB_INS4_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeEST_SR_NS0_4spanISS_XT2_EEESU_.exit.i.i: // CHECK-LOCAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR6]] // CHECK-LOCAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS3IINS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESR_SP_RSQ_SS__EXIT:%.*]] // CHECK-LOCAL: if.end.i.i: -// CHECK-LOCAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef i32 @_Z30__spirv_SubgroupBlockReadINTELIjET_PU3AS3Kj(ptr addrspace(3) noundef nonnull [[P]]) #[[ATTR6]] +// CHECK-LOCAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef i32 @_Z30__spirv_SubgroupBlockReadINTELIjET_PU3AS3Kj(ptr addrspace(3) noundef [[P]]) #[[ATTR6]] // CHECK-LOCAL-NEXT: store i32 [[CALL6_I_I]], ptr addrspace(4) [[OUT]], align 4 // CHECK-LOCAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS3IINS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESR_SP_RSQ_SS__EXIT]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS3iiNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT2_EEvE4typeESR_SP_RSQ_SS_.exit: @@ -137,15 +138,15 @@ SYCL_EXTERNAL void test_optimized(sycl::sub_group &sg, plain_ptr p, // CHECK-GLOBAL-LABEL: @_ZN7blocked27test_contiguous_auto_detectERN4sycl3_V19sub_groupEPU3AS1iRi( // CHECK-GLOBAL-NEXT: entry: -// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], null +// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(1)) // CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) -// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef i32 @_Z30__spirv_SubgroupBlockReadINTELIjET_PU3AS1Kj(ptr addrspace(1) noundef nonnull [[P]]) #[[ATTR7]] +// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef i32 @_Z30__spirv_SubgroupBlockReadINTELIjET_PU3AS1Kj(ptr addrspace(1) noundef [[P]]) #[[ATTR7]] // CHECK-GLOBAL-NEXT: store i32 [[CALL6_I_I]], ptr addrspace(4) [[OUT:%.*]], align 4 // CHECK-GLOBAL-NEXT: ret void // // CHECK-LOCAL-LABEL: @_ZN7blocked27test_contiguous_auto_detectERN4sycl3_V19sub_groupEPU3AS3iRi( // CHECK-LOCAL-NEXT: entry: -// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], null +// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(3)) // CHECK-LOCAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-LOCAL-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(3) [[P]] to i64 // CHECK-LOCAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP0]], 15 @@ -161,14 +162,14 @@ SYCL_EXTERNAL void test_optimized(sycl::sub_group &sg, plain_ptr p, // CHECK-LOCAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR7]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[CALL_I_I_I_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P]], i64 [[IDXPROM_I_I_I]] -// CHECK-LOCAL-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] -// CHECK-LOCAL-NEXT: store i32 [[TMP1]], ptr addrspace(4) [[OUT:%.*]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] +// CHECK-LOCAL-NEXT: store i32 [[TMP1]], ptr addrspace(4) [[OUT:%.*]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP14:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental6detail15group_load_implINS0_9sub_groupEPU3AS3iiLm1ENS3_10propertiesINS4_20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_14full_group_keyEJEEENSB_INS4_9naive_keyEJEEENSB_INS4_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit.i.i: // CHECK-LOCAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR6]] // CHECK-LOCAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS3IINS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESP_SN_RSO_SQ__EXIT:%.*]] // CHECK-LOCAL: if.end.i.i: -// CHECK-LOCAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef i32 @_Z30__spirv_SubgroupBlockReadINTELIjET_PU3AS3Kj(ptr addrspace(3) noundef nonnull [[P]]) #[[ATTR6]] +// CHECK-LOCAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef i32 @_Z30__spirv_SubgroupBlockReadINTELIjET_PU3AS3Kj(ptr addrspace(3) noundef [[P]]) #[[ATTR6]] // CHECK-LOCAL-NEXT: store i32 [[CALL6_I_I]], ptr addrspace(4) [[OUT]], align 4 // CHECK-LOCAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS3IINS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESP_SN_RSO_SQ__EXIT]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS3iiNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT2_EEvE4typeESP_SN_RSO_SQ_.exit: @@ -203,8 +204,8 @@ using accessor_iter_t = local_accessor::iterator; // CHECK-GLOBAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR8]] // CHECK-GLOBAL-NEXT: [[CONV3_I_I_I:%.*]] = sext i32 [[CALL_I_I_I_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ADD_PTR_I_I_I_I_I:%.*]] = getelementptr i32, ptr addrspace(4) [[TMP0]], i64 [[CONV3_I_I_I]] -// CHECK-GLOBAL-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(4) [[ADD_PTR_I_I_I_I_I]], align 4, !tbaa [[TBAA7]] -// CHECK-GLOBAL-NEXT: store i32 [[TMP1]], ptr addrspace(4) [[OUT:%.*]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(4) [[ADD_PTR_I_I_I_I_I]], align 4, !tbaa [[TBAA6]] +// CHECK-GLOBAL-NEXT: store i32 [[TMP1]], ptr addrspace(4) [[OUT:%.*]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP18:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupENS0_6detail17accessor_iteratorIKiLi1EEEiNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSD_INS3_14full_group_keyEJEEENSD_INSB_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT2_EEvE4typeESR_SP_RSQ_SS_.exit: // CHECK-GLOBAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR7]] @@ -276,7 +277,7 @@ SYCL_EXTERNAL void test_accessor_iter_force_optimized(sycl::sub_group &sg, // CHECK-GLOBAL-LABEL: @_ZN7blocked24test_runtime_align_checkERN4sycl3_V19sub_groupEPU3AS1cRc( // CHECK-GLOBAL-NEXT: entry: -// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], null +// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(1)) // CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-GLOBAL-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(1) [[P]] to i64 // CHECK-GLOBAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP0]], 3 @@ -299,7 +300,7 @@ SYCL_EXTERNAL void test_accessor_iter_force_optimized(sycl::sub_group &sg, // CHECK-GLOBAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR7]] // CHECK-GLOBAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS1CCNS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESR_SP_RSQ_SS__EXIT:%.*]] // CHECK-GLOBAL: if.end.i.i: -// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef zeroext i8 @_Z30__spirv_SubgroupBlockReadINTELIhET_PU3AS1Kh(ptr addrspace(1) noundef nonnull [[P]]) #[[ATTR7]] +// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef zeroext i8 @_Z30__spirv_SubgroupBlockReadINTELIhET_PU3AS1Kh(ptr addrspace(1) noundef [[P]]) #[[ATTR7]] // CHECK-GLOBAL-NEXT: store i8 [[CALL6_I_I]], ptr addrspace(4) [[OUT]], align 1 // CHECK-GLOBAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS1CCNS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESR_SP_RSQ_SS__EXIT]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS1ccNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT2_EEvE4typeESR_SP_RSQ_SS_.exit: @@ -307,7 +308,7 @@ SYCL_EXTERNAL void test_accessor_iter_force_optimized(sycl::sub_group &sg, // // CHECK-LOCAL-LABEL: @_ZN7blocked24test_runtime_align_checkERN4sycl3_V19sub_groupEPU3AS3cRc( // CHECK-LOCAL-NEXT: entry: -// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], null +// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(3)) // CHECK-LOCAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-LOCAL-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(3) [[P]] to i64 // CHECK-LOCAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP0]], 15 @@ -330,7 +331,7 @@ SYCL_EXTERNAL void test_accessor_iter_force_optimized(sycl::sub_group &sg, // CHECK-LOCAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR6]] // CHECK-LOCAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS3CCNS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESR_SP_RSQ_SS__EXIT:%.*]] // CHECK-LOCAL: if.end.i.i: -// CHECK-LOCAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef zeroext i8 @_Z30__spirv_SubgroupBlockReadINTELIhET_PU3AS3Kh(ptr addrspace(3) noundef nonnull [[P]]) #[[ATTR6]] +// CHECK-LOCAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef zeroext i8 @_Z30__spirv_SubgroupBlockReadINTELIhET_PU3AS3Kh(ptr addrspace(3) noundef [[P]]) #[[ATTR6]] // CHECK-LOCAL-NEXT: store i8 [[CALL6_I_I]], ptr addrspace(4) [[OUT]], align 1 // CHECK-LOCAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS3CCNS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESR_SP_RSQ_SS__EXIT]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS3ccNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT2_EEvE4typeESR_SP_RSQ_SS_.exit: @@ -347,7 +348,7 @@ SYCL_EXTERNAL void test_runtime_align_check(sycl::sub_group &sg, // CHECK-GLOBAL-NEXT: entry: // CHECK-GLOBAL-NEXT: [[TMP0:%.*]] = load i64, ptr [[OUT:%.*]], align 8, !tbaa [[TBAA27:![0-9]+]] // CHECK-GLOBAL-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4) -// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], null +// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(1)) // CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = ptrtoint ptr addrspace(1) [[P]] to i64 // CHECK-GLOBAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP2]], 3 @@ -376,7 +377,7 @@ SYCL_EXTERNAL void test_runtime_align_check(sycl::sub_group &sg, // CHECK-GLOBAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR7]] // CHECK-GLOBAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS1SSLM4ENS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT3_EEVE4TYPEESR_SP_NS0_4SPANISQ_XT2_EEESS__EXIT:%.*]] // CHECK-GLOBAL: if.end.i.i: -// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef i64 @_Z30__spirv_SubgroupBlockReadINTELImET_PU3AS1Km(ptr addrspace(1) noundef nonnull [[P]]) #[[ATTR7]] +// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef i64 @_Z30__spirv_SubgroupBlockReadINTELImET_PU3AS1Km(ptr addrspace(1) noundef [[P]]) #[[ATTR7]] // CHECK-GLOBAL-NEXT: store i64 [[CALL6_I_I]], ptr addrspace(4) [[TMP1]], align 2 // CHECK-GLOBAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS1SSLM4ENS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT3_EEVE4TYPEESR_SP_NS0_4SPANISQ_XT2_EEESS__EXIT]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS1ssLm4ENS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit: @@ -386,7 +387,7 @@ SYCL_EXTERNAL void test_runtime_align_check(sycl::sub_group &sg, // CHECK-LOCAL-NEXT: entry: // CHECK-LOCAL-NEXT: [[TMP0:%.*]] = load i64, ptr [[OUT:%.*]], align 8, !tbaa [[TBAA30:![0-9]+]] // CHECK-LOCAL-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4) -// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], null +// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(3)) // CHECK-LOCAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-LOCAL-NEXT: [[TMP2:%.*]] = ptrtoint ptr addrspace(3) [[P]] to i64 // CHECK-LOCAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP2]], 15 @@ -415,7 +416,7 @@ SYCL_EXTERNAL void test_runtime_align_check(sycl::sub_group &sg, // CHECK-LOCAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR6]] // CHECK-LOCAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS3SSLM4ENS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT3_EEVE4TYPEESR_SP_NS0_4SPANISQ_XT2_EEESS__EXIT:%.*]] // CHECK-LOCAL: if.end.i.i: -// CHECK-LOCAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef i64 @_Z30__spirv_SubgroupBlockReadINTELImET_PU3AS3Km(ptr addrspace(3) noundef nonnull [[P]]) #[[ATTR6]] +// CHECK-LOCAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef i64 @_Z30__spirv_SubgroupBlockReadINTELImET_PU3AS3Km(ptr addrspace(3) noundef [[P]]) #[[ATTR6]] // CHECK-LOCAL-NEXT: store i64 [[CALL6_I_I]], ptr addrspace(4) [[TMP1]], align 2 // CHECK-LOCAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS3SSLM4ENS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT3_EEVE4TYPEESR_SP_NS0_4SPANISQ_XT2_EEESS__EXIT]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS3ssLm4ENS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit: @@ -445,9 +446,9 @@ SYCL_EXTERNAL void test_four_shorts(sycl::sub_group &sg, plain_ptr p, // CHECK-GLOBAL-NEXT: [[ADD_I_I_I_I:%.*]] = add i32 [[MUL_I_I_I_I]], [[I_0_I_I_I]] // CHECK-GLOBAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[P:%.*]], i64 [[IDXPROM_I_I_I]] -// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP32:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS1iiLm3ENS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit: @@ -471,9 +472,9 @@ SYCL_EXTERNAL void test_four_shorts(sycl::sub_group &sg, plain_ptr p, // CHECK-LOCAL-NEXT: [[ADD_I_I_I_I:%.*]] = add i32 [[MUL_I_I_I_I]], [[I_0_I_I_I]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P:%.*]], i64 [[IDXPROM_I_I_I]] -// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP35:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS3iiLm3ENS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit: @@ -503,9 +504,9 @@ SYCL_EXTERNAL void test_non_power_of_two(sycl::sub_group &sg, plain_ptr p, // CHECK-GLOBAL-NEXT: [[ADD_I_I_I_I:%.*]] = or disjoint i32 [[MUL_I_I_I_I]], [[I_0_I_I_I]] // CHECK-GLOBAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[P:%.*]], i64 [[IDXPROM_I_I_I]] -// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP33:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS1iiLm4ENS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit: @@ -529,9 +530,9 @@ SYCL_EXTERNAL void test_non_power_of_two(sycl::sub_group &sg, plain_ptr p, // CHECK-LOCAL-NEXT: [[ADD_I_I_I_I:%.*]] = or disjoint i32 [[MUL_I_I_I_I]], [[I_0_I_I_I]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P:%.*]], i64 [[IDXPROM_I_I_I]] -// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP36:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS3iiLm4ENS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit: @@ -562,9 +563,9 @@ SYCL_EXTERNAL void test_four_ints(sycl::sub_group &sg, plain_ptr p, // CHECK-GLOBAL-NEXT: [[ADD_I_I_I_I:%.*]] = add i32 [[MUL_I_I_I_I]], [[I_0_I_I_I]] // CHECK-GLOBAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[P:%.*]], i64 [[IDXPROM_I_I_I]] -// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP34:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS1iiLm7ENS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit: @@ -588,9 +589,9 @@ SYCL_EXTERNAL void test_four_ints(sycl::sub_group &sg, plain_ptr p, // CHECK-LOCAL-NEXT: [[ADD_I_I_I_I:%.*]] = add i32 [[MUL_I_I_I_I]], [[I_0_I_I_I]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P:%.*]], i64 [[IDXPROM_I_I_I]] -// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP37:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS3iiLm7ENS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit: @@ -626,9 +627,9 @@ namespace striped { // CHECK-GLOBAL-NEXT: [[ADD_I_I_I:%.*]] = add i32 [[CALL_I_I_I_I_I]], [[MUL_I_I_I]] // CHECK-GLOBAL-NEXT: [[IDXPROM_I_I:%.*]] = sext i32 [[ADD_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[P:%.*]], i64 [[IDXPROM_I_I]] -// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(1) [[ARRAYIDX_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(1) [[ARRAYIDX_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I]] -// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[INC_I_I]] = add nuw nsw i32 [[I_0_I_I]], 1 // CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I]], !llvm.loop [[LOOP35:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS1iiLm2ENS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS9_9naive_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESP_SN_NS0_4spanISO_XT2_EEESQ_.exit: @@ -653,9 +654,9 @@ namespace striped { // CHECK-LOCAL-NEXT: [[ADD_I_I_I:%.*]] = add i32 [[CALL_I_I_I_I_I]], [[MUL_I_I_I]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I:%.*]] = sext i32 [[ADD_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P:%.*]], i64 [[IDXPROM_I_I]] -// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I]] -// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[INC_I_I]] = add nuw nsw i32 [[I_0_I_I]], 1 // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I]], !llvm.loop [[LOOP38:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS3iiLm2ENS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS9_9naive_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESP_SN_NS0_4spanISO_XT2_EEESQ_.exit: @@ -672,9 +673,9 @@ SYCL_EXTERNAL void test_naive(sycl::sub_group &sg, plain_ptr p, // CHECK-GLOBAL-NEXT: entry: // CHECK-GLOBAL-NEXT: [[TMP0:%.*]] = load i64, ptr [[OUT:%.*]], align 8, !tbaa [[TBAA13]] // CHECK-GLOBAL-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4) -// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], null +// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(1)) // CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) -// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef <2 x i32> @_Z30__spirv_SubgroupBlockReadINTELIDv2_jET_PU3AS1Kj(ptr addrspace(1) noundef nonnull [[P]]) #[[ATTR7]] +// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef <2 x i32> @_Z30__spirv_SubgroupBlockReadINTELIDv2_jET_PU3AS1Kj(ptr addrspace(1) noundef [[P]]) #[[ATTR7]] // CHECK-GLOBAL-NEXT: store <2 x i32> [[CALL6_I_I]], ptr addrspace(4) [[TMP1]], align 4 // CHECK-GLOBAL-NEXT: ret void // @@ -682,7 +683,7 @@ SYCL_EXTERNAL void test_naive(sycl::sub_group &sg, plain_ptr p, // CHECK-LOCAL-NEXT: entry: // CHECK-LOCAL-NEXT: [[TMP0:%.*]] = load i64, ptr [[OUT:%.*]], align 8, !tbaa [[TBAA15]] // CHECK-LOCAL-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4) -// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], null +// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(3)) // CHECK-LOCAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-LOCAL-NEXT: [[TMP2:%.*]] = ptrtoint ptr addrspace(3) [[P]] to i64 // CHECK-LOCAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP2]], 15 @@ -703,16 +704,16 @@ SYCL_EXTERNAL void test_naive(sycl::sub_group &sg, plain_ptr p, // CHECK-LOCAL-NEXT: [[ADD_I_I_I_I:%.*]] = add i32 [[CALL_I_I_I_I_I_I]], [[MUL_I_I_I_I]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P]], i64 [[IDXPROM_I_I_I]] -// CHECK-LOCAL-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-LOCAL-NEXT: store i32 [[TMP3]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: store i32 [[TMP3]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP39:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental6detail15group_load_implINS0_9sub_groupEPU3AS3iiLm2ENS3_10propertiesINS4_20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS4_9naive_keyEJEEENSB_INS4_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeEST_SR_NS0_4spanISS_XT2_EEESU_.exit.i.i: // CHECK-LOCAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR6]] // CHECK-LOCAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS3IILM2ENS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI1EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT3_EEVE4TYPEESR_SP_NS0_4SPANISQ_XT2_EEESS__EXIT:%.*]] // CHECK-LOCAL: if.end.i.i: -// CHECK-LOCAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef <2 x i32> @_Z30__spirv_SubgroupBlockReadINTELIDv2_jET_PU3AS3Kj(ptr addrspace(3) noundef nonnull [[P]]) #[[ATTR6]] +// CHECK-LOCAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef <2 x i32> @_Z30__spirv_SubgroupBlockReadINTELIDv2_jET_PU3AS3Kj(ptr addrspace(3) noundef [[P]]) #[[ATTR6]] // CHECK-LOCAL-NEXT: store <2 x i32> [[CALL6_I_I]], ptr addrspace(4) [[TMP1]], align 4 // CHECK-LOCAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS3IILM2ENS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI1EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT3_EEVE4TYPEESR_SP_NS0_4SPANISQ_XT2_EEESS__EXIT]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS3iiLm2ENS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit: @@ -726,15 +727,15 @@ SYCL_EXTERNAL void test_optimized(sycl::sub_group &sg, plain_ptr p, // CHECK-GLOBAL-LABEL: @_ZN7striped27test_contiguous_auto_detectERN4sycl3_V19sub_groupEPU3AS1iRi( // CHECK-GLOBAL-NEXT: entry: -// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], null +// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(1)) // CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) -// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef i32 @_Z30__spirv_SubgroupBlockReadINTELIjET_PU3AS1Kj(ptr addrspace(1) noundef nonnull [[P]]) #[[ATTR7]] +// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef i32 @_Z30__spirv_SubgroupBlockReadINTELIjET_PU3AS1Kj(ptr addrspace(1) noundef [[P]]) #[[ATTR7]] // CHECK-GLOBAL-NEXT: store i32 [[CALL6_I_I]], ptr addrspace(4) [[OUT:%.*]], align 4 // CHECK-GLOBAL-NEXT: ret void // // CHECK-LOCAL-LABEL: @_ZN7striped27test_contiguous_auto_detectERN4sycl3_V19sub_groupEPU3AS3iRi( // CHECK-LOCAL-NEXT: entry: -// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], null +// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(3)) // CHECK-LOCAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-LOCAL-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(3) [[P]] to i64 // CHECK-LOCAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP0]], 15 @@ -750,14 +751,14 @@ SYCL_EXTERNAL void test_optimized(sycl::sub_group &sg, plain_ptr p, // CHECK-LOCAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR7]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[CALL_I_I_I_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P]], i64 [[IDXPROM_I_I_I]] -// CHECK-LOCAL-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] -// CHECK-LOCAL-NEXT: store i32 [[TMP1]], ptr addrspace(4) [[OUT:%.*]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] +// CHECK-LOCAL-NEXT: store i32 [[TMP1]], ptr addrspace(4) [[OUT:%.*]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP40:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental6detail15group_load_implINS0_9sub_groupEPU3AS3iiLm1ENS3_10propertiesINS4_20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS3_14full_group_keyEJEEENSB_INS4_9naive_keyEJEEENSB_INS4_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit.i.i: // CHECK-LOCAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR6]] // CHECK-LOCAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS3IINS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI1EEEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESP_SN_RSO_SQ__EXIT:%.*]] // CHECK-LOCAL: if.end.i.i: -// CHECK-LOCAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef i32 @_Z30__spirv_SubgroupBlockReadINTELIjET_PU3AS3Kj(ptr addrspace(3) noundef nonnull [[P]]) #[[ATTR6]] +// CHECK-LOCAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef i32 @_Z30__spirv_SubgroupBlockReadINTELIjET_PU3AS3Kj(ptr addrspace(3) noundef [[P]]) #[[ATTR6]] // CHECK-LOCAL-NEXT: store i32 [[CALL6_I_I]], ptr addrspace(4) [[OUT]], align 4 // CHECK-LOCAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS3IINS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI1EEEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESP_SN_RSO_SQ__EXIT]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS3iiNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT2_EEvE4typeESP_SN_RSO_SQ_.exit: @@ -798,9 +799,9 @@ using accessor_iter_t = local_accessor::iterator; // CHECK-GLOBAL-NEXT: [[ADD_I_I_I_I:%.*]] = add i32 [[CALL_I_I_I_I_I_I]], [[MUL_I_I_I_I]] // CHECK-GLOBAL-NEXT: [[CONV3_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ADD_PTR_I_I_I_I_I:%.*]] = getelementptr i32, ptr addrspace(4) [[TMP2]], i64 [[CONV3_I_I_I]] -// CHECK-GLOBAL-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(4) [[ADD_PTR_I_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(4) [[ADD_PTR_I_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-GLOBAL-NEXT: store i32 [[TMP3]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: store i32 [[TMP3]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP36:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupENS0_6detail17accessor_iteratorIKiLi1EEEiLm2ENS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSD_INS3_14full_group_keyEJEEENSD_INSB_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit: @@ -878,7 +879,7 @@ SYCL_EXTERNAL void test_accessor_iter_force_optimized(sycl::sub_group &sg, // CHECK-GLOBAL-NEXT: entry: // CHECK-GLOBAL-NEXT: [[TMP0:%.*]] = load i64, ptr [[OUT:%.*]], align 8, !tbaa [[TBAA43:![0-9]+]] // CHECK-GLOBAL-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4) -// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], null +// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(1)) // CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = ptrtoint ptr addrspace(1) [[P]] to i64 // CHECK-GLOBAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP2]], 3 @@ -908,7 +909,7 @@ SYCL_EXTERNAL void test_accessor_iter_force_optimized(sycl::sub_group &sg, // CHECK-GLOBAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR7]] // CHECK-GLOBAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS1CCLM2ENS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI1EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT3_EEVE4TYPEESR_SP_NS0_4SPANISQ_XT2_EEESS__EXIT:%.*]] // CHECK-GLOBAL: if.end.i.i: -// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef <2 x i8> @_Z30__spirv_SubgroupBlockReadINTELIDv2_hET_PU3AS1Kh(ptr addrspace(1) noundef nonnull [[P]]) #[[ATTR7]] +// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef <2 x i8> @_Z30__spirv_SubgroupBlockReadINTELIDv2_hET_PU3AS1Kh(ptr addrspace(1) noundef [[P]]) #[[ATTR7]] // CHECK-GLOBAL-NEXT: store <2 x i8> [[CALL6_I_I]], ptr addrspace(4) [[TMP1]], align 1 // CHECK-GLOBAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS1CCLM2ENS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI1EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT3_EEVE4TYPEESR_SP_NS0_4SPANISQ_XT2_EEESS__EXIT]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS1ccLm2ENS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit: @@ -918,7 +919,7 @@ SYCL_EXTERNAL void test_accessor_iter_force_optimized(sycl::sub_group &sg, // CHECK-LOCAL-NEXT: entry: // CHECK-LOCAL-NEXT: [[TMP0:%.*]] = load i64, ptr [[OUT:%.*]], align 8, !tbaa [[TBAA51:![0-9]+]] // CHECK-LOCAL-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4) -// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], null +// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(3)) // CHECK-LOCAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-LOCAL-NEXT: [[TMP2:%.*]] = ptrtoint ptr addrspace(3) [[P]] to i64 // CHECK-LOCAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP2]], 15 @@ -948,7 +949,7 @@ SYCL_EXTERNAL void test_accessor_iter_force_optimized(sycl::sub_group &sg, // CHECK-LOCAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR6]] // CHECK-LOCAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS3CCLM2ENS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI1EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT3_EEVE4TYPEESR_SP_NS0_4SPANISQ_XT2_EEESS__EXIT:%.*]] // CHECK-LOCAL: if.end.i.i: -// CHECK-LOCAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef <2 x i8> @_Z30__spirv_SubgroupBlockReadINTELIDv2_hET_PU3AS3Kh(ptr addrspace(3) noundef nonnull [[P]]) #[[ATTR6]] +// CHECK-LOCAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef <2 x i8> @_Z30__spirv_SubgroupBlockReadINTELIDv2_hET_PU3AS3Kh(ptr addrspace(3) noundef [[P]]) #[[ATTR6]] // CHECK-LOCAL-NEXT: store <2 x i8> [[CALL6_I_I]], ptr addrspace(4) [[TMP1]], align 1 // CHECK-LOCAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS3CCLM2ENS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI1EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT3_EEVE4TYPEESR_SP_NS0_4SPANISQ_XT2_EEESS__EXIT]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS3ccLm2ENS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit: @@ -966,7 +967,7 @@ SYCL_EXTERNAL void test_runtime_align_check(sycl::sub_group &sg, // CHECK-GLOBAL-NEXT: entry: // CHECK-GLOBAL-NEXT: [[TMP0:%.*]] = load i64, ptr [[OUT:%.*]], align 8, !tbaa [[TBAA27]] // CHECK-GLOBAL-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4) -// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], null +// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(1)) // CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = ptrtoint ptr addrspace(1) [[P]] to i64 // CHECK-GLOBAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP2]], 3 @@ -996,7 +997,7 @@ SYCL_EXTERNAL void test_runtime_align_check(sycl::sub_group &sg, // CHECK-GLOBAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR7]] // CHECK-GLOBAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS1SSLM4ENS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI1EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT3_EEVE4TYPEESR_SP_NS0_4SPANISQ_XT2_EEESS__EXIT:%.*]] // CHECK-GLOBAL: if.end.i.i: -// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef <4 x i16> @_Z30__spirv_SubgroupBlockReadINTELIDv4_tET_PU3AS1Kt(ptr addrspace(1) noundef nonnull [[P]]) #[[ATTR7]] +// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef <4 x i16> @_Z30__spirv_SubgroupBlockReadINTELIDv4_tET_PU3AS1Kt(ptr addrspace(1) noundef [[P]]) #[[ATTR7]] // CHECK-GLOBAL-NEXT: store <4 x i16> [[CALL6_I_I]], ptr addrspace(4) [[TMP1]], align 2 // CHECK-GLOBAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS1SSLM4ENS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI1EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT3_EEVE4TYPEESR_SP_NS0_4SPANISQ_XT2_EEESS__EXIT]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS1ssLm4ENS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit: @@ -1006,7 +1007,7 @@ SYCL_EXTERNAL void test_runtime_align_check(sycl::sub_group &sg, // CHECK-LOCAL-NEXT: entry: // CHECK-LOCAL-NEXT: [[TMP0:%.*]] = load i64, ptr [[OUT:%.*]], align 8, !tbaa [[TBAA30]] // CHECK-LOCAL-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4) -// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], null +// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(3)) // CHECK-LOCAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-LOCAL-NEXT: [[TMP2:%.*]] = ptrtoint ptr addrspace(3) [[P]] to i64 // CHECK-LOCAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP2]], 15 @@ -1036,7 +1037,7 @@ SYCL_EXTERNAL void test_runtime_align_check(sycl::sub_group &sg, // CHECK-LOCAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR6]] // CHECK-LOCAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS3SSLM4ENS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI1EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT3_EEVE4TYPEESR_SP_NS0_4SPANISQ_XT2_EEESS__EXIT:%.*]] // CHECK-LOCAL: if.end.i.i: -// CHECK-LOCAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef <4 x i16> @_Z30__spirv_SubgroupBlockReadINTELIDv4_tET_PU3AS3Kt(ptr addrspace(3) noundef nonnull [[P]]) #[[ATTR6]] +// CHECK-LOCAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef <4 x i16> @_Z30__spirv_SubgroupBlockReadINTELIDv4_tET_PU3AS3Kt(ptr addrspace(3) noundef [[P]]) #[[ATTR6]] // CHECK-LOCAL-NEXT: store <4 x i16> [[CALL6_I_I]], ptr addrspace(4) [[TMP1]], align 2 // CHECK-LOCAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS3SSLM4ENS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI1EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT3_EEVE4TYPEESR_SP_NS0_4SPANISQ_XT2_EEESS__EXIT]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS3ssLm4ENS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit: @@ -1053,7 +1054,7 @@ SYCL_EXTERNAL void test_four_shorts(sycl::sub_group &sg, plain_ptr p, // CHECK-GLOBAL-NEXT: entry: // CHECK-GLOBAL-NEXT: [[TMP0:%.*]] = load i64, ptr [[OUT:%.*]], align 8, !tbaa [[TBAA27]] // CHECK-GLOBAL-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4) -// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], null +// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(1)) // CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = ptrtoint ptr addrspace(1) [[P]] to i64 // CHECK-GLOBAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP2]], 3 @@ -1083,7 +1084,7 @@ SYCL_EXTERNAL void test_four_shorts(sycl::sub_group &sg, plain_ptr p, // CHECK-GLOBAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR7]] // CHECK-GLOBAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS1SSLM16ENS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI1EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT3_EEVE4TYPEESR_SP_NS0_4SPANISQ_XT2_EEESS__EXIT:%.*]] // CHECK-GLOBAL: if.end.i.i: -// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef <16 x i16> @_Z30__spirv_SubgroupBlockReadINTELIDv16_tET_PU3AS1Kt(ptr addrspace(1) noundef nonnull [[P]]) #[[ATTR7]] +// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef <16 x i16> @_Z30__spirv_SubgroupBlockReadINTELIDv16_tET_PU3AS1Kt(ptr addrspace(1) noundef [[P]]) #[[ATTR7]] // CHECK-GLOBAL-NEXT: store <16 x i16> [[CALL6_I_I]], ptr addrspace(4) [[TMP1]], align 2 // CHECK-GLOBAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS1SSLM16ENS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI1EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT3_EEVE4TYPEESR_SP_NS0_4SPANISQ_XT2_EEESS__EXIT]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS1ssLm16ENS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit: @@ -1093,7 +1094,7 @@ SYCL_EXTERNAL void test_four_shorts(sycl::sub_group &sg, plain_ptr p, // CHECK-LOCAL-NEXT: entry: // CHECK-LOCAL-NEXT: [[TMP0:%.*]] = load i64, ptr [[OUT:%.*]], align 8, !tbaa [[TBAA30]] // CHECK-LOCAL-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4) -// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], null +// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(3)) // CHECK-LOCAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-LOCAL-NEXT: [[TMP2:%.*]] = ptrtoint ptr addrspace(3) [[P]] to i64 // CHECK-LOCAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP2]], 15 @@ -1123,7 +1124,7 @@ SYCL_EXTERNAL void test_four_shorts(sycl::sub_group &sg, plain_ptr p, // CHECK-LOCAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR6]] // CHECK-LOCAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS3SSLM16ENS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI1EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT3_EEVE4TYPEESR_SP_NS0_4SPANISQ_XT2_EEESS__EXIT:%.*]] // CHECK-LOCAL: if.end.i.i: -// CHECK-LOCAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef <16 x i16> @_Z30__spirv_SubgroupBlockReadINTELIDv16_tET_PU3AS3Kt(ptr addrspace(3) noundef nonnull [[P]]) #[[ATTR6]] +// CHECK-LOCAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef <16 x i16> @_Z30__spirv_SubgroupBlockReadINTELIDv16_tET_PU3AS3Kt(ptr addrspace(3) noundef [[P]]) #[[ATTR6]] // CHECK-LOCAL-NEXT: store <16 x i16> [[CALL6_I_I]], ptr addrspace(4) [[TMP1]], align 2 // CHECK-LOCAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS3SSLM16ENS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI1EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT3_EEVE4TYPEESR_SP_NS0_4SPANISQ_XT2_EEESS__EXIT]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS3ssLm16ENS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit: @@ -1152,9 +1153,9 @@ SYCL_EXTERNAL void test_sixteen_shorts(sycl::sub_group &sg, plain_ptr p, // CHECK-GLOBAL-NEXT: [[ADD_I_I_I_I:%.*]] = add i32 [[CALL_I_I_I_I_I_I]], [[MUL_I_I_I_I]] // CHECK-GLOBAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[P:%.*]], i64 [[IDXPROM_I_I_I]] -// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP48:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS1iiLm3ENS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit: @@ -1179,9 +1180,9 @@ SYCL_EXTERNAL void test_sixteen_shorts(sycl::sub_group &sg, plain_ptr p, // CHECK-LOCAL-NEXT: [[ADD_I_I_I_I:%.*]] = add i32 [[CALL_I_I_I_I_I_I]], [[MUL_I_I_I_I]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P:%.*]], i64 [[IDXPROM_I_I_I]] -// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP56:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS3iiLm3ENS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit: @@ -1212,9 +1213,9 @@ SYCL_EXTERNAL void test_non_power_of_two(sycl::sub_group &sg, plain_ptr p, // CHECK-GLOBAL-NEXT: [[ADD_I_I_I_I:%.*]] = add i32 [[CALL_I_I_I_I_I_I]], [[MUL_I_I_I_I]] // CHECK-GLOBAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[P:%.*]], i64 [[IDXPROM_I_I_I]] -// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP49:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS1iiLm16ENS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit: @@ -1239,9 +1240,9 @@ SYCL_EXTERNAL void test_non_power_of_two(sycl::sub_group &sg, plain_ptr p, // CHECK-LOCAL-NEXT: [[ADD_I_I_I_I:%.*]] = add i32 [[CALL_I_I_I_I_I_I]], [[MUL_I_I_I_I]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P:%.*]], i64 [[IDXPROM_I_I_I]] -// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP57:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS3iiLm16ENS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit: @@ -1273,9 +1274,9 @@ SYCL_EXTERNAL void test_sixteen_ints(sycl::sub_group &sg, plain_ptr p, // CHECK-GLOBAL-NEXT: [[ADD_I_I_I_I:%.*]] = add i32 [[CALL_I_I_I_I_I_I]], [[MUL_I_I_I_I]] // CHECK-GLOBAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[P:%.*]], i64 [[IDXPROM_I_I_I]] -// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP50:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS1iiLm11ENS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit: @@ -1300,9 +1301,9 @@ SYCL_EXTERNAL void test_sixteen_ints(sycl::sub_group &sg, plain_ptr p, // CHECK-LOCAL-NEXT: [[ADD_I_I_I_I:%.*]] = add i32 [[CALL_I_I_I_I_I_I]], [[MUL_I_I_I_I]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P:%.*]], i64 [[IDXPROM_I_I_I]] -// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP58:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS3iiLm11ENS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit: diff --git a/sycl/test/check_device_code/group_load_store_alignment.cpp b/sycl/test/check_device_code/group_load_store_alignment.cpp index 76fe4706380de..0943f53401558 100644 --- a/sycl/test/check_device_code/group_load_store_alignment.cpp +++ b/sycl/test/check_device_code/group_load_store_alignment.cpp @@ -22,7 +22,7 @@ using plain_ptr = typename sycl::detail::DecoratedType< // CHECK-GLOBAL-LABEL: @_Z32test_load_without_alignment_hintRN4sycl3_V19sub_groupEPU3AS1sRs( // CHECK-GLOBAL-NEXT: entry: -// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], null +// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(1)) // CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-GLOBAL-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(1) [[P]] to i64 // CHECK-GLOBAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP0]], 3 @@ -38,14 +38,14 @@ using plain_ptr = typename sycl::detail::DecoratedType< // CHECK-GLOBAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR5:[0-9]+]] // CHECK-GLOBAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[CALL_I_I_I_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i16, ptr addrspace(1) [[P]], i64 [[IDXPROM_I_I_I]] -// CHECK-GLOBAL-NEXT: [[TMP1:%.*]] = load i16, ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 2, !tbaa [[TBAA7:![0-9]+]] -// CHECK-GLOBAL-NEXT: store i16 [[TMP1]], ptr addrspace(4) [[OUT:%.*]], align 2, !tbaa [[TBAA7]] -// CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP11:![0-9]+]] +// CHECK-GLOBAL-NEXT: [[TMP1:%.*]] = load i16, ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 2, !tbaa [[TBAA11:![0-9]+]] +// CHECK-GLOBAL-NEXT: store i16 [[TMP1]], ptr addrspace(4) [[OUT:%.*]], align 2, !tbaa [[TBAA11]] +// CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP13:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental6detail15group_load_implINS0_9sub_groupEPU3AS1ssLm1ENS3_10propertiesINS4_20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS4_9naive_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_SP_NS0_4spanISQ_XT2_EEESS_.exit.i.i: // CHECK-GLOBAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR4]] // CHECK-GLOBAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS1SSNS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESP_SN_RSO_SQ__EXIT:%.*]] // CHECK-GLOBAL: if.end.i.i: -// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef zeroext i16 @_Z30__spirv_SubgroupBlockReadINTELItET_PU3AS1Kt(ptr addrspace(1) noundef nonnull [[P]]) #[[ATTR4]] +// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef zeroext i16 @_Z30__spirv_SubgroupBlockReadINTELItET_PU3AS1Kt(ptr addrspace(1) noundef [[P]]) #[[ATTR4]] // CHECK-GLOBAL-NEXT: store i16 [[CALL6_I_I]], ptr addrspace(4) [[OUT]], align 2 // CHECK-GLOBAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS1SSNS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESP_SN_RSO_SQ__EXIT]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS1ssNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT2_EEvE4typeESP_SN_RSO_SQ_.exit: @@ -59,9 +59,9 @@ SYCL_EXTERNAL void test_load_without_alignment_hint(sycl::sub_group &sg, // CHECK-GLOBAL-LABEL: @_Z29test_load_with_alignment_hintRN4sycl3_V19sub_groupEPU3AS1sRs( // CHECK-GLOBAL-NEXT: entry: -// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], null +// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(1)) // CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) -// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef zeroext i16 @_Z30__spirv_SubgroupBlockReadINTELItET_PU3AS1Kt(ptr addrspace(1) noundef nonnull [[P]]) #[[ATTR4]] +// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef zeroext i16 @_Z30__spirv_SubgroupBlockReadINTELItET_PU3AS1Kt(ptr addrspace(1) noundef [[P]]) #[[ATTR4]] // CHECK-GLOBAL-NEXT: store i16 [[CALL6_I_I]], ptr addrspace(4) [[OUT:%.*]], align 2 // CHECK-GLOBAL-NEXT: ret void // @@ -73,7 +73,7 @@ SYCL_EXTERNAL void test_load_with_alignment_hint(sycl::sub_group &sg, // CHECK-GLOBAL-LABEL: @_Z33test_store_without_alignment_hintRN4sycl3_V19sub_groupEiPU3AS1i( // CHECK-GLOBAL-NEXT: entry: -// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], null +// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(1)) // CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-GLOBAL-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(1) [[P]] to i64 // CHECK-GLOBAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP0]], 15 @@ -89,7 +89,7 @@ SYCL_EXTERNAL void test_load_with_alignment_hint(sycl::sub_group &sg, // CHECK-GLOBAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR5]] // CHECK-GLOBAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[CALL_I_I_I_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[P]], i64 [[IDXPROM_I_I_I]] -// CHECK-GLOBAL-NEXT: store i32 [[V:%.*]], ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA13:![0-9]+]] +// CHECK-GLOBAL-NEXT: store i32 [[V:%.*]], ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6:![0-9]+]] // CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP15:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental6detail16group_store_implINS0_9sub_groupEKiLm1EPU3AS1iNS3_10propertiesINS4_20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSC_INS3_21contiguous_memory_keyEJEEENSC_INS3_14full_group_keyEJEEENSC_INS4_9naive_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T2_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESS_NS0_4spanISQ_XT1_EEESR_ST_.exit.i.i: // CHECK-GLOBAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR4]] @@ -98,7 +98,7 @@ SYCL_EXTERNAL void test_load_with_alignment_hint(sycl::sub_group &sg, // CHECK-GLOBAL-NEXT: [[I_0_I_I:%.*]] = phi i1 [ false, [[FOR_COND_I_I]] ], [ true, [[ENTRY:%.*]] ] // CHECK-GLOBAL-NEXT: br i1 [[I_0_I_I]], label [[FOR_COND_I_I]], label [[FOR_COND_CLEANUP_I_I:%.*]], !llvm.loop [[LOOP16:![0-9]+]] // CHECK-GLOBAL: for.cond.cleanup.i.i: -// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS1jj(ptr addrspace(1) noundef nonnull [[P]], i32 noundef [[V]]) #[[ATTR4]] +// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS1jj(ptr addrspace(1) noundef [[P]], i32 noundef [[V]]) #[[ATTR4]] // CHECK-GLOBAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL11GROUP_STOREINS0_9SUB_GROUPEIPU3AS1INS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE18VERIFY_STORE_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESP_RKSN_SO_SQ__EXIT]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiPU3AS1iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT2_EEvE4typeESP_RKSN_SO_SQ_.exit: // CHECK-GLOBAL-NEXT: ret void @@ -110,14 +110,14 @@ SYCL_EXTERNAL void test_store_without_alignment_hint(sycl::sub_group &sg, int v, // CHECK-GLOBAL-LABEL: @_Z30test_store_with_alignment_hintRN4sycl3_V19sub_groupEiPU3AS1i( // CHECK-GLOBAL-NEXT: entry: -// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], null +// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(1)) // CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I:%.*]] // CHECK-GLOBAL: for.cond.i.i: // CHECK-GLOBAL-NEXT: [[I_0_I_I:%.*]] = phi i1 [ true, [[ENTRY:%.*]] ], [ false, [[FOR_COND_I_I]] ] // CHECK-GLOBAL-NEXT: br i1 [[I_0_I_I]], label [[FOR_COND_I_I]], label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL11GROUP_STOREINS0_9SUB_GROUPEIPU3AS1INS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_13ALIGNMENT_KEYEJST17INTEGRAL_CONSTANTIILI16EEEEENSB_INS3_18DATA_PLACEMENT_KEYEJSD_IILI0EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE18VERIFY_STORE_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESS_RKSQ_SR_ST__EXIT:%.*]], !llvm.loop [[LOOP17:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiPU3AS1iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_13alignment_keyEJSt17integral_constantIiLi16EEEEENSB_INS3_18data_placement_keyEJSD_IiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT2_EEvE4typeESS_RKSQ_SR_ST_.exit: -// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS1jj(ptr addrspace(1) noundef nonnull [[P]], i32 noundef [[V:%.*]]) #[[ATTR4]] +// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS1jj(ptr addrspace(1) noundef [[P]], i32 noundef [[V:%.*]]) #[[ATTR4]] // CHECK-GLOBAL-NEXT: ret void // SYCL_EXTERNAL void test_store_with_alignment_hint(sycl::sub_group &sg, int v, diff --git a/sycl/test/check_device_code/group_load_store_native_key.cpp b/sycl/test/check_device_code/group_load_store_native_key.cpp index 62d72255c3596..f15111cbecd47 100644 --- a/sycl/test/check_device_code/group_load_store_native_key.cpp +++ b/sycl/test/check_device_code/group_load_store_native_key.cpp @@ -32,15 +32,15 @@ using plain_ptr = typename sycl::detail::DecoratedType< // CHECK-GLOBAL-LABEL: @_Z9test_loadRN4sycl3_V19sub_groupEPU3AS1iRi( // CHECK-GLOBAL-NEXT: entry: -// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], null +// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(1)) // CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) -// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef i32 @_Z30__spirv_SubgroupBlockReadINTELIjET_PU3AS1Kj(ptr addrspace(1) noundef nonnull [[P]]) #[[ATTR4:[0-9]+]] +// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef i32 @_Z30__spirv_SubgroupBlockReadINTELIjET_PU3AS1Kj(ptr addrspace(1) noundef [[P]]) #[[ATTR4:[0-9]+]] // CHECK-GLOBAL-NEXT: store i32 [[CALL6_I_I]], ptr addrspace(4) [[OUT:%.*]], align 4 // CHECK-GLOBAL-NEXT: ret void // // CHECK-LOCAL-LABEL: @_Z9test_loadRN4sycl3_V19sub_groupEPU3AS3iRi( // CHECK-LOCAL-NEXT: entry: -// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], null +// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(3)) // CHECK-LOCAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-LOCAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR4:[0-9]+]] // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I_I:%.*]] @@ -51,8 +51,8 @@ using plain_ptr = typename sycl::detail::DecoratedType< // CHECK-LOCAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR5:[0-9]+]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[CALL_I_I_I_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P]], i64 [[IDXPROM_I_I_I]] -// CHECK-LOCAL-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7:![0-9]+]] -// CHECK-LOCAL-NEXT: store i32 [[TMP0]], ptr addrspace(4) [[OUT:%.*]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6:![0-9]+]] +// CHECK-LOCAL-NEXT: store i32 [[TMP0]], ptr addrspace(4) [[OUT:%.*]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP11:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS3iiNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT2_EEvE4typeESP_SN_RSO_SQ_.exit: // CHECK-LOCAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR4]] @@ -64,15 +64,15 @@ SYCL_EXTERNAL void test_load(sycl::sub_group &sg, plain_ptr p, int &out) { // CHECK-GLOBAL-LABEL: @_Z16test_load_nativeRN4sycl3_V19sub_groupEPU3AS1iRi( // CHECK-GLOBAL-NEXT: entry: -// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], null +// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(1)) // CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) -// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef i32 @_Z30__spirv_SubgroupBlockReadINTELIjET_PU3AS1Kj(ptr addrspace(1) noundef nonnull [[P]]) #[[ATTR4]] +// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef i32 @_Z30__spirv_SubgroupBlockReadINTELIjET_PU3AS1Kj(ptr addrspace(1) noundef [[P]]) #[[ATTR4]] // CHECK-GLOBAL-NEXT: store i32 [[CALL6_I_I]], ptr addrspace(4) [[OUT:%.*]], align 4 // CHECK-GLOBAL-NEXT: ret void // // CHECK-LOCAL-LABEL: @_Z16test_load_nativeRN4sycl3_V19sub_groupEPU3AS3iRi( // CHECK-LOCAL-NEXT: entry: -// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], null +// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(3)) // CHECK-LOCAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-LOCAL-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(3) [[P]] to i64 // CHECK-LOCAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP0]], 15 @@ -88,14 +88,14 @@ SYCL_EXTERNAL void test_load(sycl::sub_group &sg, plain_ptr p, int &out) { // CHECK-LOCAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR5]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[CALL_I_I_I_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P]], i64 [[IDXPROM_I_I_I]] -// CHECK-LOCAL-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] -// CHECK-LOCAL-NEXT: store i32 [[TMP1]], ptr addrspace(4) [[OUT:%.*]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] +// CHECK-LOCAL-NEXT: store i32 [[TMP1]], ptr addrspace(4) [[OUT:%.*]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP13:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental6detail15group_load_implINS0_9sub_groupEPU3AS3iiLm1ENS3_10propertiesINS4_20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS4_9naive_keyEJEEENSB_INS4_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeEST_SR_NS0_4spanISS_XT2_EEESU_.exit.i.i: // CHECK-LOCAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR4]] // CHECK-LOCAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS3IINS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESR_SP_RSQ_SS__EXIT:%.*]] // CHECK-LOCAL: if.end.i.i: -// CHECK-LOCAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef i32 @_Z30__spirv_SubgroupBlockReadINTELIjET_PU3AS3Kj(ptr addrspace(3) noundef nonnull [[P]]) #[[ATTR4]] +// CHECK-LOCAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef i32 @_Z30__spirv_SubgroupBlockReadINTELIjET_PU3AS3Kj(ptr addrspace(3) noundef [[P]]) #[[ATTR4]] // CHECK-LOCAL-NEXT: store i32 [[CALL6_I_I]], ptr addrspace(4) [[OUT]], align 4 // CHECK-LOCAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL10GROUP_LOADINS0_9SUB_GROUPEPU3AS3IINS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE17VERIFY_LOAD_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESR_SP_RSQ_SS__EXIT]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental10group_loadINS0_9sub_groupEPU3AS3iiNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE17verify_load_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT2_EEvE4typeESR_SP_RSQ_SS_.exit: @@ -108,7 +108,7 @@ SYCL_EXTERNAL void test_load_native(sycl::sub_group &sg, plain_ptr p, // CHECK-GLOBAL-LABEL: @_Z10test_storeRN4sycl3_V19sub_groupEiPU3AS1i( // CHECK-GLOBAL-NEXT: entry: -// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], null +// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(1)) // CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-GLOBAL-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(1) [[P]] to i64 // CHECK-GLOBAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP0]], 15 @@ -124,7 +124,7 @@ SYCL_EXTERNAL void test_load_native(sycl::sub_group &sg, plain_ptr p, // CHECK-GLOBAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR5:[0-9]+]] // CHECK-GLOBAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[CALL_I_I_I_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[P]], i64 [[IDXPROM_I_I_I]] -// CHECK-GLOBAL-NEXT: store i32 [[V:%.*]], ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7:![0-9]+]] +// CHECK-GLOBAL-NEXT: store i32 [[V:%.*]], ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6:![0-9]+]] // CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP11:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental6detail16group_store_implINS0_9sub_groupEKiLm1EPU3AS1iNS3_10propertiesINS4_20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSC_INS3_21contiguous_memory_keyEJEEENSC_INS3_14full_group_keyEJEEENSC_INS4_9naive_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T2_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESS_NS0_4spanISQ_XT1_EEESR_ST_.exit.i.i: // CHECK-GLOBAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR4]] @@ -133,14 +133,14 @@ SYCL_EXTERNAL void test_load_native(sycl::sub_group &sg, plain_ptr p, // CHECK-GLOBAL-NEXT: [[I_0_I_I:%.*]] = phi i1 [ false, [[FOR_COND_I_I]] ], [ true, [[ENTRY:%.*]] ] // CHECK-GLOBAL-NEXT: br i1 [[I_0_I_I]], label [[FOR_COND_I_I]], label [[FOR_COND_CLEANUP_I_I:%.*]], !llvm.loop [[LOOP13:![0-9]+]] // CHECK-GLOBAL: for.cond.cleanup.i.i: -// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS1jj(ptr addrspace(1) noundef nonnull [[P]], i32 noundef [[V]]) #[[ATTR4]] +// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS1jj(ptr addrspace(1) noundef [[P]], i32 noundef [[V]]) #[[ATTR4]] // CHECK-GLOBAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL11GROUP_STOREINS0_9SUB_GROUPEIPU3AS1INS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE18VERIFY_STORE_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESP_RKSN_SO_SQ__EXIT]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiPU3AS1iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT2_EEvE4typeESP_RKSN_SO_SQ_.exit: // CHECK-GLOBAL-NEXT: ret void // // CHECK-LOCAL-LABEL: @_Z10test_storeRN4sycl3_V19sub_groupEiPU3AS3i( // CHECK-LOCAL-NEXT: entry: -// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], null +// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(3)) // CHECK-LOCAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-LOCAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR4]] // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I_I:%.*]] @@ -151,7 +151,7 @@ SYCL_EXTERNAL void test_load_native(sycl::sub_group &sg, plain_ptr p, // CHECK-LOCAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR5]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[CALL_I_I_I_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P]], i64 [[IDXPROM_I_I_I]] -// CHECK-LOCAL-NEXT: store i32 [[V:%.*]], ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: store i32 [[V:%.*]], ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP14:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiPU3AS3iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT2_EEvE4typeESP_RKSN_SO_SQ_.exit: // CHECK-LOCAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR4]] @@ -163,7 +163,7 @@ SYCL_EXTERNAL void test_store(sycl::sub_group &sg, int v, plain_ptr p) { // CHECK-GLOBAL-LABEL: @_Z17test_store_nativeRN4sycl3_V19sub_groupEiPU3AS1i( // CHECK-GLOBAL-NEXT: entry: -// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], null +// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(1)) // CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-GLOBAL-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(1) [[P]] to i64 // CHECK-GLOBAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP0]], 15 @@ -179,7 +179,7 @@ SYCL_EXTERNAL void test_store(sycl::sub_group &sg, int v, plain_ptr p) { // CHECK-GLOBAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR5]] // CHECK-GLOBAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[CALL_I_I_I_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[P]], i64 [[IDXPROM_I_I_I]] -// CHECK-GLOBAL-NEXT: store i32 [[V:%.*]], ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: store i32 [[V:%.*]], ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP14:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental6detail16group_store_implINS0_9sub_groupEKiLm1EPU3AS1iNS3_10propertiesINS4_20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSC_INS3_21contiguous_memory_keyEJEEENSC_INS3_14full_group_keyEJEEENSC_INS4_9naive_keyEJEEENSC_INS4_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T2_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESU_NS0_4spanISS_XT1_EEEST_SV_.exit.i.i: // CHECK-GLOBAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR4]] @@ -188,14 +188,14 @@ SYCL_EXTERNAL void test_store(sycl::sub_group &sg, int v, plain_ptr p) { // CHECK-GLOBAL-NEXT: [[I_0_I_I:%.*]] = phi i1 [ false, [[FOR_COND_I_I]] ], [ true, [[ENTRY:%.*]] ] // CHECK-GLOBAL-NEXT: br i1 [[I_0_I_I]], label [[FOR_COND_I_I]], label [[FOR_COND_CLEANUP_I_I:%.*]], !llvm.loop [[LOOP15:![0-9]+]] // CHECK-GLOBAL: for.cond.cleanup.i.i: -// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS1jj(ptr addrspace(1) noundef nonnull [[P]], i32 noundef [[V]]) #[[ATTR4]] +// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS1jj(ptr addrspace(1) noundef [[P]], i32 noundef [[V]]) #[[ATTR4]] // CHECK-GLOBAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL11GROUP_STOREINS0_9SUB_GROUPEIPU3AS1INS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE18VERIFY_STORE_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESR_RKSP_SQ_SS__EXIT]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiPU3AS1iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT2_EEvE4typeESR_RKSP_SQ_SS_.exit: // CHECK-GLOBAL-NEXT: ret void // // CHECK-LOCAL-LABEL: @_Z17test_store_nativeRN4sycl3_V19sub_groupEiPU3AS3i( // CHECK-LOCAL-NEXT: entry: -// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], null +// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(3)) // CHECK-LOCAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-LOCAL-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(3) [[P]] to i64 // CHECK-LOCAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP0]], 15 @@ -211,7 +211,7 @@ SYCL_EXTERNAL void test_store(sycl::sub_group &sg, int v, plain_ptr p) { // CHECK-LOCAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR5]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[CALL_I_I_I_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P]], i64 [[IDXPROM_I_I_I]] -// CHECK-LOCAL-NEXT: store i32 [[V:%.*]], ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: store i32 [[V:%.*]], ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP15:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental6detail16group_store_implINS0_9sub_groupEKiLm1EPU3AS3iNS3_10propertiesINS4_20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSC_INS3_21contiguous_memory_keyEJEEENSC_INS3_14full_group_keyEJEEENSC_INS4_9naive_keyEJEEENSC_INS4_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T2_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESU_NS0_4spanISS_XT1_EEEST_SV_.exit.i.i: // CHECK-LOCAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR4]] @@ -220,7 +220,7 @@ SYCL_EXTERNAL void test_store(sycl::sub_group &sg, int v, plain_ptr p) { // CHECK-LOCAL-NEXT: [[I_0_I_I:%.*]] = phi i1 [ false, [[FOR_COND_I_I]] ], [ true, [[ENTRY:%.*]] ] // CHECK-LOCAL-NEXT: br i1 [[I_0_I_I]], label [[FOR_COND_I_I]], label [[FOR_COND_CLEANUP_I_I:%.*]], !llvm.loop [[LOOP16:![0-9]+]] // CHECK-LOCAL: for.cond.cleanup.i.i: -// CHECK-LOCAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS3jj(ptr addrspace(3) noundef nonnull [[P]], i32 noundef [[V]]) #[[ATTR4]] +// CHECK-LOCAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS3jj(ptr addrspace(3) noundef [[P]], i32 noundef [[V]]) #[[ATTR4]] // CHECK-LOCAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL11GROUP_STOREINS0_9SUB_GROUPEIPU3AS3INS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE18VERIFY_STORE_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESR_RKSP_SQ_SS__EXIT]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiPU3AS3iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT2_EEvE4typeESR_RKSP_SQ_SS_.exit: // CHECK-LOCAL-NEXT: ret void diff --git a/sycl/test/check_device_code/group_store.cpp b/sycl/test/check_device_code/group_store.cpp index 72965df838921..a070c5894e023 100644 --- a/sycl/test/check_device_code/group_store.cpp +++ b/sycl/test/check_device_code/group_store.cpp @@ -61,7 +61,7 @@ namespace blocked { // CHECK-GLOBAL-NEXT: [[CALL_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR8:[0-9]+]] // CHECK-GLOBAL-NEXT: [[IDXPROM_I_I:%.*]] = sext i32 [[CALL_I_I_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[P:%.*]], i64 [[IDXPROM_I_I]] -// CHECK-GLOBAL-NEXT: store i32 [[V:%.*]], ptr addrspace(1) [[ARRAYIDX_I_I]], align 4, !tbaa [[TBAA7:![0-9]+]] +// CHECK-GLOBAL-NEXT: store i32 [[V:%.*]], ptr addrspace(1) [[ARRAYIDX_I_I]], align 4, !tbaa [[TBAA6:![0-9]+]] // CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I]], !llvm.loop [[LOOP11:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiPU3AS1iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS9_9naive_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT2_EEvE4typeESP_RKSN_SO_SQ_.exit: // CHECK-GLOBAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR7]] @@ -78,7 +78,7 @@ namespace blocked { // CHECK-LOCAL-NEXT: [[CALL_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR8:[0-9]+]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I:%.*]] = sext i32 [[CALL_I_I_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P:%.*]], i64 [[IDXPROM_I_I]] -// CHECK-LOCAL-NEXT: store i32 [[V:%.*]], ptr addrspace(3) [[ARRAYIDX_I_I]], align 4, !tbaa [[TBAA7:![0-9]+]] +// CHECK-LOCAL-NEXT: store i32 [[V:%.*]], ptr addrspace(3) [[ARRAYIDX_I_I]], align 4, !tbaa [[TBAA6:![0-9]+]] // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I]], !llvm.loop [[LOOP11:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiPU3AS3iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS9_9naive_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT2_EEvE4typeESP_RKSN_SO_SQ_.exit: // CHECK-LOCAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR7]] @@ -91,7 +91,7 @@ SYCL_EXTERNAL void test_naive(sycl::sub_group &sg, int v, plain_ptr p) { // CHECK-GLOBAL-LABEL: @_ZN7blocked14test_optimizedERN4sycl3_V19sub_groupEiPU3AS1i( // CHECK-GLOBAL-NEXT: entry: -// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], null +// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(1)) // CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-GLOBAL-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(1) [[P]] to i64 // CHECK-GLOBAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP0]], 15 @@ -107,7 +107,7 @@ SYCL_EXTERNAL void test_naive(sycl::sub_group &sg, int v, plain_ptr p) { // CHECK-GLOBAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR8]] // CHECK-GLOBAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[CALL_I_I_I_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[P]], i64 [[IDXPROM_I_I_I]] -// CHECK-GLOBAL-NEXT: store i32 [[V:%.*]], ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: store i32 [[V:%.*]], ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP13:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental6detail16group_store_implINS0_9sub_groupEKiLm1EPU3AS1iNS3_10propertiesINS4_20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSC_INS3_21contiguous_memory_keyEJEEENSC_INS3_14full_group_keyEJEEENSC_INS4_9naive_keyEJEEENSC_INS4_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T2_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESU_NS0_4spanISS_XT1_EEEST_SV_.exit.i.i: // CHECK-GLOBAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR7]] @@ -116,14 +116,14 @@ SYCL_EXTERNAL void test_naive(sycl::sub_group &sg, int v, plain_ptr p) { // CHECK-GLOBAL-NEXT: [[I_0_I_I:%.*]] = phi i1 [ false, [[FOR_COND_I_I]] ], [ true, [[ENTRY:%.*]] ] // CHECK-GLOBAL-NEXT: br i1 [[I_0_I_I]], label [[FOR_COND_I_I]], label [[FOR_COND_CLEANUP_I_I:%.*]], !llvm.loop [[LOOP14:![0-9]+]] // CHECK-GLOBAL: for.cond.cleanup.i.i: -// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS1jj(ptr addrspace(1) noundef nonnull [[P]], i32 noundef [[V]]) #[[ATTR7]] +// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS1jj(ptr addrspace(1) noundef [[P]], i32 noundef [[V]]) #[[ATTR7]] // CHECK-GLOBAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL11GROUP_STOREINS0_9SUB_GROUPEIPU3AS1INS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE18VERIFY_STORE_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESR_RKSP_SQ_SS__EXIT]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiPU3AS1iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT2_EEvE4typeESR_RKSP_SQ_SS_.exit: // CHECK-GLOBAL-NEXT: ret void // // CHECK-LOCAL-LABEL: @_ZN7blocked14test_optimizedERN4sycl3_V19sub_groupEiPU3AS3i( // CHECK-LOCAL-NEXT: entry: -// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], null +// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(3)) // CHECK-LOCAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-LOCAL-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(3) [[P]] to i64 // CHECK-LOCAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP0]], 15 @@ -139,7 +139,7 @@ SYCL_EXTERNAL void test_naive(sycl::sub_group &sg, int v, plain_ptr p) { // CHECK-LOCAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR8]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[CALL_I_I_I_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P]], i64 [[IDXPROM_I_I_I]] -// CHECK-LOCAL-NEXT: store i32 [[V:%.*]], ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: store i32 [[V:%.*]], ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP13:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental6detail16group_store_implINS0_9sub_groupEKiLm1EPU3AS3iNS3_10propertiesINS4_20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSC_INS3_21contiguous_memory_keyEJEEENSC_INS3_14full_group_keyEJEEENSC_INS4_9naive_keyEJEEENSC_INS4_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T2_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESU_NS0_4spanISS_XT1_EEEST_SV_.exit.i.i: // CHECK-LOCAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR7]] @@ -148,7 +148,7 @@ SYCL_EXTERNAL void test_naive(sycl::sub_group &sg, int v, plain_ptr p) { // CHECK-LOCAL-NEXT: [[I_0_I_I:%.*]] = phi i1 [ false, [[FOR_COND_I_I]] ], [ true, [[ENTRY:%.*]] ] // CHECK-LOCAL-NEXT: br i1 [[I_0_I_I]], label [[FOR_COND_I_I]], label [[FOR_COND_CLEANUP_I_I:%.*]], !llvm.loop [[LOOP14:![0-9]+]] // CHECK-LOCAL: for.cond.cleanup.i.i: -// CHECK-LOCAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS3jj(ptr addrspace(3) noundef nonnull [[P]], i32 noundef [[V]]) #[[ATTR7]] +// CHECK-LOCAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS3jj(ptr addrspace(3) noundef [[P]], i32 noundef [[V]]) #[[ATTR7]] // CHECK-LOCAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL11GROUP_STOREINS0_9SUB_GROUPEIPU3AS3INS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE18VERIFY_STORE_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESR_RKSP_SQ_SS__EXIT]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiPU3AS3iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT2_EEvE4typeESR_RKSP_SQ_SS_.exit: // CHECK-LOCAL-NEXT: ret void @@ -161,7 +161,7 @@ SYCL_EXTERNAL void test_optimized(sycl::sub_group &sg, int v, // CHECK-GLOBAL-LABEL: @_ZN7blocked27test_contiguous_auto_detectERN4sycl3_V19sub_groupEiPU3AS1i( // CHECK-GLOBAL-NEXT: entry: -// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], null +// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(1)) // CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-GLOBAL-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(1) [[P]] to i64 // CHECK-GLOBAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP0]], 15 @@ -177,7 +177,7 @@ SYCL_EXTERNAL void test_optimized(sycl::sub_group &sg, int v, // CHECK-GLOBAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR8]] // CHECK-GLOBAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[CALL_I_I_I_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[P]], i64 [[IDXPROM_I_I_I]] -// CHECK-GLOBAL-NEXT: store i32 [[V:%.*]], ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: store i32 [[V:%.*]], ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP15:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental6detail16group_store_implINS0_9sub_groupEKiLm1EPU3AS1iNS3_10propertiesINS4_20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSC_INS3_14full_group_keyEJEEENSC_INS4_9naive_keyEJEEENSC_INS4_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T2_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESS_NS0_4spanISQ_XT1_EEESR_ST_.exit.i.i: // CHECK-GLOBAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR7]] @@ -186,14 +186,14 @@ SYCL_EXTERNAL void test_optimized(sycl::sub_group &sg, int v, // CHECK-GLOBAL-NEXT: [[I_0_I_I:%.*]] = phi i1 [ false, [[FOR_COND_I_I]] ], [ true, [[ENTRY:%.*]] ] // CHECK-GLOBAL-NEXT: br i1 [[I_0_I_I]], label [[FOR_COND_I_I]], label [[FOR_COND_CLEANUP_I_I:%.*]], !llvm.loop [[LOOP16:![0-9]+]] // CHECK-GLOBAL: for.cond.cleanup.i.i: -// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS1jj(ptr addrspace(1) noundef nonnull [[P]], i32 noundef [[V]]) #[[ATTR7]] +// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS1jj(ptr addrspace(1) noundef [[P]], i32 noundef [[V]]) #[[ATTR7]] // CHECK-GLOBAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL11GROUP_STOREINS0_9SUB_GROUPEIPU3AS1INS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE18VERIFY_STORE_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESP_RKSN_SO_SQ__EXIT]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiPU3AS1iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT2_EEvE4typeESP_RKSN_SO_SQ_.exit: // CHECK-GLOBAL-NEXT: ret void // // CHECK-LOCAL-LABEL: @_ZN7blocked27test_contiguous_auto_detectERN4sycl3_V19sub_groupEiPU3AS3i( // CHECK-LOCAL-NEXT: entry: -// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], null +// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(3)) // CHECK-LOCAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-LOCAL-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(3) [[P]] to i64 // CHECK-LOCAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP0]], 15 @@ -209,7 +209,7 @@ SYCL_EXTERNAL void test_optimized(sycl::sub_group &sg, int v, // CHECK-LOCAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR8]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[CALL_I_I_I_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P]], i64 [[IDXPROM_I_I_I]] -// CHECK-LOCAL-NEXT: store i32 [[V:%.*]], ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: store i32 [[V:%.*]], ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP15:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental6detail16group_store_implINS0_9sub_groupEKiLm1EPU3AS3iNS3_10propertiesINS4_20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSC_INS3_14full_group_keyEJEEENSC_INS4_9naive_keyEJEEENSC_INS4_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T2_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESS_NS0_4spanISQ_XT1_EEESR_ST_.exit.i.i: // CHECK-LOCAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR7]] @@ -218,7 +218,7 @@ SYCL_EXTERNAL void test_optimized(sycl::sub_group &sg, int v, // CHECK-LOCAL-NEXT: [[I_0_I_I:%.*]] = phi i1 [ false, [[FOR_COND_I_I]] ], [ true, [[ENTRY:%.*]] ] // CHECK-LOCAL-NEXT: br i1 [[I_0_I_I]], label [[FOR_COND_I_I]], label [[FOR_COND_CLEANUP_I_I:%.*]], !llvm.loop [[LOOP16:![0-9]+]] // CHECK-LOCAL: for.cond.cleanup.i.i: -// CHECK-LOCAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS3jj(ptr addrspace(3) noundef nonnull [[P]], i32 noundef [[V]]) #[[ATTR7]] +// CHECK-LOCAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS3jj(ptr addrspace(3) noundef [[P]], i32 noundef [[V]]) #[[ATTR7]] // CHECK-LOCAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL11GROUP_STOREINS0_9SUB_GROUPEIPU3AS3INS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_14FULL_GROUP_KEYEJEEENSB_INS9_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE18VERIFY_STORE_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESP_RKSN_SO_SQ__EXIT]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiPU3AS3iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT2_EEvE4typeESP_RKSN_SO_SQ_.exit: // CHECK-LOCAL-NEXT: ret void @@ -248,7 +248,7 @@ using accessor_iter_t = accessor v, // CHECK-GLOBAL-NEXT: [[VALUES_I_I:%.*]] = alloca [4 x i16], align 2 // CHECK-GLOBAL-NEXT: [[TMP0:%.*]] = load i64, ptr [[V:%.*]], align 8, !tbaa [[TBAA30]] // CHECK-GLOBAL-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4) -// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], null +// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(1)) // CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = ptrtoint ptr addrspace(1) [[P]] to i64 // CHECK-GLOBAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP2]], 15 @@ -493,7 +493,7 @@ SYCL_EXTERNAL void test_four_shorts(sycl::sub_group &sg, span v, // CHECK-GLOBAL-NEXT: br i1 [[CMP_I_I]], label [[FOR_BODY_I_I]], label [[FOR_COND_CLEANUP_I_I:%.*]] // CHECK-GLOBAL: for.cond.cleanup.i.i: // CHECK-GLOBAL-NEXT: [[TMP4:%.*]] = load i64, ptr [[VALUES_I_I]], align 2, !tbaa [[TBAA35]] -// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS1mm(ptr addrspace(1) noundef nonnull [[P]], i64 noundef [[TMP4]]) #[[ATTR7]] +// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS1mm(ptr addrspace(1) noundef [[P]], i64 noundef [[TMP4]]) #[[ATTR7]] // CHECK-GLOBAL-NEXT: call void @llvm.lifetime.end.p0(ptr nonnull [[VALUES_I_I]]) #[[ATTR9]] // CHECK-GLOBAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL11GROUP_STOREINS0_9SUB_GROUPEKSLM4EPU3AS1SNS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSC_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSC_INS3_14FULL_GROUP_KEYEJEEENSC_INSA_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE18VERIFY_STORE_TYPESIT0_T2_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT3_EEVE4TYPEESS_NS0_4SPANISQ_XT1_EEESR_ST__EXIT]] // CHECK-GLOBAL: for.body.i.i: @@ -512,7 +512,7 @@ SYCL_EXTERNAL void test_four_shorts(sycl::sub_group &sg, span v, // CHECK-LOCAL-NEXT: [[VALUES_I_I:%.*]] = alloca [4 x i16], align 2 // CHECK-LOCAL-NEXT: [[TMP0:%.*]] = load i64, ptr [[V:%.*]], align 8, !tbaa [[TBAA30]] // CHECK-LOCAL-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4) -// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], null +// CHECK-LOCAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(3) [[P:%.*]], addrspacecast (ptr addrspace(4) null to ptr addrspace(3)) // CHECK-LOCAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]]) // CHECK-LOCAL-NEXT: [[TMP2:%.*]] = ptrtoint ptr addrspace(3) [[P]] to i64 // CHECK-LOCAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP2]], 15 @@ -549,7 +549,7 @@ SYCL_EXTERNAL void test_four_shorts(sycl::sub_group &sg, span v, // CHECK-LOCAL-NEXT: br i1 [[CMP_I_I]], label [[FOR_BODY_I_I]], label [[FOR_COND_CLEANUP_I_I:%.*]] // CHECK-LOCAL: for.cond.cleanup.i.i: // CHECK-LOCAL-NEXT: [[TMP4:%.*]] = load i64, ptr [[VALUES_I_I]], align 2, !tbaa [[TBAA35]] -// CHECK-LOCAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS3mm(ptr addrspace(3) noundef nonnull [[P]], i64 noundef [[TMP4]]) #[[ATTR7]] +// CHECK-LOCAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS3mm(ptr addrspace(3) noundef [[P]], i64 noundef [[TMP4]]) #[[ATTR7]] // CHECK-LOCAL-NEXT: call void @llvm.lifetime.end.p0(ptr nonnull [[VALUES_I_I]]) #[[ATTR9]] // CHECK-LOCAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL11GROUP_STOREINS0_9SUB_GROUPEKSLM4EPU3AS3SNS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSC_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSC_INS3_14FULL_GROUP_KEYEJEEENSC_INSA_25NATIVE_LOCAL_BLOCK_IO_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE18VERIFY_STORE_TYPESIT0_T2_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT3_EEVE4TYPEESS_NS0_4SPANISQ_XT1_EEESR_ST__EXIT]] // CHECK-LOCAL: for.body.i.i: @@ -583,13 +583,13 @@ SYCL_EXTERNAL void test_four_const_shorts(sycl::sub_group &sg, // CHECK-GLOBAL: for.body.i.i.i: // CHECK-GLOBAL-NEXT: [[CONV_I_I_I:%.*]] = zext nneg i32 [[I_0_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR8]] // CHECK-GLOBAL-NEXT: [[MUL_I_I_I_I:%.*]] = mul i32 [[CALL_I_I_I_I_I_I]], 3 // CHECK-GLOBAL-NEXT: [[ADD_I_I_I_I:%.*]] = add i32 [[MUL_I_I_I_I]], [[I_0_I_I_I]] // CHECK-GLOBAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[P:%.*]], i64 [[IDXPROM_I_I_I]] -// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP39:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiLm3EPU3AS1iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T2_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_NS0_4spanISP_XT1_EEESQ_SS_.exit: @@ -609,13 +609,13 @@ SYCL_EXTERNAL void test_four_const_shorts(sycl::sub_group &sg, // CHECK-LOCAL: for.body.i.i.i: // CHECK-LOCAL-NEXT: [[CONV_I_I_I:%.*]] = zext nneg i32 [[I_0_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR8]] // CHECK-LOCAL-NEXT: [[MUL_I_I_I_I:%.*]] = mul i32 [[CALL_I_I_I_I_I_I]], 3 // CHECK-LOCAL-NEXT: [[ADD_I_I_I_I:%.*]] = add i32 [[MUL_I_I_I_I]], [[I_0_I_I_I]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P:%.*]], i64 [[IDXPROM_I_I_I]] -// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP39:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiLm3EPU3AS3iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T2_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_NS0_4spanISP_XT1_EEESQ_SS_.exit: @@ -641,13 +641,13 @@ SYCL_EXTERNAL void test_non_power_of_two(sycl::sub_group &sg, span v, // CHECK-GLOBAL: for.body.i.i.i: // CHECK-GLOBAL-NEXT: [[CONV_I_I_I:%.*]] = zext nneg i32 [[I_0_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR8]] // CHECK-GLOBAL-NEXT: [[MUL_I_I_I_I:%.*]] = shl i32 [[CALL_I_I_I_I_I_I]], 2 // CHECK-GLOBAL-NEXT: [[ADD_I_I_I_I:%.*]] = or disjoint i32 [[MUL_I_I_I_I]], [[I_0_I_I_I]] // CHECK-GLOBAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[P:%.*]], i64 [[IDXPROM_I_I_I]] -// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP40:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiLm4EPU3AS1iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T2_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_NS0_4spanISP_XT1_EEESQ_SS_.exit: @@ -667,13 +667,13 @@ SYCL_EXTERNAL void test_non_power_of_two(sycl::sub_group &sg, span v, // CHECK-LOCAL: for.body.i.i.i: // CHECK-LOCAL-NEXT: [[CONV_I_I_I:%.*]] = zext nneg i32 [[I_0_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR8]] // CHECK-LOCAL-NEXT: [[MUL_I_I_I_I:%.*]] = shl i32 [[CALL_I_I_I_I_I_I]], 2 // CHECK-LOCAL-NEXT: [[ADD_I_I_I_I:%.*]] = or disjoint i32 [[MUL_I_I_I_I]], [[I_0_I_I_I]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P:%.*]], i64 [[IDXPROM_I_I_I]] -// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP40:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiLm4EPU3AS3iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T2_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_NS0_4spanISP_XT1_EEESQ_SS_.exit: @@ -700,13 +700,13 @@ SYCL_EXTERNAL void test_four_ints(sycl::sub_group &sg, span v, // CHECK-GLOBAL: for.body.i.i.i: // CHECK-GLOBAL-NEXT: [[CONV_I_I_I:%.*]] = zext nneg i32 [[I_0_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR8]] // CHECK-GLOBAL-NEXT: [[MUL_I_I_I_I:%.*]] = mul i32 [[CALL_I_I_I_I_I_I]], 7 // CHECK-GLOBAL-NEXT: [[ADD_I_I_I_I:%.*]] = add i32 [[MUL_I_I_I_I]], [[I_0_I_I_I]] // CHECK-GLOBAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[P:%.*]], i64 [[IDXPROM_I_I_I]] -// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP41:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiLm7EPU3AS1iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T2_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_NS0_4spanISP_XT1_EEESQ_SS_.exit: @@ -726,13 +726,13 @@ SYCL_EXTERNAL void test_four_ints(sycl::sub_group &sg, span v, // CHECK-LOCAL: for.body.i.i.i: // CHECK-LOCAL-NEXT: [[CONV_I_I_I:%.*]] = zext nneg i32 [[I_0_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR8]] // CHECK-LOCAL-NEXT: [[MUL_I_I_I_I:%.*]] = mul i32 [[CALL_I_I_I_I_I_I]], 7 // CHECK-LOCAL-NEXT: [[ADD_I_I_I_I:%.*]] = add i32 [[MUL_I_I_I_I]], [[I_0_I_I_I]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P:%.*]], i64 [[IDXPROM_I_I_I]] -// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP41:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiLm7EPU3AS3iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T2_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_NS0_4spanISP_XT1_EEESQ_SS_.exit: @@ -760,14 +760,14 @@ namespace striped { // CHECK-GLOBAL: for.body.i.i: // CHECK-GLOBAL-NEXT: [[CONV_I_I:%.*]] = zext nneg i32 [[I_0_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I]] -// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[CALL_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR8]] // CHECK-GLOBAL-NEXT: [[CALL_I_I2_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z27__spirv_BuiltInSubgroupSizev() #[[ATTR8]] // CHECK-GLOBAL-NEXT: [[MUL_I_I_I:%.*]] = mul nuw nsw i32 [[CALL_I_I2_I_I_I]], [[I_0_I_I]] // CHECK-GLOBAL-NEXT: [[ADD_I_I_I:%.*]] = add i32 [[CALL_I_I_I_I_I]], [[MUL_I_I_I]] // CHECK-GLOBAL-NEXT: [[IDXPROM_I_I:%.*]] = sext i32 [[ADD_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[P:%.*]], i64 [[IDXPROM_I_I]] -// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[ARRAYIDX_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[ARRAYIDX_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[INC_I_I]] = add nuw nsw i32 [[I_0_I_I]], 1 // CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I]], !llvm.loop [[LOOP42:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiLm2EPU3AS1iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS9_9naive_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T2_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESP_NS0_4spanISN_XT1_EEESO_SQ_.exit: @@ -787,14 +787,14 @@ namespace striped { // CHECK-LOCAL: for.body.i.i: // CHECK-LOCAL-NEXT: [[CONV_I_I:%.*]] = zext nneg i32 [[I_0_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I]] -// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[CALL_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR8]] // CHECK-LOCAL-NEXT: [[CALL_I_I2_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z27__spirv_BuiltInSubgroupSizev() #[[ATTR8]] // CHECK-LOCAL-NEXT: [[MUL_I_I_I:%.*]] = mul nuw nsw i32 [[CALL_I_I2_I_I_I]], [[I_0_I_I]] // CHECK-LOCAL-NEXT: [[ADD_I_I_I:%.*]] = add i32 [[CALL_I_I_I_I_I]], [[MUL_I_I_I]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I:%.*]] = sext i32 [[ADD_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P:%.*]], i64 [[IDXPROM_I_I]] -// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(3) [[ARRAYIDX_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(3) [[ARRAYIDX_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[INC_I_I]] = add nuw nsw i32 [[I_0_I_I]], 1 // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I]], !llvm.loop [[LOOP42:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiLm2EPU3AS3iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS9_9naive_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T2_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESP_NS0_4spanISN_XT1_EEESO_SQ_.exit: @@ -905,14 +905,14 @@ using accessor_iter_t = accessor v, // CHECK-GLOBAL: for.body.i.i.i: // CHECK-GLOBAL-NEXT: [[CONV_I_I_I:%.*]] = zext nneg i32 [[I_0_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR8]] // CHECK-GLOBAL-NEXT: [[CALL_I_I2_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z27__spirv_BuiltInSubgroupSizev() #[[ATTR8]] // CHECK-GLOBAL-NEXT: [[MUL_I_I_I_I:%.*]] = mul i32 [[CALL_I_I2_I_I_I_I]], [[I_0_I_I_I]] // CHECK-GLOBAL-NEXT: [[ADD_I_I_I_I:%.*]] = add i32 [[CALL_I_I_I_I_I_I]], [[MUL_I_I_I_I]] // CHECK-GLOBAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[P:%.*]], i64 [[IDXPROM_I_I_I]] -// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP63:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiLm3EPU3AS1iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T2_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_NS0_4spanISP_XT1_EEESQ_SS_.exit: @@ -1119,14 +1119,14 @@ SYCL_EXTERNAL void test_sixteen_shorts(sycl::sub_group &sg, span v, // CHECK-LOCAL: for.body.i.i.i: // CHECK-LOCAL-NEXT: [[CONV_I_I_I:%.*]] = zext nneg i32 [[I_0_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR8]] // CHECK-LOCAL-NEXT: [[CALL_I_I2_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z27__spirv_BuiltInSubgroupSizev() #[[ATTR8]] // CHECK-LOCAL-NEXT: [[MUL_I_I_I_I:%.*]] = mul i32 [[CALL_I_I2_I_I_I_I]], [[I_0_I_I_I]] // CHECK-LOCAL-NEXT: [[ADD_I_I_I_I:%.*]] = add i32 [[CALL_I_I_I_I_I_I]], [[MUL_I_I_I_I]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P:%.*]], i64 [[IDXPROM_I_I_I]] -// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP63:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiLm3EPU3AS3iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T2_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_NS0_4spanISP_XT1_EEESQ_SS_.exit: @@ -1152,14 +1152,14 @@ SYCL_EXTERNAL void test_non_power_of_two(sycl::sub_group &sg, span v, // CHECK-GLOBAL: for.body.i.i.i: // CHECK-GLOBAL-NEXT: [[CONV_I_I_I:%.*]] = zext nneg i32 [[I_0_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR8]] // CHECK-GLOBAL-NEXT: [[CALL_I_I2_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z27__spirv_BuiltInSubgroupSizev() #[[ATTR8]] // CHECK-GLOBAL-NEXT: [[MUL_I_I_I_I:%.*]] = mul i32 [[CALL_I_I2_I_I_I_I]], [[I_0_I_I_I]] // CHECK-GLOBAL-NEXT: [[ADD_I_I_I_I:%.*]] = add i32 [[CALL_I_I_I_I_I_I]], [[MUL_I_I_I_I]] // CHECK-GLOBAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[P:%.*]], i64 [[IDXPROM_I_I_I]] -// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP64:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiLm16EPU3AS1iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T2_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_NS0_4spanISP_XT1_EEESQ_SS_.exit: @@ -1179,14 +1179,14 @@ SYCL_EXTERNAL void test_non_power_of_two(sycl::sub_group &sg, span v, // CHECK-LOCAL: for.body.i.i.i: // CHECK-LOCAL-NEXT: [[CONV_I_I_I:%.*]] = zext nneg i32 [[I_0_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR8]] // CHECK-LOCAL-NEXT: [[CALL_I_I2_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z27__spirv_BuiltInSubgroupSizev() #[[ATTR8]] // CHECK-LOCAL-NEXT: [[MUL_I_I_I_I:%.*]] = mul i32 [[CALL_I_I2_I_I_I_I]], [[I_0_I_I_I]] // CHECK-LOCAL-NEXT: [[ADD_I_I_I_I:%.*]] = add i32 [[CALL_I_I_I_I_I_I]], [[MUL_I_I_I_I]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P:%.*]], i64 [[IDXPROM_I_I_I]] -// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP64:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiLm16EPU3AS3iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T2_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_NS0_4spanISP_XT1_EEESQ_SS_.exit: @@ -1213,14 +1213,14 @@ SYCL_EXTERNAL void test_sixteen_ints(sycl::sub_group &sg, span v, // CHECK-GLOBAL: for.body.i.i.i: // CHECK-GLOBAL-NEXT: [[CONV_I_I_I:%.*]] = zext nneg i32 [[I_0_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR8]] // CHECK-GLOBAL-NEXT: [[CALL_I_I2_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z27__spirv_BuiltInSubgroupSizev() #[[ATTR8]] // CHECK-GLOBAL-NEXT: [[MUL_I_I_I_I:%.*]] = mul i32 [[CALL_I_I2_I_I_I_I]], [[I_0_I_I_I]] // CHECK-GLOBAL-NEXT: [[ADD_I_I_I_I:%.*]] = add i32 [[CALL_I_I_I_I_I_I]], [[MUL_I_I_I_I]] // CHECK-GLOBAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[P:%.*]], i64 [[IDXPROM_I_I_I]] -// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-GLOBAL-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-GLOBAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP65:![0-9]+]] // CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiLm11EPU3AS1iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T2_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_NS0_4spanISP_XT1_EEESQ_SS_.exit: @@ -1240,14 +1240,14 @@ SYCL_EXTERNAL void test_sixteen_ints(sycl::sub_group &sg, span v, // CHECK-LOCAL: for.body.i.i.i: // CHECK-LOCAL-NEXT: [[CONV_I_I_I:%.*]] = zext nneg i32 [[I_0_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I_I:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) [[TMP1]], i64 [[CONV_I_I_I]] -// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX_I_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR8]] // CHECK-LOCAL-NEXT: [[CALL_I_I2_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z27__spirv_BuiltInSubgroupSizev() #[[ATTR8]] // CHECK-LOCAL-NEXT: [[MUL_I_I_I_I:%.*]] = mul i32 [[CALL_I_I2_I_I_I_I]], [[I_0_I_I_I]] // CHECK-LOCAL-NEXT: [[ADD_I_I_I_I:%.*]] = add i32 [[CALL_I_I_I_I_I_I]], [[MUL_I_I_I_I]] // CHECK-LOCAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[ADD_I_I_I_I]] to i64 // CHECK-LOCAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(3) [[P:%.*]], i64 [[IDXPROM_I_I_I]] -// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA7]] +// CHECK-LOCAL-NEXT: store i32 [[TMP2]], ptr addrspace(3) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA6]] // CHECK-LOCAL-NEXT: [[INC_I_I_I]] = add nuw nsw i32 [[I_0_I_I_I]], 1 // CHECK-LOCAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP65:![0-9]+]] // CHECK-LOCAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiLm11EPU3AS3iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi1EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEENSB_INS9_25native_local_block_io_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T2_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT3_EEvE4typeESR_NS0_4spanISP_XT1_EEESQ_SS_.exit: