diff --git a/modules/compiler/vecz/source/vectorization_context.cpp b/modules/compiler/vecz/source/vectorization_context.cpp index 15b1420e6..d31b6cb3a 100644 --- a/modules/compiler/vecz/source/vectorization_context.cpp +++ b/modules/compiler/vecz/source/vectorization_context.cpp @@ -207,6 +207,13 @@ Function *VectorizationContext::getOrCreateInternalBuiltin(StringRef Name, if (!F && FT) { F = dyn_cast_or_null( Module.getOrInsertFunction(Name, FT).getCallee()); + if (F) { + // Set some default attributes on the function. + // We never use exceptions + F->addFnAttr(Attribute::NoUnwind); + // Recursion is not supported in ComputeMux + F->addFnAttr(Attribute::NoRecurse); + } } return F; diff --git a/modules/compiler/vecz/test/lit/llvm/OpaquePointers/control_flow_conversion_ptrs.ll b/modules/compiler/vecz/test/lit/llvm/OpaquePointers/control_flow_conversion_ptrs.ll index 689cf3057..853fb9229 100644 --- a/modules/compiler/vecz/test/lit/llvm/OpaquePointers/control_flow_conversion_ptrs.ll +++ b/modules/compiler/vecz/test/lit/llvm/OpaquePointers/control_flow_conversion_ptrs.ll @@ -43,10 +43,12 @@ if.end: ret void } -; CHECK: define void @__vecz_b_masked_store4_u3ptru3ptrb(ptr [[A:%.*]], ptr [[B:%.*]], i1 [[MASK:%.*]]) { +; CHECK: define void @__vecz_b_masked_store4_u3ptru3ptrb(ptr [[A:%.*]], ptr [[B:%.*]], i1 [[MASK:%.*]]) [[ATTRS:#[0-9]+]] { ; CHECK: br i1 [[MASK]], label %[[IF:.*]], label %[[EXIT:.*]] ; CHECK: [[IF]]: ; CHECK-NEXT: store ptr [[A]], ptr [[B]], align 4 ; CHECK-NEXT: br label %[[EXIT]] ; CHECK: [[EXIT]]: ; CHECK-NEXT: ret void + +; CHECK: attributes [[ATTRS]] = { norecurse nounwind } diff --git a/modules/compiler/vecz/test/lit/llvm/OpaquePointers/masked_store.ll b/modules/compiler/vecz/test/lit/llvm/OpaquePointers/masked_store.ll index 43e027f0b..cd1652f0d 100644 --- a/modules/compiler/vecz/test/lit/llvm/OpaquePointers/masked_store.ll +++ b/modules/compiler/vecz/test/lit/llvm/OpaquePointers/masked_store.ll @@ -65,7 +65,7 @@ if.end: ret void } -; CHECK: define void @__vecz_b_masked_store4_fu3ptrb(float [[A:%.*]], ptr [[B:%.*]], i1 [[MASK:%.*]]) { +; CHECK: define void @__vecz_b_masked_store4_fu3ptrb(float [[A:%.*]], ptr [[B:%.*]], i1 [[MASK:%.*]]) [[ATTRS:#[0-9]+]] { ; CHECK: br i1 [[MASK]], label %[[IF:.*]], label %[[EXIT:.*]] ; CHECK: [[IF]]: ; CHECK-NEXT: store float [[A]], ptr [[B]], align 4 @@ -73,10 +73,12 @@ if.end: ; CHECK: [[EXIT]]: ; CHECK-NEXT: ret void -; CHECK: define void @__vecz_b_masked_store4_fu3ptrU3AS3b(float [[A:%.*]], ptr addrspace(3) [[B:%.*]], i1 [[MASK:%.*]]) { +; CHECK: define void @__vecz_b_masked_store4_fu3ptrU3AS3b(float [[A:%.*]], ptr addrspace(3) [[B:%.*]], i1 [[MASK:%.*]]) [[ATTRS]] { ; CHECK: br i1 [[MASK]], label %[[IF:.*]], label %[[EXIT:.*]] ; CHECK: [[IF]]: ; CHECK-NEXT: store float [[A]], ptr addrspace(3) [[B]], align 4 ; CHECK-NEXT: br label %[[EXIT]] ; CHECK: [[EXIT]]: ; CHECK-NEXT: ret void + +; CHECK: attributes [[ATTRS]] = { norecurse nounwind } 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 895481427..b1199cf84 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 @@ -53,7 +53,7 @@ declare void @__mux_work_group_barrier(i32, i32, i32) declare <4 x double> @llvm.fmuladd.v4f64(<4 x double>, <4 x double>, <4 x double>) ; Test if the interleaved store is defined correctly -; CHECK: define void @__vecz_b_interleaved_store8_4_u5nxv4du3ptrU3AS1( %0, ptr addrspace(1) %1) { +; CHECK: define void @__vecz_b_interleaved_store8_4_u5nxv4du3ptrU3AS1( %0, ptr addrspace(1) %1) [[ATTRS:#[0-9]+]] { ; CHECK: entry: ; CHECK: %BroadcastAddr.splatinsert = insertelement poison, ptr addrspace(1) %1, {{i32|i64}} 0 ; CHECK: %BroadcastAddr.splat = shufflevector %BroadcastAddr.splatinsert, poison, zeroinitializer @@ -63,3 +63,5 @@ declare <4 x double> @llvm.fmuladd.v4f64(<4 x double>, <4 x double>, <4 x double ; CHECK: call void @llvm.masked.scatter.nxv4f64.nxv4p1( %0, %4, i32 immarg 8, shufflevector ( insertelement ( poison, i1 true, {{i32|i64}} 0), poison, zeroinitializer)) ; CHECK: ret void ; CHECK: } + +; CHECK: attributes [[ATTRS]] = { norecurse nounwind } 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 5b71067b8..1bff13b48 100644 --- a/modules/compiler/vecz/test/lit/llvm/ScalableVectors/interleaved_load.ll +++ b/modules/compiler/vecz/test/lit/llvm/ScalableVectors/interleaved_load.ll @@ -44,7 +44,7 @@ entry: declare i64 @__mux_get_global_id(i32) -; CHECK: define void @__vecz_b_interleaved_store4_V_u5nxv4ju3ptrU3AS1( [[ARG0:%.*]], ptr addrspace(1) [[ARG1:%.*]], i64 [[ARG2:%.*]]) { +; CHECK: define void @__vecz_b_interleaved_store4_V_u5nxv4ju3ptrU3AS1( [[ARG0:%.*]], ptr addrspace(1) [[ARG1:%.*]], i64 [[ARG2:%.*]]) [[ATTRS:#[0-9]+]] { ; CHECK-NEXT: entry: ; CHECK-NEXT: [[TMP0:%.*]] = insertelement poison, ptr addrspace(1) [[ARG1]], {{i32|i64}} 0 ; CHECK-NEXT: [[TMP1:%.*]] = shufflevector [[TMP0]], poison, zeroinitializer @@ -53,8 +53,9 @@ 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)) #[[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)) [[MASKED_ATTRS:#[0-9]+]] ; CHECK-NEXT: ret void ; CHECK-NEXT: } -; CHECK: attributes #[[ATTRS]] = { +; CHECK-DAG: attributes [[ATTRS]] = { norecurse nounwind } +; CHECK-DAG: attributes [[MASKED_ATTRS]] = { 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 fd80c3690..2d01057a6 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 @@ -54,7 +54,7 @@ declare <4 x double> @llvm.fmuladd.v4f64(<4 x double>, <4 x double>, <4 x double ; Test if the interleaved load is defined correctly ; Vector-predicated interleaved loads are always masked -; CHECK: define @__vecz_b_masked_interleaved_load8_vp_4_u5nxv4du3ptrU3AS1u5nxv4bj(ptr addrspace(1){{( %0)?}}, {{( %1)?}}, i32{{( %2)?}}) { +; CHECK: define @__vecz_b_masked_interleaved_load8_vp_4_u5nxv4du3ptrU3AS1u5nxv4bj(ptr addrspace(1){{( %0)?}}, {{( %1)?}}, i32{{( %2)?}}) [[ATTRS:#[0-9]+]] { ; CHECK: entry: ; CHECK: %BroadcastAddr.splatinsert = insertelement poison, ptr addrspace(1) %0, {{i32|i64}} 0 ; CHECK: %BroadcastAddr.splat = shufflevector %BroadcastAddr.splatinsert, poison, zeroinitializer @@ -68,7 +68,7 @@ declare <4 x double> @llvm.fmuladd.v4f64(<4 x double>, <4 x double>, <4 x double ; Test if the interleaved store is defined correctly ; Vector-predicated interleaved stores are always masked -; CHECK: define void @__vecz_b_masked_interleaved_store8_vp_4_u5nxv4du3ptrU3AS1u5nxv4bj({{( %0)?}}, ptr addrspace(1){{( %1)?}}, {{( %2)?}}, i32{{( %3)?}}) +; CHECK: define void @__vecz_b_masked_interleaved_store8_vp_4_u5nxv4du3ptrU3AS1u5nxv4bj({{( %0)?}}, ptr addrspace(1){{( %1)?}}, {{( %2)?}}, i32{{( %3)?}}) [[ATTRS]] ; CHECK: entry: ; CHECK: %BroadcastAddr.splatinsert = insertelement poison, ptr addrspace(1) %1, {{i32|i64}} 0 ; CHECK: %BroadcastAddr.splat = shufflevector %BroadcastAddr.splatinsert, poison, zeroinitializer @@ -78,3 +78,5 @@ declare <4 x double> @llvm.fmuladd.v4f64(<4 x double>, <4 x double>, <4 x double ; CHECK: call void @llvm.vp.scatter.nxv4f64.nxv4p1( %0, %6, %2, i32 %3) ; CHECK: ret void ; CHECK: } + +; CHECK: attributes [[ATTRS]] = { norecurse nounwind } diff --git a/modules/compiler/vecz/test/lit/llvm/VectorPredication/define_masked_load_store.ll b/modules/compiler/vecz/test/lit/llvm/VectorPredication/define_masked_load_store.ll index 210a95872..549f8cb8e 100644 --- a/modules/compiler/vecz/test/lit/llvm/VectorPredication/define_masked_load_store.ll +++ b/modules/compiler/vecz/test/lit/llvm/VectorPredication/define_masked_load_store.ll @@ -64,13 +64,15 @@ declare i64 @__mux_get_local_size(i32) declare i64 @__mux_get_group_id(i32) ; Test if the masked store is defined correctly -; CHECK: define void @__vecz_b_masked_store4_vp_Dv4_ju3ptrU3AS1Dv4_bj(<4 x i32>{{( %0)?}}, ptr addrspace(1){{( %1)?}}, <4 x i1>{{( %2)?}}, i32{{( %3)?}}) { +; CHECK: define void @__vecz_b_masked_store4_vp_Dv4_ju3ptrU3AS1Dv4_bj(<4 x i32>{{( %0)?}}, ptr addrspace(1){{( %1)?}}, <4 x i1>{{( %2)?}}, i32{{( %3)?}}) [[ATTRS:#[0-9]+]] { ; CHECK: entry: ; CHECK: call void @llvm.vp.store.v4i32.p1(<4 x i32> %0, ptr addrspace(1) %1, <4 x i1> %2, i32 %3) ; CHECK: ret void ; Test if the masked load is defined correctly -; CHECK: define <4 x i32> @__vecz_b_masked_load4_vp_Dv4_ju3ptrU3AS2Dv4_bj(ptr addrspace(2){{( %0)?}}, <4 x i1>{{( %1)?}}, i32{{( %2)?}}) +; CHECK: define <4 x i32> @__vecz_b_masked_load4_vp_Dv4_ju3ptrU3AS2Dv4_bj(ptr addrspace(2){{( %0)?}}, <4 x i1>{{( %1)?}}, i32{{( %2)?}}) [[ATTRS]] { ; CHECK: entry: ; CHECK: %3 = call <4 x i32> @llvm.vp.load.v4i32.p2(ptr addrspace(2) %0, <4 x i1> %1, i32 %2) ; CHECK: ret <4 x i32> %3 + +; CHECK: attributes [[ATTRS]] = { norecurse nounwind } diff --git a/modules/compiler/vecz/test/lit/llvm/VectorPredication/scatter_gather.ll b/modules/compiler/vecz/test/lit/llvm/VectorPredication/scatter_gather.ll index 1c4fccb05..15d66ea84 100644 --- a/modules/compiler/vecz/test/lit/llvm/VectorPredication/scatter_gather.ll +++ b/modules/compiler/vecz/test/lit/llvm/VectorPredication/scatter_gather.ll @@ -55,11 +55,11 @@ entry: ; CHECK: [[v:%.*]] = call @llvm.vp.load.nxv4i32.p1( ; CHECK: call void @__vecz_b_masked_scatter_store4_vp_u5nxv4ju14nxv4u3ptrU3AS1u5nxv4bj( [[v]], -; CHECK: define @__vecz_b_masked_gather_load4_vp_u5nxv4ju14nxv4u3ptrU3AS1u5nxv4bj( %0, %1, i32 %2) { +; CHECK: define @__vecz_b_masked_gather_load4_vp_u5nxv4ju14nxv4u3ptrU3AS1u5nxv4bj( %0, %1, i32 %2) [[ATTRS:#[0-9]+]] { ; CHECK: %3 = call @llvm.vp.gather.nxv4i32.nxv4p1( %0, %1, i32 %2) ; CHECK: ret %3 -; CHECK: define void @__vecz_b_masked_scatter_store4_vp_u5nxv4ju14nxv4u3ptrU3AS1u5nxv4bj( %0, %1, %2, i32 %3) { +; CHECK: define void @__vecz_b_masked_scatter_store4_vp_u5nxv4ju14nxv4u3ptrU3AS1u5nxv4bj( %0, %1, %2, i32 %3) [[ATTRS]] { ; CHECK: entry: ; CHECK: call void @llvm.vp.scatter.nxv4i32.nxv4p1( %0, %1, %2, i32 %3) ; CHECK: ret void diff --git a/modules/compiler/vecz/test/lit/llvm/control_flow_conversion_ptrs.ll b/modules/compiler/vecz/test/lit/llvm/control_flow_conversion_ptrs.ll index 5232baa40..457568e63 100644 --- a/modules/compiler/vecz/test/lit/llvm/control_flow_conversion_ptrs.ll +++ b/modules/compiler/vecz/test/lit/llvm/control_flow_conversion_ptrs.ll @@ -41,7 +41,7 @@ if.else: if.end: ret void -; CHECK: define void @__vecz_b_masked_store4_u3ptru3ptrb(ptr [[A:%.*]], ptr [[B:%.*]], i1 [[MASK:%.*]]) { +; CHECK: define void @__vecz_b_masked_store4_u3ptru3ptrb(ptr [[A:%.*]], ptr [[B:%.*]], i1 [[MASK:%.*]]) [[ATTRS:#[0-9]+]] { ; CHECK: br i1 [[MASK]], label %[[IF:.*]], label %[[EXIT:.*]] ; CHECK: [[IF]]: ; CHECK-NEXT: store ptr [[A]], ptr [[B]], align 4 @@ -71,3 +71,5 @@ if.else: if.end: ret void } + +; CHECK: attributes [[ATTRS]] = { norecurse nounwind } diff --git a/modules/compiler/vecz/test/lit/llvm/define_gather_load.ll b/modules/compiler/vecz/test/lit/llvm/define_gather_load.ll index 45a177ad1..eff4f12e6 100644 --- a/modules/compiler/vecz/test/lit/llvm/define_gather_load.ll +++ b/modules/compiler/vecz/test/lit/llvm/define_gather_load.ll @@ -38,6 +38,8 @@ entry: declare i64 @__mux_get_global_id(i32) ; Test if the scatter store is defined correctly -; CHECK: define <4 x i64> @__vecz_b_gather_load4_Dv4_mDv4_u3ptr(<4 x ptr>{{( %0)?}}) { +; CHECK: define <4 x i64> @__vecz_b_gather_load4_Dv4_mDv4_u3ptr(<4 x ptr>{{( %0)?}}) [[ATTRS:#[0-9]+]] { ; CHECK: %[[V1:[0-9]+]] = call <4 x i64> @llvm.masked.gather.v4i64.v4p0(<4 x ptr> %0, i32{{( immarg)?}} 4, <4 x i1> , ; CHECK: ret <4 x i64> %[[V1]] + +; CHECK: attributes [[ATTRS]] = { norecurse nounwind } diff --git a/modules/compiler/vecz/test/lit/llvm/define_gather_load_as_masked.ll b/modules/compiler/vecz/test/lit/llvm/define_gather_load_as_masked.ll index b28780408..c25c29af3 100644 --- a/modules/compiler/vecz/test/lit/llvm/define_gather_load_as_masked.ll +++ b/modules/compiler/vecz/test/lit/llvm/define_gather_load_as_masked.ll @@ -38,6 +38,8 @@ entry: declare i64 @__mux_get_global_id(i32) ; Test if the scatter store is defined correctly -; CHECK: define <4 x i64> @__vecz_b_gather_load4_Dv4_mDv4_u3ptr(<4 x ptr>{{( %0)?}}) { +; CHECK: define <4 x i64> @__vecz_b_gather_load4_Dv4_mDv4_u3ptr(<4 x ptr>{{( %0)?}}) [[ATTRS:#[0-9]+]] { ; CHECK: call <4 x i64> @llvm.masked.gather.v4i64.v4p0(<4 x ptr> %0, i32{{( immarg)?}} 4, <4 x i1> , <4 x i64> undef) ; CHECK: ret <4 x i64> + +; CHECK: attributes [[ATTRS]] = { norecurse nounwind } diff --git a/modules/compiler/vecz/test/lit/llvm/define_scatter_store.ll b/modules/compiler/vecz/test/lit/llvm/define_scatter_store.ll index e41f41d52..a035ca05d 100644 --- a/modules/compiler/vecz/test/lit/llvm/define_scatter_store.ll +++ b/modules/compiler/vecz/test/lit/llvm/define_scatter_store.ll @@ -38,7 +38,9 @@ entry: declare i64 @__mux_get_global_id(i32) ; Test if the scatter store is defined correctly -; CHECK: define void @__vecz_b_scatter_store4_Dv4_mDv4_u3ptr(<4 x i64>{{( %0)?}}, <4 x ptr>{{( %1)?}}) { +; CHECK: define void @__vecz_b_scatter_store4_Dv4_mDv4_u3ptr(<4 x i64>{{( %0)?}}, <4 x ptr>{{( %1)?}}) [[ATTRS:#[0-9]+]] { ; CHECK: entry ; CHECK: call void @llvm.masked.scatter.v4i64.v4p0(<4 x i64> %0, <4 x ptr> %1, i32{{( immarg)?}} 4, <4 x i1> ) ; CHECK: ret void + +; CHECK: attributes [[ATTRS]] = { norecurse nounwind } diff --git a/modules/compiler/vecz/test/lit/llvm/define_scatter_store_as_masked.ll b/modules/compiler/vecz/test/lit/llvm/define_scatter_store_as_masked.ll index 768599fcb..fd7a7570b 100644 --- a/modules/compiler/vecz/test/lit/llvm/define_scatter_store_as_masked.ll +++ b/modules/compiler/vecz/test/lit/llvm/define_scatter_store_as_masked.ll @@ -38,7 +38,9 @@ entry: declare i64 @__mux_get_global_id(i32) ; Test if the scatter store is defined correctly -; CHECK: define void @__vecz_b_scatter_store4_Dv4_mDv4_u3ptr(<4 x i64>{{( %0)?}}, <4 x ptr>{{( %1)?}}) { +; CHECK: define void @__vecz_b_scatter_store4_Dv4_mDv4_u3ptr(<4 x i64>{{( %0)?}}, <4 x ptr>{{( %1)?}}) [[ATTRS:#[0-9]+]] { ; CHECK: entry: ; CHECK: call void @llvm.masked.scatter.v4i64.v4p0(<4 x i64> %0, <4 x ptr> %1, i32{{( immarg)?}} 4, <4 x i1> ) ; CHECK: ret void + +; CHECK: attributes [[ATTRS]] = { norecurse nounwind } diff --git a/modules/compiler/vecz/test/lit/llvm/masked_atomics.ll b/modules/compiler/vecz/test/lit/llvm/masked_atomics.ll index 2f11e37c2..7413f6ca6 100644 --- a/modules/compiler/vecz/test/lit/llvm/masked_atomics.ll +++ b/modules/compiler/vecz/test/lit/llvm/masked_atomics.ll @@ -50,7 +50,7 @@ if.end: ; preds = %if.then, %entry ret void } -; CHECK: define <4 x i32> @__vecz_b_v4_masked_atomicrmw_add_align4_acquire_1_Dv4_u3ptrDv4_jDv4_b(<4 x ptr> [[PTRS:%0]], <4 x i32> [[VALS:%1]], <4 x i1> [[MASK:%2]]) { +; CHECK: define <4 x i32> @__vecz_b_v4_masked_atomicrmw_add_align4_acquire_1_Dv4_u3ptrDv4_jDv4_b(<4 x ptr> [[PTRS:%0]], <4 x i32> [[VALS:%1]], <4 x i1> [[MASK:%2]]) [[ATTRS:#[0-9]+]] { ; CHECK: entry: ; CHECK: br label %loopIR @@ -77,11 +77,13 @@ if.end: ; preds = %if.then, %entry ; Assume that all masked atomicrmw operations follow the logic above. Just ; check that the right atomicrmw instruction is being generated. -; CHECK: define <4 x i32> @__vecz_b_v4_masked_atomicrmw_umin_align2_monotonic_1_Dv4_u3ptrDv4_jDv4_b(<4 x ptr> [[PTRS:%0]], <4 x i32> [[VALS:%1]], <4 x i1> [[MASK:%2]]) { +; CHECK: define <4 x i32> @__vecz_b_v4_masked_atomicrmw_umin_align2_monotonic_1_Dv4_u3ptrDv4_jDv4_b(<4 x ptr> [[PTRS:%0]], <4 x i32> [[VALS:%1]], <4 x i1> [[MASK:%2]]) [[ATTRS]] { ; CHECK: atomicrmw umin ptr {{%.*}}, i32 {{%.*}} monotonic, align 2 -; CHECK: define <4 x float> @__vecz_b_v4_masked_atomicrmw_volatile_fmax_align4_seqcst_0_Dv4_u3ptrDv4_fDv4_b(<4 x ptr> [[PTRS:%0]], <4 x float> [[VALS:%1]], <4 x i1> [[MASK:%2]]) { +; CHECK: define <4 x float> @__vecz_b_v4_masked_atomicrmw_volatile_fmax_align4_seqcst_0_Dv4_u3ptrDv4_fDv4_b(<4 x ptr> [[PTRS:%0]], <4 x float> [[VALS:%1]], <4 x i1> [[MASK:%2]]) [[ATTRS]] { ; CHECK: atomicrmw volatile fmax ptr {{%.*}}, float {{%.*}} syncscope("singlethread") seq_cst, align 4 +; CHECK: attributes [[ATTRS]] = { norecurse nounwind } + declare i64 @__mux_get_global_id(i32) diff --git a/modules/compiler/vecz/test/lit/llvm/masked_cmpxchg.ll b/modules/compiler/vecz/test/lit/llvm/masked_cmpxchg.ll index 73aec6dfc..80576d6aa 100644 --- a/modules/compiler/vecz/test/lit/llvm/masked_cmpxchg.ll +++ b/modules/compiler/vecz/test/lit/llvm/masked_cmpxchg.ll @@ -64,7 +64,7 @@ if.end: ; preds = %if.then, %entry ret void } -; CHECK: define { <4 x i32>, <4 x i1> } @__vecz_b_v4_masked_cmpxchg_align4_acquire_monotonic_1_Dv4_u3ptrDv4_jDv4_jDv4_b(<4 x ptr> [[PTRS:%0]], <4 x i32> [[CMPS:%1]], <4 x i32> [[NEWS:%2]], <4 x i1> [[MASK:%3]]) { +; CHECK: define { <4 x i32>, <4 x i1> } @__vecz_b_v4_masked_cmpxchg_align4_acquire_monotonic_1_Dv4_u3ptrDv4_jDv4_jDv4_b(<4 x ptr> [[PTRS:%0]], <4 x i32> [[CMPS:%1]], <4 x i32> [[NEWS:%2]], <4 x i1> [[MASK:%3]]) [[ATTRS:#[0-9]+]] { ; CHECK: entry: ; CHECK: br label %loopIR @@ -99,7 +99,9 @@ if.end: ; preds = %if.then, %entry ; Assume that all masked cmpxchg operations follow the logic above. Just ; check that the right cmpxchg instruction is being generated. -; CHECK: define { <4 x i32>, <4 x i1> } @__vecz_b_v4_masked_cmpxchg_weak_volatile_align8_monotonic_seqcst_0_Dv4_u3ptrDv4_jDv4_jDv4_b(<4 x ptr> [[PTRS:%0]], <4 x i32> [[CMPS:%1]], <4 x i32> [[NEWS:%2]], <4 x i1> [[MASK:%3]]) { +; CHECK: define { <4 x i32>, <4 x i1> } @__vecz_b_v4_masked_cmpxchg_weak_volatile_align8_monotonic_seqcst_0_Dv4_u3ptrDv4_jDv4_jDv4_b(<4 x ptr> [[PTRS:%0]], <4 x i32> [[CMPS:%1]], <4 x i32> [[NEWS:%2]], <4 x i1> [[MASK:%3]]) [[ATTRS]] { ; CHECK: cmpxchg weak volatile ptr {{%.*}}, i32 {{%.*}}, i32 {{%.*}} syncscope("singlethread") monotonic seq_cst, align 8 +; CHECK: attributes [[ATTRS]] = { norecurse nounwind } + declare i64 @__mux_get_global_id(i32) diff --git a/modules/compiler/vecz/test/lit/llvm/masked_interleaved.ll b/modules/compiler/vecz/test/lit/llvm/masked_interleaved.ll index 4c3d9e829..43dcc6217 100644 --- a/modules/compiler/vecz/test/lit/llvm/masked_interleaved.ll +++ b/modules/compiler/vecz/test/lit/llvm/masked_interleaved.ll @@ -65,10 +65,12 @@ attributes #2 = { nobuiltin } !6 = !{!"clang version 3.8.0 "} -; CHECK: define void @__vecz_b_masked_interleaved_store4_2_Dv4_ju3ptrU3AS1Dv4_b(<4 x i32>{{( %0)?}}, ptr addrspace(1){{( %1)?}}, <4 x i1>{{( %2)?}}) { +; CHECK: define void @__vecz_b_masked_interleaved_store4_2_Dv4_ju3ptrU3AS1Dv4_b(<4 x i32>{{( %0)?}}, ptr addrspace(1){{( %1)?}}, <4 x i1>{{( %2)?}}) [[ATTRS:#[0-9]+]] { ; CHECK: entry: ; CHECK: %BroadcastAddr.splatinsert = insertelement <4 x ptr addrspace(1)> {{poison|undef}}, ptr addrspace(1) %1, {{i32|i64}} 0 ; CHECK: %BroadcastAddr.splat = shufflevector <4 x ptr addrspace(1)> %BroadcastAddr.splatinsert, <4 x ptr addrspace(1)> {{poison|undef}}, <4 x i32> zeroinitializer ; CHECK: %3 = getelementptr i32, <4 x ptr addrspace(1)> %BroadcastAddr.splat, <4 x i64> ; CHECK: call void @llvm.masked.scatter.v4i32.v4p1(<4 x i32> %0, <4 x ptr addrspace(1)> %3, i32{{( immarg)?}} 4, <4 x i1> %2) # ; CHECK: ret void + +; CHECK: attributes [[ATTRS]] = { norecurse nounwind } diff --git a/modules/compiler/vecz/test/lit/llvm/masked_interleaved_as_scatter.ll b/modules/compiler/vecz/test/lit/llvm/masked_interleaved_as_scatter.ll index 5166cab21..11d14417f 100644 --- a/modules/compiler/vecz/test/lit/llvm/masked_interleaved_as_scatter.ll +++ b/modules/compiler/vecz/test/lit/llvm/masked_interleaved_as_scatter.ll @@ -65,7 +65,7 @@ attributes #2 = { nobuiltin } !6 = !{!"clang version 3.8.0 "} -; CHECK: define void @__vecz_b_masked_interleaved_store4_2_Dv4_ju3ptrU3AS1Dv4_b(<4 x i32>{{( %0)?}}, ptr addrspace(1){{( %1)?}}, <4 x i1>{{( %2)?}}) { +; CHECK: define void @__vecz_b_masked_interleaved_store4_2_Dv4_ju3ptrU3AS1Dv4_b(<4 x i32>{{( %0)?}}, ptr addrspace(1){{( %1)?}}, <4 x i1>{{( %2)?}}) [[ATTRS:#[0-9]+]] { ; Check for the address splat ; CHECK: %[[BROADCASTADDRSPLATINSERT:.+]] = insertelement <4 x ptr addrspace(1)> {{poison|undef}}, ptr addrspace(1) %{{.+}}, {{i32|i64}} 0 @@ -73,3 +73,5 @@ attributes #2 = { nobuiltin } ; CHECK: getelementptr i32, <4 x ptr addrspace(1)> %[[BROADCASTADDRSPLAT]], <4 x i64> ; CHECK: ret void + +; CHECK: attributes [[ATTRS]] = { norecurse nounwind }