From cfb43528df2f56a6cd3e4bf80056569ba9ba1a4f Mon Sep 17 00:00:00 2001 From: Fraser Cormack Date: Tue, 2 Jan 2024 18:37:35 +0000 Subject: [PATCH] [vecz] Add nounwind/norecurse attributes to internal vecz builtins Though in practice we almost always have these attributes set on the functions that call these, this change should help codegen in the rare case that the builtins aren't inlined. --- modules/compiler/vecz/source/vectorization_context.cpp | 7 +++++++ .../llvm/OpaquePointers/control_flow_conversion_ptrs.ll | 4 +++- .../vecz/test/lit/llvm/OpaquePointers/masked_store.ll | 6 ++++-- .../ScalableVectors/define_interleaved_store_as_masked.ll | 4 +++- .../test/lit/llvm/ScalableVectors/interleaved_load.ll | 7 ++++--- .../VectorPredication/define_interleaved_load_store.ll | 6 ++++-- .../llvm/VectorPredication/define_masked_load_store.ll | 6 ++++-- .../test/lit/llvm/VectorPredication/scatter_gather.ll | 4 ++-- .../vecz/test/lit/llvm/control_flow_conversion_ptrs.ll | 4 +++- modules/compiler/vecz/test/lit/llvm/define_gather_load.ll | 4 +++- .../vecz/test/lit/llvm/define_gather_load_as_masked.ll | 4 +++- .../compiler/vecz/test/lit/llvm/define_scatter_store.ll | 4 +++- .../vecz/test/lit/llvm/define_scatter_store_as_masked.ll | 4 +++- modules/compiler/vecz/test/lit/llvm/masked_atomics.ll | 8 +++++--- modules/compiler/vecz/test/lit/llvm/masked_cmpxchg.ll | 6 ++++-- modules/compiler/vecz/test/lit/llvm/masked_interleaved.ll | 4 +++- .../vecz/test/lit/llvm/masked_interleaved_as_scatter.ll | 4 +++- 17 files changed, 61 insertions(+), 25 deletions(-) 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 }