Skip to content

Commit f3d62f8

Browse files
committed
Address nits comments
1 parent 00a45d2 commit f3d62f8

File tree

3 files changed

+10
-14
lines changed

3 files changed

+10
-14
lines changed

mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1236,7 +1236,7 @@ def NVVM_FenceProxyAcquireOp : NVVM_Op<"fence.proxy.acquire">,
12361236
let hasVerifier = 1;
12371237
}
12381238

1239-
def NVVM_MembarOp : NVVM_Op<"memory_barrier">,
1239+
def NVVM_MembarOp : NVVM_Op<"memory.barrier">,
12401240
Arguments<(ins MemScopeKindAttr:$scope)> {
12411241
let summary = "Memory barrier operation";
12421242
let description = [{

mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp

Lines changed: 5 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -293,20 +293,16 @@ static unsigned getUnidirectionalFenceProxyID(NVVM::ProxyKind fromProxy,
293293

294294
static unsigned getMembarIntrinsicID(NVVM::MemScopeKind scope) {
295295
switch (scope) {
296-
case NVVM::MemScopeKind::CTA: {
296+
case NVVM::MemScopeKind::CTA:
297297
return llvm::Intrinsic::nvvm_membar_cta;
298-
}
299-
case NVVM::MemScopeKind::CLUSTER: {
298+
case NVVM::MemScopeKind::CLUSTER:
300299
return llvm::Intrinsic::nvvm_fence_sc_cluster;
301-
}
302-
case NVVM::MemScopeKind::GPU: {
300+
case NVVM::MemScopeKind::GPU:
303301
return llvm::Intrinsic::nvvm_membar_gl;
304-
}
305-
case NVVM::MemScopeKind::SYS: {
302+
case NVVM::MemScopeKind::SYS:
306303
return llvm::Intrinsic::nvvm_membar_sys;
307304
}
308-
}
309-
llvm_unreachable("Unknown level for memory barrier");
305+
llvm_unreachable("Unknown scope for memory barrier");
310306
}
311307

312308
#define TCGEN05LD(SHAPE, NUM) llvm::Intrinsic::nvvm_tcgen05_ld_##SHAPE##_##NUM

mlir/test/Target/LLVMIR/nvvm/membar.mlir

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -3,12 +3,12 @@
33
// CHECK-lABEL: @memorybarrier()
44
llvm.func @memorybarrier() {
55
// CHECK: call void @llvm.nvvm.membar.cta()
6-
nvvm.memory_barrier #nvvm.mem_scope<cta>
6+
nvvm.memory.barrier #nvvm.mem_scope<cta>
77
// CHECK: call void @llvm.nvvm.fence.sc.cluster()
8-
nvvm.memory_barrier #nvvm.mem_scope<cluster>
8+
nvvm.memory.barrier #nvvm.mem_scope<cluster>
99
// CHECK: call void @llvm.nvvm.membar.gl()
10-
nvvm.memory_barrier #nvvm.mem_scope<gpu>
10+
nvvm.memory.barrier #nvvm.mem_scope<gpu>
1111
// CHECK: call void @llvm.nvvm.membar.sys()
12-
nvvm.memory_barrier #nvvm.mem_scope<sys>
12+
nvvm.memory.barrier #nvvm.mem_scope<sys>
1313
llvm.return
1414
}

0 commit comments

Comments
 (0)