Merge pull request #274 from frasercrmck/vecz-nounwind-builtins

[vecz] Add nounwind/norecurse attributes to internal vecz builtins
uxlfoundation · Jan 3, 2024 · ce8a356 · ce8a356
2 parents f6b5f2a + cfb4352
commit ce8a356
Show file tree

Hide file tree

Showing 17 changed files with 61 additions and 25 deletions.
diff --git 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<Function>(
         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
@@ -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
@@ -65,18 +65,20 @@ 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
 ; CHECK-NEXT:  br label %[[EXIT]]
 ; 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
@@ -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(<vscale x 4 x double> %0, ptr addrspace(1) %1) {
+; CHECK: define void @__vecz_b_interleaved_store8_4_u5nxv4du3ptrU3AS1(<vscale x 4 x double> %0, ptr addrspace(1) %1) [[ATTRS:#[0-9]+]] {
 ; CHECK: entry:
 ; CHECK:   %BroadcastAddr.splatinsert = insertelement <vscale x 4 x ptr addrspace(1)> poison, ptr addrspace(1) %1, {{i32|i64}} 0
 ; CHECK:   %BroadcastAddr.splat = shufflevector <vscale x 4 x ptr addrspace(1)> %BroadcastAddr.splatinsert, <vscale x 4 x ptr addrspace(1)> poison, <vscale x 4 x i32> 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(<vscale x 4 x double> %0, <vscale x 4 x ptr addrspace(1)> %4, i32 immarg 8, <vscale x 4 x i1> shufflevector (<vscale x 4 x i1> insertelement (<vscale x 4 x i1> poison, i1 true, {{i32|i64}} 0), <vscale x 4 x i1> poison, <vscale x 4 x i32> 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
@@ -44,7 +44,7 @@ entry:
 
 declare i64 @__mux_get_global_id(i32)
 
-; CHECK: define void @__vecz_b_interleaved_store4_V_u5nxv4ju3ptrU3AS1(<vscale x 4 x i32> [[ARG0:%.*]], ptr addrspace(1) [[ARG1:%.*]], i64 [[ARG2:%.*]]) {
+; CHECK: define void @__vecz_b_interleaved_store4_V_u5nxv4ju3ptrU3AS1(<vscale x 4 x i32> [[ARG0:%.*]], ptr addrspace(1) [[ARG1:%.*]], i64 [[ARG2:%.*]]) [[ATTRS:#[0-9]+]] {
 ; CHECK-NEXT: entry:
 ; CHECK-NEXT: [[TMP0:%.*]] = insertelement <vscale x 4 x ptr addrspace(1)> poison, ptr addrspace(1) [[ARG1]], {{i32|i64}} 0
 ; CHECK-NEXT: [[TMP1:%.*]] = shufflevector <vscale x 4 x ptr addrspace(1)> [[TMP0]], <vscale x 4 x ptr addrspace(1)> poison, <vscale x 4 x i32> zeroinitializer
@@ -53,8 +53,9 @@ declare i64 @__mux_get_global_id(i32)
 ; CHECK-NEXT: [[TMP4:%.*]] = call <vscale x 4 x i64> @llvm.experimental.stepvector.nxv4i64()
 ; CHECK-NEXT: [[TMP5:%.*]] = mul <vscale x 4 x i64> [[TMP3]], [[TMP4]]
 ; CHECK-NEXT: [[TMP6:%.*]] = getelementptr i32, <vscale x 4 x ptr addrspace(1)> [[TMP1]], <vscale x 4 x i64> [[TMP5]]
-; CHECK-NEXT: call void @llvm.masked.scatter.nxv4i32.nxv4p1(<vscale x 4 x i32> [[ARG0]], <vscale x 4 x ptr addrspace(1)> [[TMP6]], i32 immarg 4, <vscale x 4 x i1> shufflevector (<vscale x 4 x i1> insertelement (<vscale x 4 x i1> poison, i1 true, {{i32|i64}} 0), <vscale x 4 x i1> poison, <vscale x 4 x i32> zeroinitializer)) #[[ATTRS:[0-9]+]]
+; CHECK-NEXT: call void @llvm.masked.scatter.nxv4i32.nxv4p1(<vscale x 4 x i32> [[ARG0]], <vscale x 4 x ptr addrspace(1)> [[TMP6]], i32 immarg 4, <vscale x 4 x i1> shufflevector (<vscale x 4 x i1> insertelement (<vscale x 4 x i1> poison, i1 true, {{i32|i64}} 0), <vscale x 4 x i1> poison, <vscale x 4 x i32> 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
@@ -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 <vscale x 4 x double> @__vecz_b_masked_interleaved_load8_vp_4_u5nxv4du3ptrU3AS1u5nxv4bj(ptr addrspace(1){{( %0)?}}, <vscale x 4 x i1>{{( %1)?}}, i32{{( %2)?}}) {
+; CHECK: define <vscale x 4 x double> @__vecz_b_masked_interleaved_load8_vp_4_u5nxv4du3ptrU3AS1u5nxv4bj(ptr addrspace(1){{( %0)?}}, <vscale x 4 x i1>{{( %1)?}}, i32{{( %2)?}}) [[ATTRS:#[0-9]+]] {
 ; CHECK: entry:
 ; CHECK:   %BroadcastAddr.splatinsert = insertelement <vscale x 4 x ptr addrspace(1)> poison, ptr addrspace(1) %0, {{i32|i64}} 0
 ; CHECK:   %BroadcastAddr.splat = shufflevector <vscale x 4 x ptr addrspace(1)> %BroadcastAddr.splatinsert, <vscale x 4 x ptr addrspace(1)> poison, <vscale x 4 x i32> 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(<vscale x 4 x double>{{( %0)?}}, ptr addrspace(1){{( %1)?}}, <vscale x 4 x i1>{{( %2)?}}, i32{{( %3)?}})
+; CHECK: define void @__vecz_b_masked_interleaved_store8_vp_4_u5nxv4du3ptrU3AS1u5nxv4bj(<vscale x 4 x double>{{( %0)?}}, ptr addrspace(1){{( %1)?}}, <vscale x 4 x i1>{{( %2)?}}, i32{{( %3)?}}) [[ATTRS]]
 ; CHECK: entry:
 ; CHECK:  %BroadcastAddr.splatinsert = insertelement <vscale x 4 x ptr addrspace(1)> poison, ptr addrspace(1) %1, {{i32|i64}} 0
 ; CHECK:  %BroadcastAddr.splat = shufflevector <vscale x 4 x ptr addrspace(1)> %BroadcastAddr.splatinsert, <vscale x 4 x ptr addrspace(1)> poison, <vscale x 4 x i32> 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(<vscale x 4 x double> %0, <vscale x 4 x ptr addrspace(1)> %6, <vscale x 4 x i1> %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
@@ -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
@@ -55,11 +55,11 @@ entry:
 ; CHECK: [[v:%.*]] = call <vscale x 4 x i32> @llvm.vp.load.nxv4i32.p1(
 ; CHECK: call void @__vecz_b_masked_scatter_store4_vp_u5nxv4ju14nxv4u3ptrU3AS1u5nxv4bj(<vscale x 4 x i32> [[v]],
 
-; CHECK: define <vscale x 4 x i32> @__vecz_b_masked_gather_load4_vp_u5nxv4ju14nxv4u3ptrU3AS1u5nxv4bj(<vscale x 4 x ptr addrspace(1)> %0, <vscale x 4 x i1> %1, i32 %2) {
+; CHECK: define <vscale x 4 x i32> @__vecz_b_masked_gather_load4_vp_u5nxv4ju14nxv4u3ptrU3AS1u5nxv4bj(<vscale x 4 x ptr addrspace(1)> %0, <vscale x 4 x i1> %1, i32 %2) [[ATTRS:#[0-9]+]] {
 ; CHECK:   %3 = call <vscale x 4 x i32> @llvm.vp.gather.nxv4i32.nxv4p1(<vscale x 4 x ptr addrspace(1)> %0, <vscale x 4 x i1> %1, i32 %2)
 ; CHECK:   ret <vscale x 4 x i32> %3
 
-; CHECK: define void @__vecz_b_masked_scatter_store4_vp_u5nxv4ju14nxv4u3ptrU3AS1u5nxv4bj(<vscale x 4 x i32> %0, <vscale x 4 x ptr addrspace(1)> %1, <vscale x 4 x i1> %2, i32 %3) {
+; CHECK: define void @__vecz_b_masked_scatter_store4_vp_u5nxv4ju14nxv4u3ptrU3AS1u5nxv4bj(<vscale x 4 x i32> %0, <vscale x 4 x ptr addrspace(1)> %1, <vscale x 4 x i1> %2, i32 %3) [[ATTRS]] {
 ; CHECK: entry:
 ; CHECK:   call void @llvm.vp.scatter.nxv4i32.nxv4p1(<vscale x 4 x i32> %0, <vscale x 4 x ptr addrspace(1)> %1, <vscale x 4 x i1> %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
@@ -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
@@ -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> <i1 true, i1 true, i1 true, i1 true>,
 ; 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
@@ -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> <i1 true, i1 true, i1 true, i1 true>, <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
@@ -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> <i1 true, i1 true, i1 true, i1 true>)
 ; 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
@@ -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> <i1 true, i1 true, i1 true, i1 true>)
 ; 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
@@ -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
@@ -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
@@ -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> <i64 0, i64 2, i64 4, i64 6>
 ; 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
@@ -65,11 +65,13 @@ 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
 ; CHECK: %[[BROADCASTADDRSPLAT:.+]] = shufflevector <4 x ptr addrspace(1)> %[[BROADCASTADDRSPLATINSERT]], <4 x ptr addrspace(1)> {{poison|undef}}, <4 x i32> zeroinitializer
 ; CHECK: getelementptr i32, <4 x ptr addrspace(1)> %[[BROADCASTADDRSPLAT]], <4 x i64> <i64 0, i64 2, i64 4, i64 6>
 
 ; CHECK: ret void
+
+; CHECK: attributes [[ATTRS]] = { norecurse nounwind }