diff --git a/modules/compiler/source/base/source/module.cpp b/modules/compiler/source/base/source/module.cpp index 53b5f4dad..340535053 100644 --- a/modules/compiler/source/base/source/module.cpp +++ b/modules/compiler/source/base/source/module.cpp @@ -1438,6 +1438,9 @@ std::unique_ptr BaseModule::compileOpenCLCToIR( #endif // CA_ENABLE_DEBUG_SUPPORT instance.createDiagnostics( +#if LLVM_VERSION_GREATER_EQUAL(20, 0) + *llvm::vfs::getRealFileSystem(), +#endif new FrontendDiagnosticPrinter(*this, &instance.getDiagnosticOpts())); // Write a copy of the kernel source out to disk and update the debug info diff --git a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/broadcast_vector.ll b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/broadcast_vector.ll index 519bb696c..5e0520a55 100644 --- a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/broadcast_vector.ll +++ b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/broadcast_vector.ll @@ -100,7 +100,7 @@ entry: ; CHECK-NEXT: entry: ; CHECK-NEXT: [[CALL:%.*]] = tail call i64 @__mux_get_global_id(i32 0) ; CHECK-NEXT: [[ARRAYIDX3:%.*]] = getelementptr <4 x float>, ptr addrspace(1) [[OUT:%.*]], i64 [[CALL]] -; CHECK-NEXT: store shufflevector ( insertelement ( {{(undef|poison)}}, float 0x7FF8000020000000, {{(i32|i64)}} 0), {{(undef|poison)}}, zeroinitializer), ptr addrspace(1) [[ARRAYIDX3]], align 16 +; CHECK-NEXT: store {{shufflevector \( insertelement \( (undef|poison), float 0x7FF8000020000000, (i32|i64) 0\), (undef|poison), zeroinitializer\)|splat \(float 0x7FF8000020000000\)}}, ptr addrspace(1) [[ARRAYIDX3]], align 16 ; CHECK-NEXT: ret void ; CHECK-LABEL: @__vecz_nxv4_vector_broadcast( @@ -108,10 +108,10 @@ entry: ; CHECK-NEXT: [[FIXLEN_ALLOC:%.*]] = alloca <4 x float>, align 16 ; CHECK-NEXT: store <4 x float> [[ADDEND:%.*]], ptr [[FIXLEN_ALLOC]], align 16 ; CHECK-NEXT: [[IDX0:%.*]] = call @llvm.{{(experimental\.)?}}stepvector.nxv16i32() -; CHECK-NEXT: [[IDX1:%.*]] = and [[IDX0]], shufflevector ( insertelement ( {{(undef|poison)}}, i32 3, {{(i32|i64)}} 0), {{(undef|poison)}}, zeroinitializer) +; CHECK-NEXT: [[IDX1:%.*]] = and [[IDX0]], {{shufflevector \( insertelement \( (undef|poison), i32 3, (i32|i64) 0\), (undef|poison), zeroinitializer\)|splat \(i32 3\)}} ; CHECK-NEXT: [[TMP0:%.*]] = {{s|z}}ext{{( nneg)?}} [[IDX1]] to ; CHECK-NEXT: [[VEC_ALLOC:%.*]] = getelementptr inbounds float, ptr [[FIXLEN_ALLOC]], [[TMP0]] -; CHECK-NEXT: [[TMP1:%.*]] = call @llvm.masked.gather.nxv16f32.nxv16p0( [[VEC_ALLOC]], i32 4, shufflevector ( insertelement ( poison, i1 true, {{(i32|i64)}} 0), poison, zeroinitializer), undef) +; CHECK-NEXT: [[TMP1:%.*]] = call @llvm.masked.gather.nxv16f32.nxv16p0( [[VEC_ALLOC]], i32 4, {{shufflevector \( insertelement \( poison, i1 true, (i32|i64) 0\), poison, zeroinitializer\)|splat \(i1 true\)}}, undef) ; CHECK-NEXT: [[CALL:%.*]] = tail call i64 @__mux_get_global_id(i32 0) ; CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr <4 x float>, ptr addrspace(1) [[IN:%.*]], i64 [[CALL]] ; CHECK-NEXT: [[TMP3:%.*]] = load , ptr addrspace(1) [[ARRAYIDX]], align 16 @@ -125,13 +125,13 @@ entry: ; CHECK-NEXT: [[CALL:%.*]] = tail call i64 @__mux_get_global_id(i32 0) ; CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr <4 x float>, ptr addrspace(1) [[IN:%.*]], i64 [[CALL]] ; CHECK-NEXT: [[TMP1:%.*]] = load , ptr addrspace(1) [[ARRAYIDX]], align 16 -; CHECK-NEXT: [[AND1_I_I_I1_I1:%.*]] = and [[TMP1]], shufflevector ( insertelement ( {{(undef|poison)}}, i32 2139095040, {{i32|i64}} 0), {{(undef|poison)}}, zeroinitializer) -; CHECK-NEXT: [[CMP_I_I_I2_I2:%.*]] = icmp ne [[AND1_I_I_I1_I1]], shufflevector ( insertelement ( {{(undef|poison)}}, i32 2139095040, {{i32|i64}} 0), {{(undef|poison)}}, zeroinitializer) -; CHECK-NEXT: [[AND2_I_I_I3_I3:%.*]] = and [[TMP1]], shufflevector ( insertelement ( {{(undef|poison)}}, i32 8388607, {{i32|i64}} 0), {{(undef|poison)}}, zeroinitializer) +; CHECK-NEXT: [[AND1_I_I_I1_I1:%.*]] = and [[TMP1]], {{shufflevector \( insertelement \( (undef|poison), i32 2139095040, (i32|i64) 0\), (undef|poison), zeroinitializer\)|splat \(i32 2139095040\)}} +; CHECK-NEXT: [[CMP_I_I_I2_I2:%.*]] = icmp ne [[AND1_I_I_I1_I1]], {{shufflevector \( insertelement \( (undef|poison), i32 2139095040, (i32|i64) 0\), (undef|poison), zeroinitializer\)|splat \(i32 2139095040\)}} +; CHECK-NEXT: [[AND2_I_I_I3_I3:%.*]] = and [[TMP1]], {{shufflevector \( insertelement \( (undef|poison), i32 8388607, (i32|i64) 0\), (undef|poison), zeroinitializer\)|splat \(i32 8388607\)}} ; CHECK-NEXT: [[CMP3_I_I_I4_I4:%.*]] = icmp eq [[AND2_I_I_I3_I3]], zeroinitializer ; CHECK-NEXT: [[TMP2:%.*]] = or [[CMP_I_I_I2_I2]], [[CMP3_I_I_I4_I4]] ; CHECK-NEXT: [[TMP3:%.*]] = bitcast [[TMP1]] to -; CHECK-NEXT: [[TMP4:%.*]] = select [[TMP2]], [[TMP3]], shufflevector ( insertelement ( {{(undef|poison)}}, float 0x7FF0000020000000, {{i32|i64}} 0), {{(undef|poison)}}, zeroinitializer) +; CHECK-NEXT: [[TMP4:%.*]] = select [[TMP2]], [[TMP3]], {{shufflevector \( insertelement \( (undef|poison), float 0x7FF0000020000000, (i32|i64) 0\), (undef|poison), zeroinitializer\)|splat \(float 0x7FF0000020000000\)}} ; CHECK-NEXT: [[ARRAYIDX3:%.*]] = getelementptr <4 x float>, ptr addrspace(1) [[OUT:%.*]], i64 [[CALL]] ; CHECK-NEXT: store [[TMP4]], ptr addrspace(1) [[ARRAYIDX3]], align 16 ; CHECK-NEXT: ret void @@ -144,10 +144,10 @@ entry: ; CHECK-NEXT: [[FIXLEN_ALLOC1:%.*]] = alloca <4 x float>, align 16 ; CHECK-NEXT: store <4 x float> [[ADDEND:%.*]], ptr [[FIXLEN_ALLOC1]], align 16 ; CHECK-NEXT: [[IDX03:%.*]] = call @llvm.{{(experimental\.)?}}stepvector.nxv16i32() -; CHECK-NEXT: [[IDX14:%.*]] = and [[IDX03]], shufflevector ( insertelement ( {{(undef|poison)}}, i32 3, {{i32|i64}} 0), {{(undef|poison)}}, zeroinitializer) +; CHECK-NEXT: [[IDX14:%.*]] = and [[IDX03]], {{shufflevector \( insertelement \( (undef|poison), i32 3, (i32|i64) 0\), (undef|poison), zeroinitializer\)|splat \(i32 3\)}} ; CHECK-NEXT: [[TMP0:%.*]] = {{s|z}}ext{{( nneg)?}} [[IDX14]] to ; CHECK-NEXT: [[VEC_ALLOC5:%.*]] = getelementptr inbounds float, ptr [[FIXLEN_ALLOC1]], [[TMP0]] -; CHECK-NEXT: [[TMP1:%.*]] = call @llvm.masked.gather.nxv16f32.nxv16p0( [[VEC_ALLOC5]], i32 4, shufflevector ( insertelement ( poison, i1 true, {{i32|i64}} 0), poison, zeroinitializer), {{(undef|poison)}}) +; CHECK-NEXT: [[TMP1:%.*]] = call @llvm.masked.gather.nxv16f32.nxv16p0( [[VEC_ALLOC5]], i32 4, {{shufflevector \( insertelement \( poison, i1 true, (i32|i64) 0\), poison, zeroinitializer\)|splat \(i1 true\)}}, {{(undef|poison)}}) ; CHECK-NEXT: [[CALL:%.*]] = tail call i64 @__mux_get_global_id(i32 0) ; CHECK-NEXT: store <4 x i32> zeroinitializer, ptr [[EXISTING_ALLOC]], align 16 ; CHECK-NEXT: store i32 1, ptr [[EXISTING_ALLOC]], align @@ -155,7 +155,7 @@ entry: ; CHECK-NEXT: store <4 x i32> [[V]], ptr [[FIXLEN_ALLOC]], align 16 ; CHECK-NEXT: [[TMP2:%.*]] = {{s|z}}ext{{( nneg)?}} [[IDX14]] to ; CHECK-NEXT: [[VEC_ALLOC:%.*]] = getelementptr inbounds i32, ptr [[FIXLEN_ALLOC]], [[TMP2]] -; CHECK-NEXT: [[TMP3:%.*]] = call @llvm.masked.gather.nxv16i32.nxv16p0( [[VEC_ALLOC]], i32 4, shufflevector ( insertelement ( poison, i1 true, {{i32|i64}} 0), poison, zeroinitializer), {{(undef|poison)}}) +; CHECK-NEXT: [[TMP3:%.*]] = call @llvm.masked.gather.nxv16i32.nxv16p0( [[VEC_ALLOC]], i32 4, {{shufflevector \( insertelement \( poison, i1 true, (i32|i64) 0\), poison, zeroinitializer\)|splat \(i1 true\)}}, {{(undef|poison)}}) ; CHECK-NEXT: [[ARRAYIDX4:%.*]] = getelementptr <4 x i32>, ptr addrspace(1) [[OUT2:%.*]], i64 [[CALL]] ; CHECK-NEXT: store [[TMP3]], ptr addrspace(1) [[ARRAYIDX4]], align 16 ; CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr <4 x float>, ptr addrspace(1) [[IN:%.*]], i64 [[CALL]] @@ -169,11 +169,11 @@ entry: ; CHECK-NEXT: entry: ; CHECK: [[FIXLEN_MASK_ALLOC:%.*]] = alloca <4 x i8>, align 4 ; CHECK: [[IDX0:%.*]] = call @llvm.{{(experimental\.)?}}stepvector.nxv16i32() -; CHECK: [[IDX1:%.*]] = and [[IDX0]], shufflevector ( insertelement ( {{(undef|poison)}}, i32 3, {{i32|i64}} 0), {{(undef|poison)}}, zeroinitializer) +; CHECK: [[IDX1:%.*]] = and [[IDX0]], {{shufflevector \( insertelement \( (undef|poison), i32 3, (i32|i64) 0\), (undef|poison), zeroinitializer\)|splat \(i32 3\)}} ; CHECK: [[SEXT:%.*]] = sext <4 x i1> [[INPUT:%.*]] to <4 x i8> ; CHECK: store <4 x i8> [[SEXT]], ptr [[FIXLEN_MASK_ALLOC]], align 4 ; CHECK: [[TMP0:%.*]] = {{s|z}}ext{{( nneg)?}} [[IDX1]] to ; CHECK: [[VEC_ALLOC:%.*]] = getelementptr inbounds i8, ptr [[FIXLEN_MASK_ALLOC]], [[TMP0]] -; CHECK: [[TMP1:%.*]] = call @llvm.masked.gather.nxv16i8.nxv16p0( [[VEC_ALLOC]], i32 1, shufflevector ( insertelement ( poison, i1 true, {{i32|i64}} 0), poison, zeroinitializer), {{(undef|poison)}}) +; CHECK: [[TMP1:%.*]] = call @llvm.masked.gather.nxv16i8.nxv16p0( [[VEC_ALLOC]], i32 1, {{shufflevector \( insertelement \( poison, i1 true, (i32|i64) 0\), poison, zeroinitializer\)|splat \(i1 true\)}}, {{(undef|poison)}}) ; CHECK: [[BMASK:%.*]] = trunc [[TMP1]] to ; CHECK: {{.*}} = and {{.*}}, [[BMASK]] diff --git a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/cmpxchg.ll b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/cmpxchg.ll index 85b4c865d..bfa7f6933 100644 --- a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/cmpxchg.ll +++ b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/cmpxchg.ll @@ -29,9 +29,9 @@ entry: ; Test that this cmpxchg is packetized by generating a call to an all-true masked version. ; CHECK: [[A0:%.*]] = call { , } @__vecz_b_nxv4_masked_cmpxchg_align4_acquire_monotonic_1_u9nxv4u3ptru5nxv4ju5nxv4ju5nxv4b( ; CHECK-SAME: [[SPLAT_PTR]], -; CHECK-SAME: shufflevector ( insertelement ( poison, i32 1, i64 0), poison, zeroinitializer) -; CHECK-SAME: shufflevector ( insertelement ( poison, i32 2, i64 0), poison, zeroinitializer) -; CHECK-SAME: shufflevector ( insertelement ( poison, i1 true, i64 0), poison, zeroinitializer) +; CHECK-SAME: {{shufflevector \( insertelement \( poison, i32 1, i64 0\), poison, zeroinitializer\)|splat \(i32 1\)}} +; CHECK-SAME: {{shufflevector \( insertelement \( poison, i32 2, i64 0\), poison, zeroinitializer\)|splat \(i32 2\)}} +; CHECK-SAME: {{shufflevector \( insertelement \( poison, i1 true, i64 0\), poison, zeroinitializer\)|splat \(i1 true\)}} %old0 = cmpxchg ptr %p, i32 1, i32 2 acquire monotonic ; CHECK: [[EXT0:%.*]] = extractvalue { , } [[A0]], 0 %val0 = extractvalue { i32, i1 } %old0, 0 diff --git a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/define_interleaved_store.ll b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/define_interleaved_store.ll index 6c0be7573..497e9a54c 100644 --- a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/define_interleaved_store.ll +++ b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/define_interleaved_store.ll @@ -59,7 +59,7 @@ declare <4 x double> @llvm.fmuladd.v4f64(<4 x double>, <4 x double>, <4 x double ; CHECK: %BroadcastAddr.splatinsert = insertelement {{poison|undef}}, ptr addrspace(1) %1, {{i32|i64}} 0 ; CHECK: %BroadcastAddr.splat = shufflevector %BroadcastAddr.splatinsert, {{poison|undef}}, zeroinitializer ; CHECK: %2 = call @llvm.{{(experimental\.)?}}stepvector.nxv4i64() -; CHECK: %3 = mul shufflevector ( insertelement ( poison, i64 4, {{i32|i64}} 0), poison, zeroinitializer), %2 +; CHECK: %3 = mul {{shufflevector \( insertelement \( poison, i64 4, (i32|i64) 0\), poison, zeroinitializer\)|splat \(i64 4\)}}, %2 ; CHECK: %4 = getelementptr double, %BroadcastAddr.splat, %3 -; CHECK: call void @llvm.masked.scatter.nxv4f64.nxv4p1( %0, %4, i32{{( immarg)?}} 8, shufflevector ( insertelement ( poison, i1 true, {{i32|i64}} 0), poison, zeroinitializer)) +; CHECK: call void @llvm.masked.scatter.nxv4f64.nxv4p1( %0, %4, i32{{( immarg)?}} 8, {{shufflevector \( insertelement \( poison, i1 true, (i32|i64) 0\), poison, zeroinitializer\)|splat \(i1 true\)}}) ; CHECK: ret void diff --git a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/define_interleaved_store_as_masked.ll b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/define_interleaved_store_as_masked.ll index f6a350a65..9d8a46850 100644 --- a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/define_interleaved_store_as_masked.ll +++ b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/define_interleaved_store_as_masked.ll @@ -58,9 +58,9 @@ declare <4 x double> @llvm.fmuladd.v4f64(<4 x double>, <4 x double>, <4 x double ; CHECK: %BroadcastAddr.splatinsert = insertelement poison, ptr addrspace(1) %1, {{i32|i64}} 0 ; CHECK: %BroadcastAddr.splat = shufflevector %BroadcastAddr.splatinsert, poison, zeroinitializer ; CHECK: %2 = call @llvm.{{(experimental\.)?}}stepvector.nxv4i64() -; CHECK: %3 = mul shufflevector ( insertelement ( poison, i64 4, {{i32|i64}} 0), poison, zeroinitializer), %2 +; CHECK: %3 = mul {{shufflevector \( insertelement \( poison, i64 4, (i32|i64) 0\), poison, zeroinitializer\)|splat \(i64 4\)}}, %2 ; CHECK: %4 = getelementptr double, %BroadcastAddr.splat, %3 -; CHECK: call void @llvm.masked.scatter.nxv4f64.nxv4p1( %0, %4, i32 immarg 8, shufflevector ( insertelement ( poison, i1 true, {{i32|i64}} 0), poison, zeroinitializer)) +; CHECK: call void @llvm.masked.scatter.nxv4f64.nxv4p1( %0, %4, i32 immarg 8, {{shufflevector \( insertelement \( poison, i1 true, (i32|i64) 0\), poison, zeroinitializer\)|splat \(i1 true\)}}) ; CHECK: ret void ; CHECK: } diff --git a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/define_subgroup_scans.ll b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/define_subgroup_scans.ll index 12b2856ce..6a8a686d0 100644 --- a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/define_subgroup_scans.ll +++ b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/define_subgroup_scans.ll @@ -58,7 +58,7 @@ declare @__vecz_b_sub_group_scan_inclusive_add_u5nxv4j( %[[VEC]], {{(\*)|(ptr)}} %[[SHUFFLE_ALLOC]] ;------- there will be a bitcast here if pointers are typed ; CHECK: %[[INDEX:.+]] = getelementptr inbounds i32, [[PTRTY:(i32\*)|ptr]] %{{.+}}, %[[MASK]] -; CHECK: %[[SHUFFLE:.+]] = call @llvm.masked.gather.nxv4i32.nxv4p0{{(i32)?}}( %[[INDEX]], i32 4, shufflevector ( insertelement ( poison, i1 true, {{i32|i64}} 0), poison, zeroinitializer), undef) +; CHECK: %[[SHUFFLE:.+]] = call @llvm.masked.gather.nxv4i32.nxv4p0{{(i32)?}}( %[[INDEX]], i32 4, {{shufflevector \( insertelement \( poison, i1 true, (i32|i64) 0\), poison, zeroinitializer\)|splat \(i1 true\)}}, undef) ; CHECK: %[[ACCUM:.+]] = add %[[VEC]], %{{.+}} ; CHECK: %[[BIT:.+]] = and %[[MASKPHI]], %[[N_SPLAT]] @@ -93,7 +93,7 @@ declare @__vecz_b_sub_group_scan_exclusive_add_u5nxv4j( %[[VEC]], {{(\*)|(ptr)}} %[[SHUFFLE_ALLOC]] ;------- there will be a bitcast here if pointers are typed ; CHECK: %[[INDEX:.+]] = getelementptr inbounds i32, [[PTRTY:(i32\*)|ptr]] %{{.+}}, %[[MASK]] -; CHECK: %[[SHUFFLE:.+]] = call @llvm.masked.gather.nxv4i32.nxv4p0{{(i32)?}}( %[[INDEX]], i32 4, shufflevector ( insertelement ( poison, i1 true, {{i32|i64}} 0), poison, zeroinitializer), undef) +; CHECK: %[[SHUFFLE:.+]] = call @llvm.masked.gather.nxv4i32.nxv4p0{{(i32)?}}( %[[INDEX]], i32 4, {{shufflevector \( insertelement \( poison, i1 true, (i32|i64) 0\), poison, zeroinitializer\)|splat \(i1 true\)}}, undef) ; CHECK: %[[ACCUM:.+]] = add %[[VEC]], %{{.+}} ; CHECK: %[[BIT:.+]] = and %[[MASKPHI]], %[[N_SPLAT]] diff --git a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/extract_element.ll b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/extract_element.ll index 91c989df3..77ea2d3f7 100644 --- a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/extract_element.ll +++ b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/extract_element.ll @@ -114,8 +114,8 @@ entry: ; EE-UNI-VEC: [[T4:%.*]] = shufflevector [[T3]], poison, zeroinitializer ; EE-UNI-VEC: [[STEP:%.*]] = call @llvm.{{(experimental\.)?}}stepvector.nxv4i64() ; EE-UNI-VEC: [[T5:%.*]] = add [[T4]], [[STEP]] -; EE-UNI-VEC: [[MOD:%.*]] = and [[T5]], shufflevector ( insertelement ( {{(undef|poison)}}, i64 3, {{(i32|i64)}} 0), {{(undef|poison)}}, zeroinitializer) -; EE-UNI-VEC: [[T6:%.*]] = shl [[STEP]], shufflevector ( insertelement ( {{(undef|poison)}}, i64 2, {{(i32|i64)}} 0), {{(undef|poison)}}, zeroinitializer) +; EE-UNI-VEC: [[MOD:%.*]] = and [[T5]], {{shufflevector \( insertelement \( (undef|poison), i64 3, (i32|i64) 0\), (undef|poison), zeroinitializer\)|splat \(i64 3\)}} +; EE-UNI-VEC: [[T6:%.*]] = shl [[STEP]], {{shufflevector \( insertelement \( (undef|poison), i64 2, (i32|i64) 0\), (undef|poison), zeroinitializer\)|splat \(i64 2\)}} ; LLVM 16 deduces add/or equivalence and uses `or` instead. ; EE-UNI-VEC: [[T7:%.*]] = {{add|or}} {{(disjoint )?}} [[T6]], [[MOD]] @@ -128,10 +128,10 @@ entry: ; EE-INDICES: [[ALLOC:%.*]] = alloca , align 64 ; EE-INDICES: [[T0:%.*]] = getelementptr i32, ptr addrspace(1) %idxs, i64 %call ; EE-INDICES: [[T2:%.*]] = load , ptr addrspace(1) [[T0]], align 4 -; EE-INDICES: [[T3:%.*]] = and [[T2]], shufflevector ( insertelement ( {{(undef|poison)}}, i32 3, {{i32|i64}} 0), {{(undef|poison)}}, zeroinitializer) +; EE-INDICES: [[T3:%.*]] = and [[T2]], {{shufflevector \( insertelement \( (undef|poison), i32 3, (i32|i64) 0\), (undef|poison), zeroinitializer\)|splat \(i32 3\)}} ; EE-INDICES: store {{.*}}, ptr [[ALLOC]], align 64 ; EE-INDICES: [[STEP:%.*]] = call @llvm.{{(experimental\.)?}}stepvector.nxv4i32() -; EE-INDICES: [[T4:%.*]] = shl [[STEP]], shufflevector ( insertelement ( {{(undef|poison)}}, i32 2, {{i32|i64}} 0), {{(undef|poison)}}, zeroinitializer) +; EE-INDICES: [[T4:%.*]] = shl [[STEP]], {{shufflevector \( insertelement \( (undef|poison), i32 2, (i32|i64) 0\), (undef|poison), zeroinitializer\)|splat \(i32 2\)}} ; EE-INDICES: [[T5:%.*]] = {{add|or}} {{(disjoint )?}} [[T4]], [[T3]] ; EE-INDICES: [[IDX:%.*]] = sext [[T5]] to ; EE-INDICES: [[ADDR:%.*]] = getelementptr float, ptr [[ALLOC]], [[IDX]] diff --git a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/insert_element.ll b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/insert_element.ll index b2dcb47b5..d2ed9cef9 100644 --- a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/insert_element.ll +++ b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/insert_element.ll @@ -99,7 +99,7 @@ entry: ; IE-INDICES: [[VAL:%.*]] = uitofp {{%.*}} to ; IE-INDICES: store {{%.*}}, ptr [[ALLOC]], align 64 ; IE-INDICES: [[T1:%.*]] = call @llvm.{{(experimental\.)?}}stepvector.nxv4i32() -; IE-INDICES: [[T2:%.*]] = shl [[T1]], shufflevector ( insertelement ( {{(undef|poison)}}, i32 2, {{(i32|i64)}} 0), {{(undef|poison)}}, zeroinitializer) +; IE-INDICES: [[T2:%.*]] = shl [[T1]], {{shufflevector \( insertelement \( (undef|poison), i32 2, (i32|i64) 0\), (undef|poison), zeroinitializer\)|splat \(i32 2\)}} ; LLVM 16 deduces add/or equivalence and uses `or` instead. ; IE-INDICES: [[T3:%.*]] = {{add|or}} {{(disjoint )?}} [[T2]], {{%.*}} diff --git a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/interleaved_load.ll b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/interleaved_load.ll index d4f4c5339..708edd894 100644 --- a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/interleaved_load.ll +++ b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/interleaved_load.ll @@ -53,7 +53,7 @@ declare i64 @__mux_get_global_id(i32) ; CHECK-NEXT: [[TMP4:%.*]] = call @llvm.{{(experimental\.)?}}stepvector.nxv4i64() ; CHECK-NEXT: [[TMP5:%.*]] = mul [[TMP3]], [[TMP4]] ; CHECK-NEXT: [[TMP6:%.*]] = getelementptr i32, [[TMP1]], [[TMP5]] -; CHECK-NEXT: call void @llvm.masked.scatter.nxv4i32.nxv4p1( [[ARG0]], [[TMP6]], i32 immarg 4, shufflevector ( insertelement ( poison, i1 true, {{i32|i64}} 0), poison, zeroinitializer)) [[MASKED_ATTRS:#[0-9]+]] +; CHECK-NEXT: call void @llvm.masked.scatter.nxv4i32.nxv4p1( [[ARG0]], [[TMP6]], i32 immarg 4, {{shufflevector \( insertelement \( poison, i1 true, (i32|i64) 0\), poison, zeroinitializer\)|splat \(i1 true\)}}) [[MASKED_ATTRS:#[0-9]+]] ; CHECK-NEXT: ret void ; CHECK-NEXT: } diff --git a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/packetize_mask_varying.ll b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/packetize_mask_varying.ll index 61682c1ba..b391a57c2 100644 --- a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/packetize_mask_varying.ll +++ b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/packetize_mask_varying.ll @@ -39,7 +39,7 @@ if.end: ret void ; CHECK: define spir_kernel void @__vecz_nxv4_mask_varying ; CHECK: [[idx0:%.*]] = call @llvm.{{(experimental\.)?}}stepvector.nxv16i32() -; CHECK: [[idx1:%.*]] = lshr [[idx0]], shufflevector ( insertelement ( {{(undef|poison)}}, i32 2, {{(i32|i64)}} 0), {{(undef|poison)}}, zeroinitializer) +; CHECK: [[idx1:%.*]] = lshr [[idx0]], {{shufflevector \( insertelement \( (undef|poison), i32 2, (i32|i64) 0\), (undef|poison), zeroinitializer\)|splat \(i32 2\)}} ; Note that since we just did a lshr 2 on the input of the extend, it doesn't ; make any difference whether it's a zext or sext, but LLVM 16 prefers zext. diff --git a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/select.ll b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/select.ll index f6a8addb3..9a693646c 100644 --- a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/select.ll +++ b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/select.ll @@ -55,7 +55,7 @@ declare i64 @__mux_get_global_id(i32) ; CHECK: [[lhs:%[0-9a-z]+]] = load , ptr ; CHECK: [[rhs:%[0-9a-z]+]] = load , ptr ; CHECK: [[cmp:%[0-9a-z]+]] = icmp slt [[lhs]], [[rhs]] -; CHECK: [[sel:%[0-9a-z]+]] = select [[cmp]], [[rhs]], shufflevector ( insertelement ( {{(undef|poison)}}, i32 4, {{(i32|i64)}} 0), {{(undef|poison)}}, zeroinitializer) +; CHECK: [[sel:%[0-9a-z]+]] = select [[cmp]], [[rhs]], {{shufflevector \( insertelement \( (undef|poison), i32 4, (i32|i64) 0\), (undef|poison), zeroinitializer\)|splat \(i32 4\)}} ; CHECK: store [[sel]], ; CHECK: define spir_kernel void @__vecz_nxv4_select_vector_vector @@ -63,5 +63,5 @@ declare i64 @__mux_get_global_id(i32) ; CHECK: [[y:%[0-9a-z]+]] = load , ptr ; CHECK: [[z:%[0-9a-z]+]] = load , ptr ; CHECK: [[cmp:%[0-9a-z]+]] = icmp slt [[x]], [[y]] -; CHECK: [[sel:%[0-9a-z]+]] = select [[cmp]], [[z]], shufflevector ( insertelement ( {{(undef|poison)}}, i32 4, {{(i32|i64)}} 0), {{(undef|poison)}}, zeroinitializer) +; CHECK: [[sel:%[0-9a-z]+]] = select [[cmp]], [[z]], {{shufflevector \( insertelement \( (undef|poison), i32 4, (i32|i64) 0\), (undef|poison), zeroinitializer\)|splat \(i32 4\)}} ; CHECK: store [[sel]], diff --git a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/select_scalar_vector.ll b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/select_scalar_vector.ll index f4fa88cb1..0d58887f9 100644 --- a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/select_scalar_vector.ll +++ b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/select_scalar_vector.ll @@ -44,7 +44,7 @@ entry: ; CHECK: [[sext:%.*]] = sext [[cmp1]] to ; CHECK: store [[sext]], ptr [[alloc:%.*]], align 4 ; CHECK: [[idx0:%.*]] = call @llvm.{{(experimental\.)?}}stepvector.nxv8i32() -; CHECK: [[idx1:%.*]] = lshr [[idx0]], shufflevector ( insertelement ( {{(undef|poison)}}, i32 1, {{(i32|i64)}} 0), {{(undef|poison)}}, zeroinitializer) +; CHECK: [[idx1:%.*]] = lshr [[idx0]], {{shufflevector \( insertelement \( (undef|poison), i32 1, (i32|i64) 0\), (undef|poison), zeroinitializer\)|splat \(i32 1\)}} ; Note that since we just did a lshr 1 on the input of the extend, it doesn't ; make any difference whether it's a zext or sext, but LLVM 16 prefers zext. @@ -53,5 +53,5 @@ entry: ; CHECK: [[addrs:%.*]] = getelementptr i8, ptr [[alloc]], [[sext2]] ; CHECK: [[gather:%.*]] = call @llvm.masked.gather.nxv8i8.nxv8p0( [[addrs]], ; CHECK: [[cmp:%.*]] = trunc [[gather]] to -; CHECK: [[sel:%.*]] = select [[cmp]], [[rhs]], shufflevector ( insertelement ( {{(undef|poison)}}, i32 4, {{(i32|i64)}} 0), {{(undef|poison)}}, zeroinitializer) +; CHECK: [[sel:%.*]] = select [[cmp]], [[rhs]], {{shufflevector \( insertelement \( (undef|poison), i32 4, (i32|i64) 0\), (undef|poison), zeroinitializer\)|splat \(i32 4\)}} ; CHECK: store [[sel]], diff --git a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/shuffle.ll b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/shuffle.ll index 074502779..110fc935b 100644 --- a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/shuffle.ll +++ b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/shuffle.ll @@ -33,7 +33,7 @@ define spir_kernel void @do_shuffle_splat(i32* %aptr, <4 x i32>* %bptr, <4 x i32 ret void ; CHECK: define spir_kernel void @__vecz_nxv4_do_shuffle_splat ; CHECK: [[idx0:%.*]] = call @llvm.{{(experimental\.)?}}stepvector.nxv16i32() -; CHECK: [[idx1:%.*]] = lshr [[idx0]], shufflevector ( insertelement ( {{(undef|poison)}}, i32 2, {{(i32|i64)}} 0), {{(undef|poison)}}, zeroinitializer) +; CHECK: [[idx1:%.*]] = lshr [[idx0]], {{shufflevector \( insertelement \( (undef|poison), i32 2, (i32|i64) 0\), (undef|poison), zeroinitializer\)|splat \(i32 2\)}} ; Note that since we just did a lshr 2 on the input of the extend, it doesn't ; make any difference whether it's a zext or sext, but LLVM 16 prefers zext. diff --git a/modules/compiler/vecz/test/lit/llvm/VectorPredication/define_interleaved_load_store.ll b/modules/compiler/vecz/test/lit/llvm/VectorPredication/define_interleaved_load_store.ll index bf8a3f08a..4ea882d80 100644 --- a/modules/compiler/vecz/test/lit/llvm/VectorPredication/define_interleaved_load_store.ll +++ b/modules/compiler/vecz/test/lit/llvm/VectorPredication/define_interleaved_load_store.ll @@ -59,7 +59,7 @@ declare <4 x double> @llvm.fmuladd.v4f64(<4 x double>, <4 x double>, <4 x double ; CHECK: %BroadcastAddr.splatinsert = insertelement poison, ptr addrspace(1) %0, {{i32|i64}} 0 ; CHECK: %BroadcastAddr.splat = shufflevector %BroadcastAddr.splatinsert, poison, zeroinitializer ; CHECK: %3 = call @llvm.{{(experimental\.)?}}stepvector.nxv4i64() -; CHECK: %4 = mul shufflevector ( insertelement ( poison, i64 4, {{i32|i64}} 0), poison, zeroinitializer), %3 +; CHECK: %4 = mul {{shufflevector \( insertelement \( poison, i64 4, (i32|i64) 0\), poison, zeroinitializer\)|splat \(i64 4\)}}, %3 ; CHECK: %5 = getelementptr double, %BroadcastAddr.splat, %4 ; CHECK: %6 = call @llvm.vp.gather.nxv4f64.nxv4p1( %5, %1, i32 %2) ; CHECK: ret %6 @@ -73,7 +73,7 @@ declare <4 x double> @llvm.fmuladd.v4f64(<4 x double>, <4 x double>, <4 x double ; CHECK: %BroadcastAddr.splatinsert = insertelement poison, ptr addrspace(1) %1, {{i32|i64}} 0 ; CHECK: %BroadcastAddr.splat = shufflevector %BroadcastAddr.splatinsert, poison, zeroinitializer ; CHECK: %4 = call @llvm.{{(experimental\.)?}}stepvector.nxv4i64() -; CHECK: %5 = mul shufflevector ( insertelement ( poison, i64 4, {{i32|i64}} 0), poison, zeroinitializer), %4 +; CHECK: %5 = mul {{shufflevector \( insertelement \( poison, i64 4, (i32|i64) 0\), poison, zeroinitializer\)|splat \(i64 4\)}}, %4 ; CHECK: %6 = getelementptr double, %BroadcastAddr.splat, %5 ; CHECK: call void @llvm.vp.scatter.nxv4f64.nxv4p1( %0, %6, %2, i32 %3) ; CHECK: ret void diff --git a/modules/compiler/vecz/test/lit/llvm/VectorPredication/load_add_store.ll b/modules/compiler/vecz/test/lit/llvm/VectorPredication/load_add_store.ll index 8da3755e1..7e8f0770d 100644 --- a/modules/compiler/vecz/test/lit/llvm/VectorPredication/load_add_store.ll +++ b/modules/compiler/vecz/test/lit/llvm/VectorPredication/load_add_store.ll @@ -58,7 +58,7 @@ entry: ; CHECK_1S: [[T1:%.*]] = shl i64 [[T0]], 2 ; CHECK_1S: [[T2:%.*]] = call i64 @llvm.umin.i64(i64 [[WREM]], i64 [[T1]]) ; CHECK_1S: [[VL:%.*]] = trunc {{(nuw )?(nsw )?}}i64 [[T2]] to i32 -; CHECK_1S: [[LHS:%.*]] = call @llvm.vp.load.nxv4i32.p0(ptr {{%.*}}, [[TRUEMASK: shufflevector \( insertelement \( (undef|poison), i1 true, (i32|i64) 0\), (undef|poison), zeroinitializer\)]], i32 [[VL]]) +; CHECK_1S: [[LHS:%.*]] = call @llvm.vp.load.nxv4i32.p0(ptr {{%.*}}, [[TRUEMASK: (shufflevector \( insertelement \( (undef|poison), i1 true, (i32|i64) 0\), (undef|poison), zeroinitializer\)|splat \(i1 true\))]], i32 [[VL]]) ; CHECK_1S: [[RHS:%.*]] = call @llvm.vp.load.nxv4i32.p0(ptr {{%.*}}, [[TRUEMASK]], i32 [[VL]]) ; CHECK_1S: [[ADD:%.*]] = call @llvm.vp.add.nxv4i32( [[LHS]], [[RHS]], [[TRUEMASK]], i32 [[VL]]) ; CHECK_1S: call void @llvm.vp.store.nxv4i32.p0( [[ADD]], ptr {{%.*}}, [[TRUEMASK]], i32 [[VL]]) @@ -99,7 +99,7 @@ entry: ; CHECK_V4_1S: [[VL:%.*]] = trunc {{(nuw )?(nsw )?}}i64 [[T2]] to i32 ; Each WI performs 4 elements, so multiply the VL by 4 ; CHECK_V4_1S: [[SVL:%.*]] = shl i32 [[VL]], 2 -; CHECK_V4_1S: [[LHS:%.*]] = call @llvm.vp.load.nxv16i32.p0(ptr {{%.*}}, [[TRUEMASK: shufflevector \( insertelement \( (undef|poison), i1 true, (i32|i64) 0\), (undef|poison), zeroinitializer\)]], i32 [[SVL]]) +; CHECK_V4_1S: [[LHS:%.*]] = call @llvm.vp.load.nxv16i32.p0(ptr {{%.*}}, [[TRUEMASK: (shufflevector \( insertelement \( (undef|poison), i1 true, (i32|i64) 0\), (undef|poison), zeroinitializer\)|splat \(i1 true\))]], i32 [[SVL]]) ; CHECK_V4_1S: [[RHS:%.*]] = call @llvm.vp.load.nxv16i32.p0(ptr {{%.*}}, [[TRUEMASK]], i32 [[SVL]]) ; CHECK_V4_1S: [[ADD:%.*]] = call @llvm.vp.add.nxv16i32( [[LHS]], [[RHS]], [[TRUEMASK]], i32 [[SVL]]) ; CHECK_V4_1S: call void @llvm.vp.store.nxv16i32.p0( [[ADD]], ptr {{%.*}}, [[TRUEMASK]], i32 [[SVL]]) diff --git a/modules/compiler/vecz/test/lit/llvm/VectorPredication/udiv.ll b/modules/compiler/vecz/test/lit/llvm/VectorPredication/udiv.ll index 7cd87a3cd..bf082b453 100644 --- a/modules/compiler/vecz/test/lit/llvm/VectorPredication/udiv.ll +++ b/modules/compiler/vecz/test/lit/llvm/VectorPredication/udiv.ll @@ -44,7 +44,7 @@ entry: ; CHECK: [[T1:%.*]] = shl i64 [[T0]], 1 ; CHECK: [[T2:%.*]] = call i64 @llvm.umin.i64(i64 [[WREM]], i64 [[T1]]) ; CHECK: [[VL:%.*]] = trunc i64 [[T2]] to i32 -; CHECK: [[LHS:%.*]] = call @llvm.vp.load.nxv2i32.p0(ptr {{%.*}}, [[TRUEMASK: shufflevector \( insertelement \( (undef|poison), i1 true, (i32|i64) 0\), (undef|poison), zeroinitializer\)]], i32 [[VL]]) +; CHECK: [[LHS:%.*]] = call @llvm.vp.load.nxv2i32.p0(ptr {{%.*}}, [[TRUEMASK: (shufflevector \( insertelement \( (undef|poison), i1 true, (i32|i64) 0\), (undef|poison), zeroinitializer\)|splat \(i1 true\))]], i32 [[VL]]) ; CHECK: [[RHS:%.*]] = call @llvm.vp.load.nxv2i32.p0(ptr {{%.*}}, [[TRUEMASK]], i32 [[VL]]) ; CHECK: [[ADD:%.*]] = call @llvm.vp.udiv.nxv2i32( [[LHS]], [[RHS]], [[TRUEMASK]], i32 [[VL]]) ; CHECK: call void @llvm.vp.store.nxv2i32.p0( [[ADD]], ptr {{%.*}}, [[TRUEMASK]], i32 [[VL]])