shylie/llvm-project
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[flang][cuda] Use NVVM op for clock64 (#149223)

Browse Source
This commit is contained in:
Valentin Clement (バレンタイン クレメン) 2025-07-16 18:12:18 -07:00 committed by GitHub
parent b9adc4a59c
commit 34951f7de8
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
3 changed files with 6 additions and 15 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
flang/include/flang/Optimizer/Builder/IntrinsicCall.h
View File

@ -241,7 +241,6 @@ struct IntrinsicLibrary {
void genCFProcPointer(llvm::ArrayRef<fir::ExtendedValue>);
fir::ExtendedValue genCFunLoc(mlir::Type, llvm::ArrayRef<fir::ExtendedValue>);
fir::ExtendedValue genCLoc(mlir::Type, llvm::ArrayRef<fir::ExtendedValue>);
mlir::Value genClock64(mlir::Type, llvm::ArrayRef<mlir::Value>);
template <mlir::arith::CmpIPredicate pred>
fir::ExtendedValue genCPtrCompare(mlir::Type,
llvm::ArrayRef<fir::ExtendedValue>);

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

@ -386,7 +386,10 @@ static constexpr IntrinsicHandler handlers[]{
{{{"name", asAddr}, {"status", asAddr, handleDynamicOptional}}},
/*isElemental=*/false},
{"clock", &I::genNVVMTime<mlir::NVVM::ClockOp>, {}, /*isElemental=*/false},
{"clock64", &I::genClock64, {}, /*isElemental=*/false},
{"clock64",
&I::genNVVMTime<mlir::NVVM::Clock64Op>,
{},
/*isElemental=*/false},
{"cmplx",
&I::genCmplx,
{{{"x", asValue}, {"y", asValue, handleDynamicOptional}}}},
@ -3565,16 +3568,6 @@ IntrinsicLibrary::genChdir(std::optional<mlir::Type> resultType,
return {};
}
// CLOCK64
mlir::Value IntrinsicLibrary::genClock64(mlir::Type resultType,
llvm::ArrayRef<mlir::Value> args) {
constexpr llvm::StringLiteral funcName = "llvm.nvvm.read.ptx.sreg.clock64";
mlir::MLIRContext *context = builder.getContext();
mlir::FunctionType ftype = mlir::FunctionType::get(context, {}, {resultType});
auto funcOp = builder.createFunction(loc, funcName, ftype);
return builder.create<fir::CallOp>(loc, funcOp, args).getResult(0);
}
// CMPLX
mlir::Value IntrinsicLibrary::genCmplx(mlir::Type resultType,
llvm::ArrayRef<mlir::Value> args) {
@ -7204,7 +7197,7 @@ IntrinsicLibrary::genNull(mlir::Type, llvm::ArrayRef<fir::ExtendedValue> args) {
return fir::MutableBoxValue(boxStorage, mold->nonDeferredLenParams(), {});
}
// CLOCK, GLOBALTIMER
// CLOCK, CLOCK64, GLOBALTIMER
template <typename OpTy>
mlir::Value IntrinsicLibrary::genNVVMTime(mlir::Type resultType,
llvm::ArrayRef<mlir::Value> args) {

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

@ -48,7 +48,6 @@ attributes(global) subroutine devsub()
smalltime = clock()
time = clock64()
time = globalTimer()
end
@ -87,7 +86,7 @@ end
! CHECK: %{{.*}} = llvm.atomicrmw udec_wrap %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i32
! CHECK: %{{.*}} = nvvm.read.ptx.sreg.clock : i32
! CHECK: fir.call @llvm.nvvm.read.ptx.sreg.clock64()
! CHECK: %{{.*}} = nvvm.read.ptx.sreg.clock64 : i64
! CHECK: %{{.*}} = nvvm.read.ptx.sreg.globaltimer : i64
subroutine host1()