caio/clang-p2996
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[flang][cuda] Use NVVM op for barrier0 intrinsic (#140947)

Browse Source 
The simple form of `Barrier0Op` is available in the NVVM dialect. It is
needed to use it instead of the string version since
https://github.com/llvm/llvm-project/pull/140615
This commit is contained in:
Valentin Clement (バレンタイン クレメン)
2025-05-21 13:05:14 -07:00
committed by GitHub
parent 5ba57a81f2
commit 89d9a83b70
2 changed files with 3 additions and 8 deletions
Download Patch File Download Diff File Expand all files Collapse all files

7
flang/lib/Optimizer/Builder/IntrinsicCall.cpp
View File

@@ -8332,12 +8332,7 @@ IntrinsicLibrary::genSum(mlir::Type resultType,
// SYNCTHREADS
void IntrinsicLibrary::genSyncThreads(llvm::ArrayRef<fir::ExtendedValue> args) {
constexpr llvm::StringLiteral funcName = "llvm.nvvm.barrier0";
mlir::FunctionType funcType =
mlir::FunctionType::get(builder.getContext(), {}, {});
auto funcOp = builder.createFunction(loc, funcName, funcType);
llvm::SmallVector<mlir::Value> noArgs;
builder.create<fir::CallOp>(loc, funcOp, noArgs);
builder.create<mlir::NVVM::Barrier0Op>(loc);
}
// SYNCTHREADS_AND

4
flang/test/Lower/CUDA/cuda-device-proc.cuf
View File

@@ -49,7 +49,7 @@ attributes(global) subroutine devsub()
end
! CHECK-LABEL: func.func @_QPdevsub() attributes {cuf.proc_attr = #cuf.cuda_proc<global>}
! CHECK: fir.call @llvm.nvvm.barrier0() fastmath<contract> : () -> ()
! CHECK: nvvm.barrier0
! CHECK: fir.call @llvm.nvvm.bar.warp.sync(%c1{{.*}}) fastmath<contract> : (i32) -> ()
! CHECK: fir.call @llvm.nvvm.membar.gl() fastmath<contract> : () -> ()
! CHECK: fir.call @llvm.nvvm.membar.cta() fastmath<contract> : () -> ()
@@ -106,7 +106,7 @@ end
! CHECK-LABEL: func.func @_QPhost1()
! CHECK: cuf.kernel
! CHECK: fir.call @llvm.nvvm.barrier0() fastmath<contract> : () -> ()
! CHECK: nvvm.barrier0
! CHECK: fir.call @llvm.nvvm.bar.warp.sync(%c1{{.*}}) fastmath<contract> : (i32) -> ()
! CHECK: fir.call @llvm.nvvm.barrier0.and(%c1{{.*}}) fastmath<contract> : (i32) -> i32
! CHECK: fir.call @llvm.nvvm.barrier0.popc(%c1{{.*}}) fastmath<contract> : (i32) -> i32