Skip to content

[mlir][nvvm] Rename mbarrier.test.wait -> mbarrier.test_wait. #96503

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 1 commit into
base: main
Choose a base branch
from
Open

[mlir][nvvm] Rename mbarrier.test.wait -> mbarrier.test_wait. #96503

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -361,7 +361,7 @@ def NVVM_MBarrierTryWaitParitySharedOp : NVVM_PTXBuilder_Op<"mbarrier.try_wait.p
}];
}

def NVVM_MBarrierTestWaitOp : NVVM_Op<"mbarrier.test.wait">,
def NVVM_MBarrierTestWaitOp : NVVM_Op<"mbarrier.test_wait">,
Results<(outs LLVM_Type:$res)>,
Arguments<(ins LLVM_AnyPointer:$addr, LLVM_Type:$state)> {
string llvmBuilder = [{
Expand All @@ -370,7 +370,7 @@ def NVVM_MBarrierTestWaitOp : NVVM_Op<"mbarrier.test.wait">,
let assemblyFormat = "$addr `,` $state attr-dict `:` type(operands) `->` type($res)";
}

def NVVM_MBarrierTestWaitSharedOp : NVVM_Op<"mbarrier.test.wait.shared">,
def NVVM_MBarrierTestWaitSharedOp : NVVM_Op<"mbarrier.test_wait.shared">,
Results<(outs LLVM_Type:$res)>,
Arguments<(ins LLVM_PointerShared:$addr, LLVM_Type:$state)> {
string llvmBuilder = [{
Expand Down
2 changes: 1 addition & 1 deletion mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -896,7 +896,7 @@ struct NVGPUMBarrierArriveNoCompleteLowering
}
};

/// Lowers `nvgpu.mbarrier.test.wait` to `nvvm.mbarrier.test.wait`
/// Lowers `nvgpu.mbarrier.test.wait` to `nvvm.mbarrier.test_wait`
struct NVGPUMBarrierTestWaitLowering
: public MBarrierBasePattern<nvgpu::MBarrierTestWaitOp> {
using MBarrierBasePattern<nvgpu::MBarrierTestWaitOp>::MBarrierBasePattern;
Expand Down
4 changes: 2 additions & 2 deletions mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -496,7 +496,7 @@ func.func @mbarrier() {

// CHECK: %[[base3:.+]] = llvm.extractvalue %[[barStr]][1] : !llvm.struct<(ptr<3>, ptr<3>, i64, array<1 x i64>, array<1 x i64>)>
// CHECK: %[[barPtr3:.+]] = llvm.getelementptr %[[base3]][%[[mid]]] : (!llvm.ptr<3>, i64) -> !llvm.ptr<3>, i64
// CHECK: nvvm.mbarrier.test.wait.shared %[[barPtr3]], %[[token]]
// CHECK: nvvm.mbarrier.test_wait.shared %[[barPtr3]], %[[token]]
%isDone = nvgpu.mbarrier.test.wait %barrier[%c0], %token : !barrierType, !tokenType

func.return
Expand Down Expand Up @@ -526,7 +526,7 @@ func.func @mbarrier_nocomplete() {

// CHECK: %[[base3:.+]] = llvm.extractvalue %[[barStr]][1] : !llvm.struct<(ptr<3>, ptr<3>, i64, array<1 x i64>, array<1 x i64>)>
// CHECK: %[[barPtr3:.+]] = llvm.getelementptr %[[base3]][%[[mid]]] : (!llvm.ptr<3>, i64) -> !llvm.ptr<3>, i64
// CHECK: nvvm.mbarrier.test.wait.shared %[[barPtr3]], %[[token]]
// CHECK: nvvm.mbarrier.test_wait.shared %[[barPtr3]], %[[token]]
%isDone = nvgpu.mbarrier.test.wait %barrier[%c0], %token : !barrierType, !tokenType

func.return
Expand Down
8 changes: 4 additions & 4 deletions mlir/test/Dialect/LLVMIR/nvvm.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -452,15 +452,15 @@ llvm.func private @mbarrier_arrive_nocomplete_shared(%barrier: !llvm.ptr<3>) {
}

llvm.func private @mbarrier_test_wait(%barrier: !llvm.ptr, %token : i64) -> i1 {
// CHECK: nvvm.mbarrier.test.wait %{{.*}}
%isComplete = nvvm.mbarrier.test.wait %barrier, %token : !llvm.ptr, i64 -> i1
// CHECK: nvvm.mbarrier.test_wait %{{.*}}
%isComplete = nvvm.mbarrier.test_wait %barrier, %token : !llvm.ptr, i64 -> i1
llvm.return %isComplete : i1
}

llvm.func private @mbarrier_test_wait_shared(%barrier: !llvm.ptr<3>, %token : i64) {
%count = nvvm.read.ptx.sreg.ntid.x : i32
// CHECK: nvvm.mbarrier.test.wait.shared %{{.*}}
%isComplete = nvvm.mbarrier.test.wait.shared %barrier, %token : !llvm.ptr<3>, i64 -> i1
// CHECK: nvvm.mbarrier.test_wait.shared %{{.*}}
%isComplete = nvvm.mbarrier.test_wait.shared %barrier, %token : !llvm.ptr<3>, i64 -> i1
llvm.return
}

Expand Down