diff --git a/include/warpforth/Dialect/Forth/ForthOps.td b/include/warpforth/Dialect/Forth/ForthOps.td
index 5d36f64..973fba6 100644
--- a/include/warpforth/Dialect/Forth/ForthOps.td
+++ b/include/warpforth/Dialect/Forth/ForthOps.td
@@ -11,6 +11,19 @@ include "warpforth/Dialect/Forth/ForthDialect.td"
 include "mlir/Interfaces/ControlFlowInterfaces.td"
 include "mlir/Interfaces/SideEffectInterfaces.td"
 
+//===----------------------------------------------------------------------===//
+// Base class for standard stack-to-stack operations.
+//===----------------------------------------------------------------------===//
+
+class Forth_StackOpBase<string mnemonic, list<Trait> traits = []>
+    : Forth_Op<mnemonic, !listconcat([Pure], traits)> {
+  let arguments = (ins Forth_StackType:$input_stack);
+  let results = (outs Forth_StackType:$output_stack);
+  let assemblyFormat = [{
+    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
+  }];
+}
+
 //===----------------------------------------------------------------------===//
 // Stack manipulation operations.
 //===----------------------------------------------------------------------===//
@@ -28,140 +41,77 @@ def Forth_StackOp : Forth_Op<"stack", [Pure]> {
   }];
 }
 
