[NFC][SYCL] Regen group_shuffle.cpp checks

Author: jsji

sycl/test/check_device_code/group_shuffle.cpp

@@ -14,10 +14,10 @@ using namespace sycl::ext::oneapi;
 // CHECK-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds nuw %"class.sycl::_V1::vec", ptr addrspace(4) [[BUF:%.*]], i64 [[ID:%.*]]
 // CHECK-NEXT:    [[TMP0:%.*]] = load i64, ptr addrspace(4) [[ARRAYIDX]], align 8
 // CHECK-NEXT:    call void @llvm.lifetime.start.p0(ptr nonnull [[REF_TMP]]) #[[ATTR5:[0-9]+]]
-// CHECK-NEXT:    tail call void @llvm.experimental.noalias.scope.decl(metadata [[META7:![0-9]+]])
-// CHECK-NEXT:    tail call void @llvm.experimental.noalias.scope.decl(metadata [[META10:![0-9]+]])
-// CHECK-NEXT:    call void @llvm.lifetime.start.p0(ptr nonnull [[AGG_TMP14_I]]), !noalias [[META7]]
-// CHECK-NEXT:    store i64 [[TMP0]], ptr [[AGG_TMP14_I]], align 8, !noalias [[META7]]
+// CHECK-NEXT:    tail call void @llvm.experimental.noalias.scope.decl(metadata [[META11:![0-9]+]])
+// CHECK-NEXT:    tail call void @llvm.experimental.noalias.scope.decl(metadata [[META14:![0-9]+]])
+// CHECK-NEXT:    call void @llvm.lifetime.start.p0(ptr nonnull [[AGG_TMP14_I]]), !noalias [[META11]]
+// CHECK-NEXT:    store i64 [[TMP0]], ptr [[AGG_TMP14_I]], align 8, !noalias [[META11]]
 // CHECK-NEXT:    br label [[FOR_COND_I_I:%.*]]
 // CHECK:       for.cond.i.i:
 // CHECK-NEXT:    [[S_0_I_I:%.*]] = phi i32 [ 0, [[ENTRY:%.*]] ], [ [[INC_I_I:%.*]], [[FOR_BODY_I_I:%.*]] ]
@@ -26,14 +26,14 @@ using namespace sycl::ext::oneapi;
 // CHECK:       for.body.i.i:
 // CHECK-NEXT:    [[IDXPROM_I_I_I:%.*]] = zext nneg i32 [[S_0_I_I]] to i64
 // CHECK-NEXT:    [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds %"class.sycl::_V1::ext::oneapi::bfloat16", ptr [[AGG_TMP14_I]], i64 [[IDXPROM_I_I_I]]
-// CHECK-NEXT:    [[TMP1:%.*]] = load i16, ptr [[ARRAYIDX_I_I_I]], align 2, !tbaa [[TBAA13:![0-9]+]], !noalias [[META17:![0-9]+]]
-// CHECK-NEXT:    [[CALL3_I_I_I_I:%.*]] = tail call spir_func noundef zeroext i16 @_Z28__spirv_SubgroupShuffleINTELtj(i16 noundef zeroext [[TMP1]], i32 noundef 1) #[[ATTR6:[0-9]+]], !noalias [[META18:![0-9]+]]
+// CHECK-NEXT:    [[TMP1:%.*]] = load i16, ptr [[ARRAYIDX_I_I_I]], align 2, !tbaa [[TBAA17:![0-9]+]], !noalias [[META19:![0-9]+]]
+// CHECK-NEXT:    [[CALL3_I_I_I_I:%.*]] = tail call spir_func noundef zeroext i16 @_Z28__spirv_SubgroupShuffleINTELtj(i16 noundef zeroext [[TMP1]], i32 noundef 1) #[[ATTR6:[0-9]+]], !noalias [[META20:![0-9]+]]
 // CHECK-NEXT:    [[ARRAYIDX_I12_I_I:%.*]] = getelementptr inbounds %"class.sycl::_V1::ext::oneapi::bfloat16", ptr [[REF_TMP]], i64 [[IDXPROM_I_I_I]]
