diff --git a/include/warpforth/Dialect/Forth/ForthOps.td b/include/warpforth/Dialect/Forth/ForthOps.td
index f13f789..15a7d4e 100644
--- a/include/warpforth/Dialect/Forth/ForthOps.td
+++ b/include/warpforth/Dialect/Forth/ForthOps.td
@@ -384,6 +384,20 @@ def Forth_GlobalIdOp : Forth_StackOpBase<"global_id"> {
   }];
 }
 
+//===----------------------------------------------------------------------===//
+// Synchronization operations.
+//===----------------------------------------------------------------------===//
+
+def Forth_BarrierOp : Forth_Op<"barrier", []> {
+  let summary = "GPU thread synchronization barrier";
+  let description = [{
+    Synchronizes all threads in a block.
+    Corresponds to CUDA's __syncthreads().
+    Forth semantics: ( -- )
+  }];
+  let assemblyFormat = "attr-dict";
+}
+
 //===----------------------------------------------------------------------===//
 // Comparison operations.
 //===----------------------------------------------------------------------===//
diff --git a/lib/Conversion/ForthToGPU/ForthToGPU.cpp b/lib/Conversion/ForthToGPU/ForthToGPU.cpp
index 17d7be4..84557c0 100644
--- a/lib/Conversion/ForthToGPU/ForthToGPU.cpp
+++ b/lib/Conversion/ForthToGPU/ForthToGPU.cpp
@@ -78,19 +78,32 @@ struct IntrinsicOpConversion : public OpConversionPattern<forth::IntrinsicOp> {
   }
 };
 
+/// Conversion pattern for forth.barrier to gpu.barrier.
+struct BarrierOpConversion : public OpConversionPattern<forth::BarrierOp> {
+  using OpConversionPattern::OpConversionPattern;
+
+  LogicalResult
+  matchAndRewrite(forth::BarrierOp op, OpAdaptor adaptor,
+                  ConversionPatternRewriter &rewriter) const override {
+    rewriter.replaceOpWithNewOp<gpu::BarrierOp>(op);
+    return success();
+  }
+};
+
 /// Applies conversion patterns to a function.
 static void applyConversionPatterns(Operation *op, MLIRContext *context) {
   ConversionTarget target(*context);
 
-  // Mark forth.intrinsic as illegal - it must be converted
+  // Mark forth.intrinsic and forth.barrier as illegal - they must be converted
   target.addIllegalOp<forth::IntrinsicOp>();
+  target.addIllegalOp<forth::BarrierOp>();
 
   // GPU, Arith, Memref and LLVM dialect ops are legal
   target.addLegalDialect<gpu::GPUDialect, arith::ArithDialect,
                          memref::MemRefDialect, LLVM::LLVMDialect>();
 
   RewritePatternSet patterns(context);
-  patterns.add<IntrinsicOpConversion>(context);
+  patterns.add<IntrinsicOpConversion, BarrierOpConversion>(context);
 
   if (failed(applyPartialConversion(op, target, std::move(patterns)))) {
     return;
diff --git a/lib/Conversion/ForthToMemRef/ForthToMemRef.cpp b/lib/Conversion/ForthToMemRef/ForthToMemRef.cpp
index f0031d1..080fe16 100644
--- a/lib/Conversion/ForthToMemRef/ForthToMemRef.cpp
+++ b/lib/Conversion/ForthToMemRef/ForthToMemRef.cpp
@@ -976,8 +976,9 @@ struct ConvertForthToMemRefPass
     target.addLegalDialect<memref::MemRefDialect, arith::ArithDialect,
                            LLVM::LLVMDialect, cf::ControlFlowDialect>();
 
-    // Mark IntrinsicOp as legal (to be lowered later)
+    // Mark IntrinsicOp and BarrierOp as legal (to be lowered later)
     target.addLegalOp<forth::IntrinsicOp>();
+    target.addLegalOp<forth::BarrierOp>();
 
     // Use dynamic legality for func operations to ensure they're properly
     // converted
diff --git a/lib/Translation/ForthToMLIR/ForthToMLIR.cpp b/lib/Translation/ForthToMLIR/ForthToMLIR.cpp
index 766e21c..b176843 100644
--- a/lib/Translation/ForthToMLIR/ForthToMLIR.cpp
+++ b/lib/Translation/ForthToMLIR/ForthToMLIR.cpp
@@ -526,6 +526,9 @@ Value ForthParser::emitOperation(StringRef word, Value inputStack,
   } else if (word == "GLOBAL-ID") {
     return builder.create<forth::GlobalIdOp>(loc, stackType, inputStack)
         .getResult();
+  } else if (word == "BARRIER") {
+    builder.create<forth::BarrierOp>(loc);
+    return inputStack;
   } else if (word == "=") {
     return builder.create<forth::EqOp>(loc, stackType, inputStack).getResult();
   } else if (word == "<") {
diff --git a/test/Conversion/ForthToMemRef/barrier.mlir b/test/Conversion/ForthToMemRef/barrier.mlir
new file mode 100644
index 0000000..68799a6
--- /dev/null
+++ b/test/Conversion/ForthToMemRef/barrier.mlir
@@ -0,0 +1,9 @@
+// RUN: %warpforth-opt --convert-forth-to-memref %s | %FileCheck %s
+// CHECK-LABEL: func.func private @main
+// CHECK: forth.barrier
+func.func private @main() {
+  %0 = forth.stack !forth.stack
+  forth.barrier
+  forth.drop %0 : !forth.stack -> !forth.stack
+  return
+}
diff --git a/test/Pipeline/barrier.forth b/test/Pipeline/barrier.forth
new file mode 100644
index 0000000..ef1c2b8
--- /dev/null
+++ b/test/Pipeline/barrier.forth
@@ -0,0 +1,10 @@
+\ RUN: %warpforth-translate --forth-to-mlir %s | %warpforth-opt --warpforth-pipeline | %FileCheck %s
+\ RUN: %warpforth-translate --forth-to-mlir %s | %warpforth-opt --convert-forth-to-memref --convert-forth-to-gpu | %FileCheck %s --check-prefix=MID
+\ CHECK: gpu.binary @warpforth_module
+\ MID: gpu.func @main
+\ MID: gpu.barrier
+\ MID: gpu.return
+\! kernel main
+\! param DATA i64[256]
+\! shared SCRATCH i64[256]
+GLOBAL-ID CELLS SCRATCH + @ BARRIER GLOBAL-ID CELLS DATA + !
diff --git a/test/Translation/Forth/barrier.forth b/test/Translation/Forth/barrier.forth
new file mode 100644
index 0000000..963c38d
--- /dev/null
+++ b/test/Translation/Forth/barrier.forth
@@ -0,0 +1,5 @@
+\ RUN: %warpforth-translate --forth-to-mlir %s | %FileCheck %s
+\ CHECK: forth.barrier
+\! kernel main
+\! param DATA i64[256]
+GLOBAL-ID CELLS DATA + @ BARRIER DROP