-def Forth_DupOp : Forth_Op<"dup", [Pure]> {
+def Forth_DupOp : Forth_StackOpBase<"dup"> {
   let summary = "Duplicate top stack element";
   let description = [{
     Duplicates the top element of the stack.
     Forth semantics: ( a -- a a )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_DropOp : Forth_Op<"drop", [Pure]> {
+def Forth_DropOp : Forth_StackOpBase<"drop"> {
   let summary = "Remove top stack element";
   let description = [{
     Removes the top element from the stack.
     Forth semantics: ( a -- )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_SwapOp : Forth_Op<"swap", [Pure]> {
+def Forth_SwapOp : Forth_StackOpBase<"swap"> {
   let summary = "Swap top two stack elements";
   let description = [{
     Swaps the top two elements of the stack.
     Forth semantics: ( a b -- b a )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_OverOp : Forth_Op<"over", [Pure]> {
+def Forth_OverOp : Forth_StackOpBase<"over"> {
   let summary = "Copy second element to top";
   let description = [{
     Copies the second element to the top of the stack.
     Forth semantics: ( a b -- a b a )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_RotOp : Forth_Op<"rot", [Pure]> {
+def Forth_RotOp : Forth_StackOpBase<"rot"> {
   let summary = "Rotate top three elements";
   let description = [{
     Rotates the top three elements of the stack.
     Forth semantics: ( a b c -- b c a )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_NipOp : Forth_Op<"nip", [Pure]> {
+def Forth_NipOp : Forth_StackOpBase<"nip"> {
   let summary = "Remove second stack element";
   let description = [{
     Removes the second element from the stack, keeping the top.
     Forth semantics: ( a b -- b )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_TuckOp : Forth_Op<"tuck", [Pure]> {
+def Forth_TuckOp : Forth_StackOpBase<"tuck"> {
   let summary = "Copy top element before second";
   let description = [{
     Copies the top element and inserts it before the second element.
     Forth semantics: ( a b -- b a b )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_PickOp : Forth_Op<"pick", [Pure]> {
+def Forth_PickOp : Forth_StackOpBase<"pick"> {
   let summary = "Copy nth element to top";
   let description = [{
     Pops n from the stack, then copies the nth element to the top.
     Forth semantics: ( xn ... x0 n -- xn ... x0 xn )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_RollOp : Forth_Op<"roll", [Pure]> {
+def Forth_RollOp : Forth_StackOpBase<"roll"> {
   let summary = "Rotate nth element to top";
   let description = [{
     Pops n from the stack, then rotates the nth element to the top,
     shifting elements above it down.
     Forth semantics: ( xn ... x0 n -- xn-1 ... x0 xn )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
 //===----------------------------------------------------------------------===//
@@ -187,209 +137,118 @@ def Forth_LiteralOp : Forth_Op<"literal", [Pure]> {
 // Arithmetic operations.
 //===----------------------------------------------------------------------===//
 
-def Forth_AddOp : Forth_Op<"add", [Pure]> {
+def Forth_AddOp : Forth_StackOpBase<"add"> {
   let summary = "Add top two stack elements";
   let description = [{
     Pops the top two elements, adds them, and pushes the result.
     Forth semantics: ( a b -- a+b )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_SubOp : Forth_Op<"sub", [Pure]> {
+def Forth_SubOp : Forth_StackOpBase<"sub"> {
   let summary = "Subtract top two stack elements";
   let description = [{
     Pops the top two elements, subtracts them (a - b), and pushes the result.
     Forth semantics: ( a b -- a-b )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_MulOp : Forth_Op<"mul", [Pure]> {
+def Forth_MulOp : Forth_StackOpBase<"mul"> {
   let summary = "Multiply top two stack elements";
   let description = [{
     Pops the top two elements, multiplies them, and pushes the result.
     Forth semantics: ( a b -- a*b )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_DivOp : Forth_Op<"div", [Pure]> {
+def Forth_DivOp : Forth_StackOpBase<"div"> {
   let summary = "Divide top two stack elements";
   let description = [{
     Pops the top two elements, divides them (a / b), and pushes the result.
     Forth semantics: ( a b -- a/b )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_ModOp : Forth_Op<"mod", [Pure]> {
+def Forth_ModOp : Forth_StackOpBase<"mod"> {
   let summary = "Modulo of top two stack elements";
   let description = [{
     Pops the top two elements, computes modulo (a % b), and pushes the result.
     Forth semantics: ( a b -- a%b )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
 //===----------------------------------------------------------------------===//
 // Bitwise operations.
 //===----------------------------------------------------------------------===//
 
-def Forth_AndOp : Forth_Op<"and", [Pure]> {
+def Forth_AndOp : Forth_StackOpBase<"and"> {
   let summary = "Bitwise AND of top two stack elements";
   let description = [{
     Pops the top two elements, performs bitwise AND, and pushes the result.
     Forth semantics: ( a b -- a&b )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_OrOp : Forth_Op<"or", [Pure]> {
+def Forth_OrOp : Forth_StackOpBase<"or"> {
   let summary = "Bitwise OR of top two stack elements";
   let description = [{
     Pops the top two elements, performs bitwise OR, and pushes the result.
     Forth semantics: ( a b -- a|b )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_XorOp : Forth_Op<"xor", [Pure]> {
+def Forth_XorOp : Forth_StackOpBase<"xor"> {
   let summary = "Bitwise XOR of top two stack elements";
   let description = [{
     Pops the top two elements, performs bitwise XOR, and pushes the result.
     Forth semantics: ( a b -- a^b )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_NotOp : Forth_Op<"not", [Pure]> {
+def Forth_NotOp : Forth_StackOpBase<"not"> {
   let summary = "Bitwise NOT of top stack element";
   let description = [{
     Pops the top element, performs bitwise NOT (complement), and pushes the result.
     Forth semantics: ( a -- ~a )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_LshiftOp : Forth_Op<"lshift", [Pure]> {
+def Forth_LshiftOp : Forth_StackOpBase<"lshift"> {
   let summary = "Left shift";
   let description = [{
     Pops shift amount and value, shifts value left, and pushes the result.
     Forth semantics: ( a n -- a<<n )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_RshiftOp : Forth_Op<"rshift", [Pure]> {
+def Forth_RshiftOp : Forth_StackOpBase<"rshift"> {
   let summary = "Logical right shift";
   let description = [{
     Pops shift amount and value, shifts value right (logical/unsigned), and pushes the result.
     Forth semantics: ( a n -- a>>n )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
 //===----------------------------------------------------------------------===//
 // Memory operations.
 //===----------------------------------------------------------------------===//
 
-def Forth_LoadOp : Forth_Op<"load", [Pure]> {
+def Forth_LoadOp : Forth_StackOpBase<"load"> {
   let summary = "Load value from memory buffer";
   let description = [{
     Pops an address from the stack, loads a value from the memory buffer at that address,
     and pushes the loaded value onto the stack.
     Forth semantics: ( addr -- value )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_StoreOp : Forth_Op<"store", [Pure]> {
+def Forth_StoreOp : Forth_StackOpBase<"store"> {
   let summary = "Store value to memory buffer";
   let description = [{
     Pops an address and value from the stack, stores the value to the memory buffer
     at the specified address.
     Forth semantics: ( x addr -- )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
 def Forth_ParamRefOp : Forth_Op<"param_ref", [Pure]> {
@@ -409,295 +268,181 @@ def Forth_ParamRefOp : Forth_Op<"param_ref", [Pure]> {
 // GPU indexing operations.
 //===----------------------------------------------------------------------===//
 
-def Forth_ThreadIdXOp : Forth_Op<"thread_id_x", [Pure]> {
+def Forth_ThreadIdXOp : Forth_StackOpBase<"thread_id_x"> {
   let summary = "Push x-dimension thread index";
   let description = [{
     Pushes the x-dimension thread index within the current block.
     Corresponds to CUDA's threadIdx.x.
     Forth semantics: ( -- thread_idx_x )
   }];
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_ThreadIdYOp : Forth_Op<"thread_id_y", [Pure]> {
+def Forth_ThreadIdYOp : Forth_StackOpBase<"thread_id_y"> {
   let summary = "Push y-dimension thread index";
   let description = [{
     Pushes the y-dimension thread index within the current block.
     Corresponds to CUDA's threadIdx.y.
     Forth semantics: ( -- thread_idx_y )
   }];
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_ThreadIdZOp : Forth_Op<"thread_id_z", [Pure]> {
+def Forth_ThreadIdZOp : Forth_StackOpBase<"thread_id_z"> {
   let summary = "Push z-dimension thread index";
   let description = [{
     Pushes the z-dimension thread index within the current block.
     Corresponds to CUDA's threadIdx.z.
     Forth semantics: ( -- thread_idx_z )
   }];
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_BlockIdXOp : Forth_Op<"block_id_x", [Pure]> {
+def Forth_BlockIdXOp : Forth_StackOpBase<"block_id_x"> {
   let summary = "Push x-dimension block index";
   let description = [{
     Pushes the x-dimension block index within the grid.
     Corresponds to CUDA's blockIdx.x.
     Forth semantics: ( -- block_idx_x )
   }];
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_BlockIdYOp : Forth_Op<"block_id_y", [Pure]> {
+def Forth_BlockIdYOp : Forth_StackOpBase<"block_id_y"> {
   let summary = "Push y-dimension block index";
   let description = [{
     Pushes the y-dimension block index within the grid.
     Corresponds to CUDA's blockIdx.y.
     Forth semantics: ( -- block_idx_y )
   }];
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_BlockIdZOp : Forth_Op<"block_id_z", [Pure]> {
+def Forth_BlockIdZOp : Forth_StackOpBase<"block_id_z"> {
   let summary = "Push z-dimension block index";
   let description = [{
     Pushes the z-dimension block index within the grid.
     Corresponds to CUDA's blockIdx.z.
     Forth semantics: ( -- block_idx_z )
   }];
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_BlockDimXOp : Forth_Op<"block_dim_x", [Pure]> {
+def Forth_BlockDimXOp : Forth_StackOpBase<"block_dim_x"> {
   let summary = "Push x-dimension block dimension";
   let description = [{
     Pushes the number of threads in the x-dimension of a block.
     Corresponds to CUDA's blockDim.x.
     Forth semantics: ( -- block_dim_x )
   }];
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_BlockDimYOp : Forth_Op<"block_dim_y", [Pure]> {
+def Forth_BlockDimYOp : Forth_StackOpBase<"block_dim_y"> {
   let summary = "Push y-dimension block dimension";
   let description = [{
     Pushes the number of threads in the y-dimension of a block.
     Corresponds to CUDA's blockDim.y.
     Forth semantics: ( -- block_dim_y )
   }];
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_BlockDimZOp : Forth_Op<"block_dim_z", [Pure]> {
+def Forth_BlockDimZOp : Forth_StackOpBase<"block_dim_z"> {
   let summary = "Push z-dimension block dimension";
   let description = [{
     Pushes the number of threads in the z-dimension of a block.
     Corresponds to CUDA's blockDim.z.
     Forth semantics: ( -- block_dim_z )
   }];
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_GridDimXOp : Forth_Op<"grid_dim_x", [Pure]> {
+def Forth_GridDimXOp : Forth_StackOpBase<"grid_dim_x"> {
   let summary = "Push x-dimension grid dimension";
   let description = [{
     Pushes the number of blocks in the x-dimension of the grid.
     Corresponds to CUDA's gridDim.x.
     Forth semantics: ( -- grid_dim_x )
   }];
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_GridDimYOp : Forth_Op<"grid_dim_y", [Pure]> {
+def Forth_GridDimYOp : Forth_StackOpBase<"grid_dim_y"> {
   let summary = "Push y-dimension grid dimension";
   let description = [{
     Pushes the number of blocks in the y-dimension of the grid.
     Corresponds to CUDA's gridDim.y.
     Forth semantics: ( -- grid_dim_y )
   }];
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_GridDimZOp : Forth_Op<"grid_dim_z", [Pure]> {
+def Forth_GridDimZOp : Forth_StackOpBase<"grid_dim_z"> {
   let summary = "Push z-dimension grid dimension";
   let description = [{
     Pushes the number of blocks in the z-dimension of the grid.
     Corresponds to CUDA's gridDim.z.
     Forth semantics: ( -- grid_dim_z )
   }];
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_GlobalIdOp : Forth_Op<"global_id", [Pure]> {
+def Forth_GlobalIdOp : Forth_StackOpBase<"global_id"> {
   let summary = "Push 1D global thread ID";
   let description = [{
     Computes and pushes the 1D global thread ID.
     Equivalent to blockIdx.x * blockDim.x + threadIdx.x.
     Forth semantics: ( -- global_id )
   }];
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
 //===----------------------------------------------------------------------===//
 // Comparison operations.
 //===----------------------------------------------------------------------===//
 
-def Forth_EqOp : Forth_Op<"eq", [Pure]> {
+def Forth_EqOp : Forth_StackOpBase<"eq"> {
   let summary = "Test equality of top two stack elements";
   let description = [{
     Pops two values, pushes -1 (true) if equal, 0 (false) otherwise.
     Forth semantics: ( a b -- flag )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_LtOp : Forth_Op<"lt", [Pure]> {
+def Forth_LtOp : Forth_StackOpBase<"lt"> {
   let summary = "Test less-than of top two stack elements";
   let description = [{
     Pops two values, pushes -1 (true) if a < b, 0 (false) otherwise.
     Forth semantics: ( a b -- flag )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_GtOp : Forth_Op<"gt", [Pure]> {
+def Forth_GtOp : Forth_StackOpBase<"gt"> {
   let summary = "Test greater-than of top two stack elements";
   let description = [{
     Pops two values, pushes -1 (true) if a > b, 0 (false) otherwise.
     Forth semantics: ( a b -- flag )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_NeOp : Forth_Op<"ne", [Pure]> {
+def Forth_NeOp : Forth_StackOpBase<"ne"> {
   let summary = "Test inequality of top two stack elements";
   let description = [{
     Pops two values, pushes -1 (true) if not equal, 0 (false) otherwise.
     Forth semantics: ( a b -- flag )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_LeOp : Forth_Op<"le", [Pure]> {
+def Forth_LeOp : Forth_StackOpBase<"le"> {
   let summary = "Test less-than-or-equal of top two stack elements";
   let description = [{
     Pops two values, pushes -1 (true) if a <= b, 0 (false) otherwise.
     Forth semantics: ( a b -- flag )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_GeOp : Forth_Op<"ge", [Pure]> {
+def Forth_GeOp : Forth_StackOpBase<"ge"> {
   let summary = "Test greater-than-or-equal of top two stack elements";
   let description = [{
     Pops two values, pushes -1 (true) if a >= b, 0 (false) otherwise.
     Forth semantics: ( a b -- flag )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
-def Forth_ZeroEqOp : Forth_Op<"zero_eq", [Pure]> {
+def Forth_ZeroEqOp : Forth_StackOpBase<"zero_eq"> {
   let summary = "Test if top of stack is zero";
   let description = [{
     Pops one value, pushes -1 (true) if zero, 0 (false) otherwise.
     Forth semantics: ( a -- flag )
   }];
-
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
 //===----------------------------------------------------------------------===//
@@ -784,14 +529,9 @@ def Forth_BeginWhileRepeatOp : Forth_Op<"begin_while_repeat",
   let hasCustomAssemblyFormat = 1;
 }
 
-def Forth_LoopIndexOp : Forth_Op<"loop_index", [Pure]> {
+def Forth_LoopIndexOp : Forth_StackOpBase<"loop_index"> {
   let summary = "Push loop index onto stack (I word)";
   let description = [{ Only valid inside a forth.do_loop body. ( -- i ) }];
-  let arguments = (ins Forth_StackType:$input_stack);
-  let results = (outs Forth_StackType:$output_stack);
-  let assemblyFormat = [{
-    $input_stack attr-dict `:` type($input_stack) `->` type($output_stack)
-  }];
 }
 
 //===----------------------------------------------------------------------===//
diff --git a/lib/Conversion/ForthToMemRef/ForthToMemRef.cpp b/lib/Conversion/ForthToMemRef/ForthToMemRef.cpp
index 00ca8e1..ed2e2fd 100644
--- a/lib/Conversion/ForthToMemRef/ForthToMemRef.cpp
+++ b/lib/Conversion/ForthToMemRef/ForthToMemRef.cpp
@@ -48,6 +48,15 @@ class ForthToMemRefTypeConverter : public TypeConverter {
   }
 };
 
+/// Push an i64 value onto the stack. Returns the new stack pointer.
+static Value pushValue(Location loc, ConversionPatternRewriter &rewriter,
+                       Value memref, Value stackPtr, Value value) {
+  Value one = rewriter.create<arith::ConstantIndexOp>(loc, 1);
+  Value newSP = rewriter.create<arith::AddIOp>(loc, stackPtr, one);
+  rewriter.create<memref::StoreOp>(loc, value, memref, newSP);
+  return newSP;
+}
+
 /// Conversion pattern for forth.stack operation.
 /// Allocates a memref<256xi64> on the stack and initializes SP to 0.
 struct StackOpConversion : public OpConversionPattern<forth::StackOp> {
@@ -88,14 +97,9 @@ struct LiteralOpConversion : public OpConversionPattern<forth::LiteralOp> {
     Value memref = inputStack[0];
     Value stackPtr = inputStack[1];
 
-    // Increment stack pointer
-    Value one = rewriter.create<arith::ConstantIndexOp>(loc, 1);
-    Value newSP = rewriter.create<arith::AddIOp>(loc, stackPtr, one);
-
-    // Store literal value at new SP position
     Value literalValue = rewriter.create<arith::ConstantOp>(
         loc, rewriter.getI64Type(), op.getValueAttr());
-    rewriter.create<memref::StoreOp>(loc, literalValue, memref, newSP);
+    Value newSP = pushValue(loc, rewriter, memref, stackPtr, literalValue);
 
     rewriter.replaceOpWithMultiple(op, {{memref, newSP}});
     return success();
@@ -117,15 +121,9 @@ struct DupOpConversion : public OpConversionPattern<forth::DupOp> {
     Value memref = inputStack[0];
     Value stackPtr = inputStack[1];
 
-    // Load value at current SP
+    // Load value at current SP and push duplicate
     Value topValue = rewriter.create<memref::LoadOp>(loc, memref, stackPtr);
-
-    // Increment SP
-    Value one = rewriter.create<arith::ConstantIndexOp>(loc, 1);
-    Value newSP = rewriter.create<arith::AddIOp>(loc, stackPtr, one);
-
-    // Store duplicate at new SP
-    rewriter.create<memref::StoreOp>(loc, topValue, memref, newSP);
+    Value newSP = pushValue(loc, rewriter, memref, stackPtr, topValue);
 
     rewriter.replaceOpWithMultiple(op, {{memref, newSP}});
     return success();
@@ -202,16 +200,11 @@ struct OverOpConversion : public OpConversionPattern<forth::OverOp> {
     Value memref = inputStack[0];
     Value stackPtr = inputStack[1];
 
-    // Load second element (SP - 1)
+    // Load second element (SP - 1) and push copy
     Value one = rewriter.create<arith::ConstantIndexOp>(loc, 1);
     Value spMinus1 = rewriter.create<arith::SubIOp>(loc, stackPtr, one);
     Value secondValue = rewriter.create<memref::LoadOp>(loc, memref, spMinus1);
-
-    // Increment SP
-    Value newSP = rewriter.create<arith::AddIOp>(loc, stackPtr, one);
-
-    // Store second value at new SP
-    rewriter.create<memref::StoreOp>(loc, secondValue, memref, newSP);
+    Value newSP = pushValue(loc, rewriter, memref, stackPtr, secondValue);
 
     rewriter.replaceOpWithMultiple(op, {{memref, newSP}});
     return success();
@@ -623,9 +616,7 @@ struct ParamRefOpConversion : public OpConversionPattern<forth::ParamRefOp> {
         loc, rewriter.getI64Type(), ptrIndex);
 
     // Push onto stack
-    Value one = rewriter.create<arith::ConstantIndexOp>(loc, 1);
-    Value newSP = rewriter.create<arith::AddIOp>(loc, stackPtr, one);
-    rewriter.create<memref::StoreOp>(loc, ptrI64, memref, newSP);
+    Value newSP = pushValue(loc, rewriter, memref, stackPtr, ptrI64);
 
     rewriter.replaceOpWithMultiple(op, {{memref, newSP}});
     return success();
@@ -723,18 +714,12 @@ struct IntrinsicOpConversion : public OpConversionPattern<ForthOp> {
     Value memref = inputStack[0];
     Value stackPtr = inputStack[1];
 
-    // Increment stack pointer
-    Value one = rewriter.create<arith::ConstantIndexOp>(loc, 1);
-    Value newSP = rewriter.create<arith::AddIOp>(loc, stackPtr, one);
-
-    // Create intrinsic op
+    // Create intrinsic op and push onto stack
     Value intrinsicValue = rewriter.create<forth::IntrinsicOp>(
         loc, rewriter.getIndexType(), rewriter.getStringAttr(intrinsicName));
     Value intrinsicI64 = rewriter.create<arith::IndexCastOp>(
         loc, rewriter.getI64Type(), intrinsicValue);
-
-    // Store at new SP
-    rewriter.create<memref::StoreOp>(loc, intrinsicI64, memref, newSP);
+    Value newSP = pushValue(loc, rewriter, memref, stackPtr, intrinsicI64);
 
     rewriter.replaceOpWithMultiple(op, {{memref, newSP}});
     return success();
@@ -771,11 +756,7 @@ struct GlobalIdOpConversion : public OpConversionPattern<forth::GlobalIdOp> {
     Value globalId = rewriter.create<arith::AddIOp>(loc, product, tidX);
     Value globalIdI64 = rewriter.create<arith::IndexCastOp>(
         loc, rewriter.getI64Type(), globalId);
-
-    // Increment SP and store result
-    Value one = rewriter.create<arith::ConstantIndexOp>(loc, 1);
-    Value newSP = rewriter.create<arith::AddIOp>(loc, stackPtr, one);
-    rewriter.create<memref::StoreOp>(loc, globalIdI64, memref, newSP);
+    Value newSP = pushValue(loc, rewriter, memref, stackPtr, globalIdI64);
 
     rewriter.replaceOpWithMultiple(op, {{memref, newSP}});
     return success();
@@ -1100,10 +1081,8 @@ struct LoopIndexOpConversion : public OpConversionPattern<forth::LoopIndexOp> {
     Value ivI64 =
         rewriter.create<arith::IndexCastOp>(loc, rewriter.getI64Type(), iv);
 
-    // Push onto stack: SP += 1, store at new SP
-    Value one = rewriter.create<arith::ConstantIndexOp>(loc, 1);
-    Value newSP = rewriter.create<arith::AddIOp>(loc, stackPtr, one);
-    rewriter.create<memref::StoreOp>(loc, ivI64, memref, newSP);
+    // Push onto stack
+    Value newSP = pushValue(loc, rewriter, memref, stackPtr, ivI64);
 
     rewriter.replaceOpWithMultiple(op, {{memref, newSP}});
     return success();
diff --git a/test/Conversion/ForthToMemRef/control-flow.mlir b/test/Conversion/ForthToMemRef/control-flow.mlir
index 03c5834..8c77c64 100644
--- a/test/Conversion/ForthToMemRef/control-flow.mlir
+++ b/test/Conversion/ForthToMemRef/control-flow.mlir
@@ -12,16 +12,16 @@
 
 // Then branch: drop (subi) + literal push + yield
 // CHECK:   arith.subi
-// CHECK:   arith.addi
 // CHECK:   arith.constant 42 : i64
+// CHECK:   arith.addi
 // CHECK:   memref.store
 // CHECK:   scf.yield %{{.*}} : index
 
 // Else branch: drop (subi) + literal push + yield
 // CHECK: } else {
 // CHECK:   arith.subi
-// CHECK:   arith.addi
 // CHECK:   arith.constant 99 : i64
+// CHECK:   arith.addi
 // CHECK:   memref.store
 // CHECK:   scf.yield %{{.*}} : index
 // CHECK: }
diff --git a/test/Conversion/ForthToMemRef/literal.mlir b/test/Conversion/ForthToMemRef/literal.mlir
index 2b86093..d6dc84f 100644
--- a/test/Conversion/ForthToMemRef/literal.mlir
+++ b/test/Conversion/ForthToMemRef/literal.mlir
@@ -3,9 +3,9 @@
 // CHECK-LABEL: func.func private @main
 // CHECK: memref.alloca() : memref<256xi64>
 // CHECK: arith.constant 0 : index
+// CHECK: arith.constant 42 : i64
 // CHECK: arith.constant 1 : index
 // CHECK: arith.addi %{{.*}}, %{{.*}} : index
-// CHECK: arith.constant 42 : i64
 // CHECK: memref.store %{{.*}}, %{{.*}}[%{{.*}}] : memref<256xi64>
 
 module {