-// CHECK-NEXT:    store i16 [[CALL3_I_I_I_I]], ptr [[ARRAYIDX_I12_I_I]], align 2, !tbaa [[TBAA13]], !alias.scope [[META17]]
+// CHECK-NEXT:    store i16 [[CALL3_I_I_I_I]], ptr [[ARRAYIDX_I12_I_I]], align 2, !tbaa [[TBAA17]], !alias.scope [[META19]]
 // CHECK-NEXT:    [[INC_I_I]] = add nuw nsw i32 [[S_0_I_I]], 1
-// CHECK-NEXT:    br label [[FOR_COND_I_I]], !llvm.loop [[LOOP21:![0-9]+]]
+// CHECK-NEXT:    br label [[FOR_COND_I_I]], !llvm.loop [[LOOP23:![0-9]+]]
 // CHECK:       _ZN4sycl3_V117select_from_groupINS0_9sub_groupENS0_3vecINS0_3ext6oneapi8bfloat16ELi4EEEEENSt9enable_ifIXaaoosr3stdE9is_same_vINSt5decayIT_E4typeES2_Esr4sycl3ext6oneapi12experimentalE27is_user_constructed_group_vISC_Eoosr3stdE23is_trivially_copyable_vIT0_Esr6detail6is_vecISD_EE5valueESD_E4typeESA_SD_NSA_7id_typeE.exit:
-// CHECK-NEXT:    call void @llvm.lifetime.end.p0(ptr nonnull [[AGG_TMP14_I]]), !noalias [[META7]]
+// CHECK-NEXT:    call void @llvm.lifetime.end.p0(ptr nonnull [[AGG_TMP14_I]]), !noalias [[META11]]
 // CHECK-NEXT:    [[TMP2:%.*]] = load i64, ptr [[REF_TMP]], align 8
 // CHECK-NEXT:    store i64 [[TMP2]], ptr addrspace(4) [[ARRAYIDX]], align 8
 // CHECK-NEXT:    call void @llvm.lifetime.end.p0(ptr nonnull [[REF_TMP]]) #[[ATTR5]]
