Skip to content

Commit a9512f2

Browse files
jsjijsji
committed
[NFC][SYCL] Auto regen check_device_code tests (#32059)
Due to SimplifyCFG change in llvm/llvm-project@ea2f539
1 parent d76519c commit a9512f2

File tree

9 files changed

+723
-495
lines changed

9 files changed

+723
-495
lines changed

‎sycl/test/check_device_code/group_load.cpp‎

Lines changed: 142 additions & 86 deletions
Large diffs are not rendered by default.

‎sycl/test/check_device_code/group_load_store_alignment.cpp‎

Lines changed: 44 additions & 23 deletions
Original file line numberDiff line numberDiff line change
@@ -26,20 +26,27 @@ using plain_ptr = typename sycl::detail::DecoratedType<
2626
// CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]])
2727
// CHECK-GLOBAL-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(1) [[P]] to i64
2828
// CHECK-GLOBAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP0]], 3
29-
// CHECK-GLOBAL-NEXT: [[CMP1_I_NOT_I_I:%.*]] = icmp eq i64 [[REM_I_I_I]], 0
30-
// CHECK-GLOBAL-NEXT: br i1 [[CMP1_I_NOT_I_I]], label [[IF_END_I_I:%.*]], label [[IF_THEN_I_I:%.*]]
29+
// CHECK-GLOBAL-NEXT: [[CMP_I15_I_I:%.*]] = icmp eq i64 [[REM_I_I_I]], 0
30+
// CHECK-GLOBAL-NEXT: br i1 [[CMP_I15_I_I]], label [[IF_END_I_I:%.*]], label [[IF_THEN_I_I:%.*]]
3131
// CHECK-GLOBAL: if.then.i.i:
32-
// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR3:[0-9]+]]
33-
// CHECK-GLOBAL-NEXT: [[TMP1:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv()
34-
// CHECK-GLOBAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[TMP1]] to i64
32+
// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR4:[0-9]+]]
33+
// CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I:%.*]]
34+
// CHECK-GLOBAL: for.cond.i.i.i:
35+
// CHECK-GLOBAL-NEXT: [[CMP_I19_I_I:%.*]] = phi i1 [ true, [[IF_THEN_I_I]] ], [ false, [[FOR_BODY_I_I_I:%.*]] ]
36+
// CHECK-GLOBAL-NEXT: br i1 [[CMP_I19_I_I]], label [[FOR_BODY_I_I_I]], label [[_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:%.*]]
37+
// CHECK-GLOBAL: for.body.i.i.i:
38+
// CHECK-GLOBAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR5:[0-9]+]]
39+
// CHECK-GLOBAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[CALL_I_I_I_I_I_I]] to i64
3540
// CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i16, ptr addrspace(1) [[P]], i64 [[IDXPROM_I_I_I]]
36-
// CHECK-GLOBAL-NEXT: [[TMP2:%.*]] = load i16, ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 2, !tbaa [[TBAA8:![0-9]+]]
37-
// CHECK-GLOBAL-NEXT: store i16 [[TMP2]], ptr addrspace(4) [[OUT:%.*]], align 2, !tbaa [[TBAA8]]
38-
// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR3]]
41+
// CHECK-GLOBAL-NEXT: [[TMP1:%.*]] = load i16, ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 2, !tbaa [[TBAA7:![0-9]+]]
42+
// CHECK-GLOBAL-NEXT: store i16 [[TMP1]], ptr addrspace(4) [[OUT:%.*]], align 2, !tbaa [[TBAA7]]
43+
// CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP11:![0-9]+]]
44+
// 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:
45+
// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR4]]
3946
// 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:%.*]]
4047
// CHECK-GLOBAL: if.end.i.i:
41-
// CHECK-GLOBAL-NEXT: [[CALL4_I_I:%.*]] = tail call spir_func noundef zeroext i16 @_Z30__spirv_SubgroupBlockReadINTELItET_PU3AS1Kt(ptr addrspace(1) noundef nonnull [[P]]) #[[ATTR3]]
42-
// CHECK-GLOBAL-NEXT: store i16 [[CALL4_I_I]], ptr addrspace(4) [[OUT]], align 2
48+
// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef zeroext i16 @_Z30__spirv_SubgroupBlockReadINTELItET_PU3AS1Kt(ptr addrspace(1) noundef nonnull [[P]]) #[[ATTR4]]
49+
// CHECK-GLOBAL-NEXT: store i16 [[CALL6_I_I]], ptr addrspace(4) [[OUT]], align 2
4350
// 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]]
4451
// 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:
4552
// CHECK-GLOBAL-NEXT: ret void
@@ -50,13 +57,12 @@ SYCL_EXTERNAL void test_load_without_alignment_hint(sycl::sub_group &sg,
5057
group_load(sg, p, out, opt_blocked{});
5158
}
5259

53-
//
5460
// CHECK-GLOBAL-LABEL: @_Z29test_load_with_alignment_hintRN4sycl3_V19sub_groupEPU3AS1sRs(
5561
// CHECK-GLOBAL-NEXT: entry:
5662
// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], null
5763
// CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]])
58-
// CHECK-GLOBAL-NEXT: [[CALL1_I_I:%.*]] = tail call spir_func noundef zeroext i16 @_Z30__spirv_SubgroupBlockReadINTELItET_PU3AS1Kt(ptr addrspace(1) noundef nonnull [[P]]) #[[ATTR3]]
59-
// CHECK-GLOBAL-NEXT: store i16 [[CALL1_I_I]], ptr addrspace(4) [[OUT:%.*]], align 2
64+
// CHECK-GLOBAL-NEXT: [[CALL6_I_I:%.*]] = tail call spir_func noundef zeroext i16 @_Z30__spirv_SubgroupBlockReadINTELItET_PU3AS1Kt(ptr addrspace(1) noundef nonnull [[P]]) #[[ATTR4]]
65+
// CHECK-GLOBAL-NEXT: store i16 [[CALL6_I_I]], ptr addrspace(4) [[OUT:%.*]], align 2
6066
// CHECK-GLOBAL-NEXT: ret void
6167
//
6268
SYCL_EXTERNAL void test_load_with_alignment_hint(sycl::sub_group &sg,
@@ -71,18 +77,28 @@ SYCL_EXTERNAL void test_load_with_alignment_hint(sycl::sub_group &sg,
7177
// CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]])
7278
// CHECK-GLOBAL-NEXT: [[TMP0:%.*]] = ptrtoint ptr addrspace(1) [[P]] to i64
7379
// CHECK-GLOBAL-NEXT: [[REM_I_I_I:%.*]] = and i64 [[TMP0]], 15
74-
// CHECK-GLOBAL-NEXT: [[CMP1_I_NOT_I_I:%.*]] = icmp eq i64 [[REM_I_I_I]], 0
75-
// CHECK-GLOBAL-NEXT: br i1 [[CMP1_I_NOT_I_I]], label [[IF_END_I_I:%.*]], label [[IF_THEN_I_I:%.*]]
80+
// CHECK-GLOBAL-NEXT: [[CMP_I25_I_I:%.*]] = icmp eq i64 [[REM_I_I_I]], 0
81+
// CHECK-GLOBAL-NEXT: br i1 [[CMP_I25_I_I]], label [[FOR_COND_I_I:%.*]], label [[IF_THEN_I_I:%.*]]
7682
// CHECK-GLOBAL: if.then.i.i:
77-
// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR3]]
78-
// CHECK-GLOBAL-NEXT: tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv()
79-
// CHECK-GLOBAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[TMP1]] to i64
83+
// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR4]]
84+
// CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I:%.*]]
85+
// CHECK-GLOBAL: for.cond.i.i.i:
86+
// CHECK-GLOBAL-NEXT: [[CMP_I29_I_I:%.*]] = phi i1 [ true, [[IF_THEN_I_I]] ], [ false, [[FOR_BODY_I_I_I:%.*]] ]
87+
// CHECK-GLOBAL-NEXT: br i1 [[CMP_I29_I_I]], label [[FOR_BODY_I_I_I]], label [[_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:%.*]]
88+
// CHECK-GLOBAL: for.body.i.i.i:
89+
// CHECK-GLOBAL-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #[[ATTR5]]
90+
// CHECK-GLOBAL-NEXT: [[IDXPROM_I_I_I:%.*]] = sext i32 [[CALL_I_I_I_I_I_I]] to i64
8091
// CHECK-GLOBAL-NEXT: [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds i32, ptr addrspace(1) [[P]], i64 [[IDXPROM_I_I_I]]
81-
// CHECK-GLOBAL-NEXT: store i32 [[V:%.*]], ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA14:![0-9]+]]
82-
// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR3]]
92+
// CHECK-GLOBAL-NEXT: store i32 [[V:%.*]], ptr addrspace(1) [[ARRAYIDX_I_I_I]], align 4, !tbaa [[TBAA13:![0-9]+]]
93+
// CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I_I]], !llvm.loop [[LOOP15:![0-9]+]]
94+
// 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:
95+
// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR4]]
8396
// 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:%.*]]
84-
// CHECK-GLOBAL: if.end.i.i:
85-
// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS1jj(ptr addrspace(1) noundef nonnull [[P]], i32 noundef [[V]]) #[[ATTR3]]
97+
// CHECK-GLOBAL: for.cond.i.i:
98+
// CHECK-GLOBAL-NEXT: [[I_0_I_I:%.*]] = phi i1 [ false, [[FOR_COND_I_I]] ], [ true, [[ENTRY:%.*]] ]
99+
// 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]+]]
100+
// CHECK-GLOBAL: for.cond.cleanup.i.i:
101+
// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS1jj(ptr addrspace(1) noundef nonnull [[P]], i32 noundef [[V]]) #[[ATTR4]]
86102
// 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]]
87103
// 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:
88104
// CHECK-GLOBAL-NEXT: ret void
@@ -96,7 +112,12 @@ SYCL_EXTERNAL void test_store_without_alignment_hint(sycl::sub_group &sg, int v,
96112
// CHECK-GLOBAL-NEXT: entry:
97113
// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], null
98114
// CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]])
99-
// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS1jj(ptr addrspace(1) noundef nonnull [[P]], i32 noundef [[V:%.*]]) #[[ATTR3]]
115+
// CHECK-GLOBAL-NEXT: br label [[FOR_COND_I_I:%.*]]
116+
// CHECK-GLOBAL: for.cond.i.i:
117+
// CHECK-GLOBAL-NEXT: [[I_0_I_I:%.*]] = phi i1 [ true, [[ENTRY:%.*]] ], [ false, [[FOR_COND_I_I]] ]
118+
// 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]+]]
119+
// 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:
120+
// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS1jj(ptr addrspace(1) noundef nonnull [[P]], i32 noundef [[V:%.*]]) #[[ATTR4]]
100121
// CHECK-GLOBAL-NEXT: ret void
101122
//
102123
SYCL_EXTERNAL void test_store_with_alignment_hint(sycl::sub_group &sg, int v,

0 commit comments

Comments
 (0)