Add BARRIER word for GPU thread synchronization

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.
 //===----------------------------------------------------------------------===//

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;

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

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 == "<") {

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
+}

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 + !

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