@@ -50,37 +50,39 @@ SYCL_EXTERNAL void test_shuffle1(sycl::sub_group &sg, vec<bfloat16, 4> *buf,
 // CHECK-NEXT:    [[AGG_TMP14_I:%.*]] = alloca %"class.sycl::_V1::marray", align 8
 // CHECK-NEXT:    [[REF_TMP:%.*]] = alloca %"class.sycl::_V1::marray", align 2
 // CHECK-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds nuw %"class.sycl::_V1::marray", ptr addrspace(4) [[BUF:%.*]], i64 [[ID:%.*]]
-// CHECK-NEXT:    [[TMP0:%.*]] = load i64, ptr addrspace(4) [[ARRAYIDX]], align 2, !tbaa [[TBAA23:![0-9]+]]
+// CHECK-NEXT:    [[TMP0:%.*]] = load i64, ptr addrspace(4) [[ARRAYIDX]], align 2, !tbaa [[TBAA25:![0-9]+]]
 // CHECK-NEXT:    call void @llvm.lifetime.start.p0(ptr nonnull [[REF_TMP]]) #[[ATTR5]]
-// CHECK-NEXT:    tail call void @llvm.experimental.noalias.scope.decl(metadata [[META24:![0-9]+]])
-// CHECK-NEXT:    tail call void @llvm.experimental.noalias.scope.decl(metadata [[META27:![0-9]+]])
-// CHECK-NEXT:    call void @llvm.lifetime.start.p0(ptr nonnull [[AGG_TMP14_I]]), !noalias [[META24]]
-// CHECK-NEXT:    store i64 [[TMP0]], ptr [[AGG_TMP14_I]], align 8, !noalias [[META24]]
+// CHECK-NEXT:    tail call void @llvm.experimental.noalias.scope.decl(metadata [[META26:![0-9]+]])
+// CHECK-NEXT:    tail call void @llvm.experimental.noalias.scope.decl(metadata [[META29:![0-9]+]])
+// CHECK-NEXT:    call void @llvm.lifetime.start.p0(ptr nonnull [[AGG_TMP14_I]]), !noalias [[META26]]
+// CHECK-NEXT:    store i64 [[TMP0]], ptr [[AGG_TMP14_I]], align 8, !noalias [[META26]]
 // CHECK-NEXT:    br label [[ARRAYINIT_BODY_I_I_I:%.*]]
 // CHECK:       arrayinit.body.i.i.i:
 // CHECK-NEXT:    [[ARRAYINIT_CUR_IDX_I_I_I:%.*]] = phi i64 [ 0, [[ENTRY:%.*]] ], [ [[ARRAYINIT_CUR_ADD_I_I_I:%.*]], [[ARRAYINIT_BODY_I_I_I]] ]
 // CHECK-NEXT:    [[ARRAYINIT_CUR_PTR_I_I_I:%.*]] = getelementptr inbounds i8, ptr [[REF_TMP]], i64 [[ARRAYINIT_CUR_IDX_I_I_I]]
-// CHECK-NEXT:    store i16 0, ptr [[ARRAYINIT_CUR_PTR_I_I_I]], align 2, !alias.scope [[META30:![0-9]+]]
+// CHECK-NEXT:    store i16 0, ptr [[ARRAYINIT_CUR_PTR_I_I_I]], align 2, !alias.scope [[META32:![0-9]+]]
 // CHECK-NEXT:    [[ARRAYINIT_CUR_ADD_I_I_I]] = add nuw nsw i64 [[ARRAYINIT_CUR_IDX_I_I_I]], 2
 // CHECK-NEXT:    [[ARRAYINIT_DONE_I_I_I:%.*]] = icmp eq i64 [[ARRAYINIT_CUR_ADD_I_I_I]], 8
-// CHECK-NEXT:    br i1 [[ARRAYINIT_DONE_I_I_I]], label [[FOR_COND_I_I:%.*]], label [[ARRAYINIT_BODY_I_I_I]]
+// CHECK-NEXT:    br i1 [[ARRAYINIT_DONE_I_I_I]], label [[FOR_COND_PREHEADER_I_I:%.*]], label [[ARRAYINIT_BODY_I_I_I]]
+// CHECK:       for.cond.preheader.i.i:
+// CHECK-NEXT:    br label [[FOR_COND_I_I:%.*]]
 // CHECK:       for.cond.i.i:
-// CHECK-NEXT:    [[S_0_I_I:%.*]] = phi i32 [ [[INC_I_I:%.*]], [[FOR_BODY_I_I:%.*]] ], [ 0, [[ARRAYINIT_BODY_I_I_I]] ]
+// CHECK-NEXT:    [[S_0_I_I:%.*]] = phi i32 [ [[INC_I_I:%.*]], [[FOR_BODY_I_I:%.*]] ], [ 0, [[FOR_COND_PREHEADER_I_I]] ]
 // CHECK-NEXT:    [[CMP_I_I:%.*]] = icmp samesign ult i32 [[S_0_I_I]], 4
 // CHECK-NEXT:    br i1 [[CMP_I_I]], label [[FOR_BODY_I_I]], label [[_ZN4SYCL3_V117SELECT_FROM_GROUPINS0_9SUB_GROUPENS0_6MARRAYINS0_3EXT6ONEAPI8BFLOAT16ELM4EEEEENST9ENABLE_IFIXAAOOSR3STDE9IS_SAME_VINST5DECAYIT_E4TYPEES2_ESR4SYCL3EXT6ONEAPI12EXPERIMENTALE27IS_USER_CONSTRUCTED_GROUP_VISC_EOOSR3STDE23IS_TRIVIALLY_COPYABLE_VIT0_ESR6DETAIL6IS_VECISD_EE5VALUEESD_E4TYPEESA_SD_NSA_7ID_TYPEE_EXIT:%.*]]
 // CHECK:       for.body.i.i:
 // CHECK-NEXT:    [[CONV_I_I:%.*]] = zext nneg i32 [[S_0_I_I]] to i64
 // CHECK-NEXT:    [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds %"class.sycl::_V1::ext::oneapi::bfloat16", ptr [[AGG_TMP14_I]], i64 [[CONV_I_I]]
-// CHECK-NEXT:    [[TMP1:%.*]] = load i16, ptr [[ARRAYIDX_I_I_I]], align 2, !tbaa [[TBAA13]], !noalias [[META30]]
-// CHECK-NEXT:    [[CALL3_I_I_I_I:%.*]] = tail call spir_func noundef zeroext i16 @_Z28__spirv_SubgroupShuffleINTELtj(i16 noundef zeroext [[TMP1]], i32 noundef 1) #[[ATTR6]], !noalias [[META31:![0-9]+]]
+// CHECK-NEXT:    [[TMP1:%.*]] = load i16, ptr [[ARRAYIDX_I_I_I]], align 2, !tbaa [[TBAA17]], !noalias [[META32]]
+// CHECK-NEXT:    [[CALL3_I_I_I_I:%.*]] = tail call spir_func noundef zeroext i16 @_Z28__spirv_SubgroupShuffleINTELtj(i16 noundef zeroext [[TMP1]], i32 noundef 1) #[[ATTR6]], !noalias [[META33:![0-9]+]]
 // CHECK-NEXT:    [[ARRAYIDX_I13_I_I:%.*]] = getelementptr inbounds %"class.sycl::_V1::ext::oneapi::bfloat16", ptr [[REF_TMP]], i64 [[CONV_I_I]]
-// CHECK-NEXT:    store i16 [[CALL3_I_I_I_I]], ptr [[ARRAYIDX_I13_I_I]], align 2, !tbaa [[TBAA13]], !alias.scope [[META30]]
+// CHECK-NEXT:    store i16 [[CALL3_I_I_I_I]], ptr [[ARRAYIDX_I13_I_I]], align 2, !tbaa [[TBAA17]], !alias.scope [[META32]]
 // CHECK-NEXT:    [[INC_I_I]] = add nuw nsw i32 [[S_0_I_I]], 1
-// CHECK-NEXT:    br label [[FOR_COND_I_I]], !llvm.loop [[LOOP34:![0-9]+]]
+// CHECK-NEXT:    br label [[FOR_COND_I_I]], !llvm.loop [[LOOP36:![0-9]+]]
 // CHECK:       _ZN4sycl3_V117select_from_groupINS0_9sub_groupENS0_6marrayINS0_3ext6oneapi8bfloat16ELm4EEEEENSt9enable_ifIXaaoosr3stdE9is_same_vINSt5decayIT_E4typeES2_Esr4sycl3ext6oneapi12experimentalE27is_user_constructed_group_vISC_Eoosr3stdE23is_trivially_copyable_vIT0_Esr6detail6is_vecISD_EE5valueESD_E4typeESA_SD_NSA_7id_typeE.exit:
-// CHECK-NEXT:    call void @llvm.lifetime.end.p0(ptr nonnull [[AGG_TMP14_I]]), !noalias [[META24]]
-// CHECK-NEXT:    [[TMP2:%.*]] = load i64, ptr [[REF_TMP]], align 2, !tbaa [[TBAA23]]
-// CHECK-NEXT:    store i64 [[TMP2]], ptr addrspace(4) [[ARRAYIDX]], align 2, !tbaa [[TBAA23]]
+// CHECK-NEXT:    call void @llvm.lifetime.end.p0(ptr nonnull [[AGG_TMP14_I]]), !noalias [[META26]]
+// CHECK-NEXT:    [[TMP2:%.*]] = load i64, ptr [[REF_TMP]], align 2, !tbaa [[TBAA25]]
+// CHECK-NEXT:    store i64 [[TMP2]], ptr addrspace(4) [[ARRAYIDX]], align 2, !tbaa [[TBAA25]]
 // CHECK-NEXT:    call void @llvm.lifetime.end.p0(ptr nonnull [[REF_TMP]]) #[[ATTR5]]
 // CHECK-NEXT:    ret void
 //
@@ -96,34 +98,36 @@ SYCL_EXTERNAL void test_shuffle2(sycl::sub_group &sg, marray<bfloat16, 4> *buf,
 // CHECK-NEXT:    [[REF_TMP:%.*]] = alloca %"class.sycl::_V1::marray.32", align 2
 // CHECK-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds nuw %"class.sycl::_V1::marray.32", ptr addrspace(4) [[BUF:%.*]], i64 [[ID:%.*]]
 // CHECK-NEXT:    call void @llvm.lifetime.start.p0(ptr nonnull [[REF_TMP]]) #[[ATTR5]]
-// CHECK-NEXT:    call void @llvm.lifetime.start.p0(ptr nonnull [[AGG_TMP14_I]]), !noalias [[META35:![0-9]+]]
+// CHECK-NEXT:    call void @llvm.lifetime.start.p0(ptr nonnull [[AGG_TMP14_I]]), !noalias [[META37:![0-9]+]]
 // CHECK-NEXT:    call void @llvm.memcpy.p0.p4.i64(ptr noundef nonnull align 8 dereferenceable(10) [[AGG_TMP14_I]], ptr addrspace(4) noundef align 2 dereferenceable(10) [[ARRAYIDX]], i64 10, i1 false)
-// CHECK-NEXT:    tail call void @llvm.experimental.noalias.scope.decl(metadata [[META35]])
-// CHECK-NEXT:    tail call void @llvm.experimental.noalias.scope.decl(metadata [[META38:![0-9]+]])
+// CHECK-NEXT:    tail call void @llvm.experimental.noalias.scope.decl(metadata [[META37]])
+// CHECK-NEXT:    tail call void @llvm.experimental.noalias.scope.decl(metadata [[META40:![0-9]+]])
 // CHECK-NEXT:    br label [[ARRAYINIT_BODY_I_I_I:%.*]]
 // CHECK:       arrayinit.body.i.i.i:
 // CHECK-NEXT:    [[ARRAYINIT_CUR_IDX_I_I_I:%.*]] = phi i64 [ 0, [[ENTRY:%.*]] ], [ [[ARRAYINIT_CUR_ADD_I_I_I:%.*]], [[ARRAYINIT_BODY_I_I_I]] ]
 // CHECK-NEXT:    [[ARRAYINIT_CUR_PTR_I_I_I:%.*]] = getelementptr inbounds i8, ptr [[REF_TMP]], i64 [[ARRAYINIT_CUR_IDX_I_I_I]]
-// CHECK-NEXT:    store i16 0, ptr [[ARRAYINIT_CUR_PTR_I_I_I]], align 2, !alias.scope [[META41:![0-9]+]]
+// CHECK-NEXT:    store i16 0, ptr [[ARRAYINIT_CUR_PTR_I_I_I]], align 2, !alias.scope [[META43:![0-9]+]]
 // CHECK-NEXT:    [[ARRAYINIT_CUR_ADD_I_I_I]] = add nuw nsw i64 [[ARRAYINIT_CUR_IDX_I_I_I]], 2
 // CHECK-NEXT:    [[ARRAYINIT_DONE_I_I_I:%.*]] = icmp eq i64 [[ARRAYINIT_CUR_ADD_I_I_I]], 10
-// CHECK-NEXT:    br i1 [[ARRAYINIT_DONE_I_I_I]], label [[FOR_COND_I_I:%.*]], label [[ARRAYINIT_BODY_I_I_I]]
+// CHECK-NEXT:    br i1 [[ARRAYINIT_DONE_I_I_I]], label [[FOR_COND_PREHEADER_I_I:%.*]], label [[ARRAYINIT_BODY_I_I_I]]
+// CHECK:       for.cond.preheader.i.i:
+// CHECK-NEXT:    br label [[FOR_COND_I_I:%.*]]
 // CHECK:       for.cond.i.i:
-// CHECK-NEXT:    [[S_0_I_I:%.*]] = phi i32 [ [[INC_I_I:%.*]], [[FOR_BODY_I_I:%.*]] ], [ 0, [[ARRAYINIT_BODY_I_I_I]] ]
+// CHECK-NEXT:    [[S_0_I_I:%.*]] = phi i32 [ [[INC_I_I:%.*]], [[FOR_BODY_I_I:%.*]] ], [ 0, [[FOR_COND_PREHEADER_I_I]] ]
 // CHECK-NEXT:    [[CMP_I_I:%.*]] = icmp samesign ult i32 [[S_0_I_I]], 5
 // CHECK-NEXT:    br i1 [[CMP_I_I]], label [[FOR_BODY_I_I]], label [[_ZN4SYCL3_V117SELECT_FROM_GROUPINS0_9SUB_GROUPENS0_6MARRAYINS0_3EXT6ONEAPI8BFLOAT16ELM5EEEEENST9ENABLE_IFIXAAOOSR3STDE9IS_SAME_VINST5DECAYIT_E4TYPEES2_ESR4SYCL3EXT6ONEAPI12EXPERIMENTALE27IS_USER_CONSTRUCTED_GROUP_VISC_EOOSR3STDE23IS_TRIVIALLY_COPYABLE_VIT0_ESR6DETAIL6IS_VECISD_EE5VALUEESD_E4TYPEESA_SD_NSA_7ID_TYPEE_EXIT:%.*]]
 // CHECK:       for.body.i.i:
 // CHECK-NEXT:    [[CONV_I_I:%.*]] = zext nneg i32 [[S_0_I_I]] to i64
 // CHECK-NEXT:    [[ARRAYIDX_I_I_I:%.*]] = getelementptr inbounds %"class.sycl::_V1::ext::oneapi::bfloat16", ptr [[AGG_TMP14_I]], i64 [[CONV_I_I]]
-// CHECK-NEXT:    [[TMP0:%.*]] = load i16, ptr [[ARRAYIDX_I_I_I]], align 2, !tbaa [[TBAA13]], !noalias [[META41]]
-// CHECK-NEXT:    [[CALL3_I_I_I_I:%.*]] = tail call spir_func noundef zeroext i16 @_Z28__spirv_SubgroupShuffleINTELtj(i16 noundef zeroext [[TMP0]], i32 noundef 1) #[[ATTR6]], !noalias [[META42:![0-9]+]]
+// CHECK-NEXT:    [[TMP0:%.*]] = load i16, ptr [[ARRAYIDX_I_I_I]], align 2, !tbaa [[TBAA17]], !noalias [[META43]]
+// CHECK-NEXT:    [[CALL3_I_I_I_I:%.*]] = tail call spir_func noundef zeroext i16 @_Z28__spirv_SubgroupShuffleINTELtj(i16 noundef zeroext [[TMP0]], i32 noundef 1) #[[ATTR6]], !noalias [[META44:![0-9]+]]
 // CHECK-NEXT:    [[ARRAYIDX_I13_I_I:%.*]] = getelementptr inbounds %"class.sycl::_V1::ext::oneapi::bfloat16", ptr [[REF_TMP]], i64 [[CONV_I_I]]
-// CHECK-NEXT:    store i16 [[CALL3_I_I_I_I]], ptr [[ARRAYIDX_I13_I_I]], align 2, !tbaa [[TBAA13]], !alias.scope [[META41]]
+// CHECK-NEXT:    store i16 [[CALL3_I_I_I_I]], ptr [[ARRAYIDX_I13_I_I]], align 2, !tbaa [[TBAA17]], !alias.scope [[META43]]
 // CHECK-NEXT:    [[INC_I_I]] = add nuw nsw i32 [[S_0_I_I]], 1
-// CHECK-NEXT:    br label [[FOR_COND_I_I]], !llvm.loop [[LOOP45:![0-9]+]]
+// CHECK-NEXT:    br label [[FOR_COND_I_I]], !llvm.loop [[LOOP47:![0-9]+]]
 // CHECK:       _ZN4sycl3_V117select_from_groupINS0_9sub_groupENS0_6marrayINS0_3ext6oneapi8bfloat16ELm5EEEEENSt9enable_ifIXaaoosr3stdE9is_same_vINSt5decayIT_E4typeES2_Esr4sycl3ext6oneapi12experimentalE27is_user_constructed_group_vISC_Eoosr3stdE23is_trivially_copyable_vIT0_Esr6detail6is_vecISD_EE5valueESD_E4typeESA_SD_NSA_7id_typeE.exit:
-// CHECK-NEXT:    call void @llvm.lifetime.end.p0(ptr nonnull [[AGG_TMP14_I]]), !noalias [[META35]]
-// CHECK-NEXT:    call void @llvm.memcpy.p4.p0.i64(ptr addrspace(4) align 2 [[ARRAYIDX]], ptr align 2 [[REF_TMP]], i64 10, i1 false), !tbaa.struct [[TBAA_STRUCT46:![0-9]+]]
+// CHECK-NEXT:    call void @llvm.lifetime.end.p0(ptr nonnull [[AGG_TMP14_I]]), !noalias [[META37]]
+// CHECK-NEXT:    call void @llvm.memcpy.p4.p0.i64(ptr addrspace(4) align 2 [[ARRAYIDX]], ptr align 2 [[REF_TMP]], i64 10, i1 false), !tbaa.struct [[TBAA_STRUCT48:![0-9]+]]
 // CHECK-NEXT:    call void @llvm.lifetime.end.p0(ptr nonnull [[REF_TMP]]) #[[ATTR5]]
 // CHECK-NEXT:    ret void
 //
@@ -136,11 +140,11 @@ SYCL_EXTERNAL void test_shuffle3(sycl::sub_group &sg, marray<bfloat16, 5> *buf,
 // CHECK-LABEL: @_Z13test_shuffle4RN4sycl3_V19sub_groupEPPim(
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds nuw ptr addrspace(4), ptr addrspace(4) [[BUF:%.*]], i64 [[ID:%.*]]
-// CHECK-NEXT:    [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[ARRAYIDX]], align 8, !tbaa [[TBAA47:![0-9]+]]
+// CHECK-NEXT:    [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[ARRAYIDX]], align 8, !tbaa [[TBAA49:![0-9]+]]
 // CHECK-NEXT:    [[TMP1:%.*]] = ptrtoint ptr addrspace(4) [[TMP0]] to i64
 // CHECK-NEXT:    [[CALL3_I_I_I:%.*]] = tail call spir_func noundef i64 @_Z28__spirv_SubgroupShuffleINTELmj(i64 noundef [[TMP1]], i32 noundef 1) #[[ATTR6]]
 // CHECK-NEXT:    [[TMP2:%.*]] = inttoptr i64 [[CALL3_I_I_I]] to ptr addrspace(4)
-// CHECK-NEXT:    store ptr addrspace(4) [[TMP2]], ptr addrspace(4) [[ARRAYIDX]], align 8, !tbaa [[TBAA47]]
+// CHECK-NEXT:    store ptr addrspace(4) [[TMP2]], ptr addrspace(4) [[ARRAYIDX]], align 8, !tbaa [[TBAA49]]
 // CHECK-NEXT:    ret void
 //
 SYCL_EXTERNAL void test_shuffle4(sycl::sub_group &sg, int **buf, size_t id) {
@@ -151,11 +155,11 @@ SYCL_EXTERNAL void test_shuffle4(sycl::sub_group &sg, int **buf, size_t id) {
 // CHECK-LABEL: @_Z13test_shuffle5RN4sycl3_V19sub_groupEPPVim(
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[ARRAYIDX:%.*]] = getelementptr inbounds nuw ptr addrspace(4), ptr addrspace(4) [[BUF:%.*]], i64 [[ID:%.*]]
-// CHECK-NEXT:    [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[ARRAYIDX]], align 8, !tbaa [[TBAA47]]
+// CHECK-NEXT:    [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[ARRAYIDX]], align 8, !tbaa [[TBAA49]]
 // CHECK-NEXT:    [[TMP1:%.*]] = ptrtoint ptr addrspace(4) [[TMP0]] to i64
 // CHECK-NEXT:    [[CALL3_I_I_I:%.*]] = tail call spir_func noundef i64 @_Z28__spirv_SubgroupShuffleINTELmj(i64 noundef [[TMP1]], i32 noundef 1) #[[ATTR6]]
 // CHECK-NEXT:    [[TMP2:%.*]] = inttoptr i64 [[CALL3_I_I_I]] to ptr addrspace(4)
-// CHECK-NEXT:    store ptr addrspace(4) [[TMP2]], ptr addrspace(4) [[ARRAYIDX]], align 8, !tbaa [[TBAA47]]
+// CHECK-NEXT:    store ptr addrspace(4) [[TMP2]], ptr addrspace(4) [[ARRAYIDX]], align 8, !tbaa [[TBAA49]]
 // CHECK-NEXT:    ret void
 //
 SYCL_EXTERNAL void test_shuffle5(sycl::sub_group &sg, volatile int **buf,