[Mlir-commits] [mlir] [MLIR][NVVM] Add `nvvm.sin` OP (PR #193775)
Guray Ozen
llvmlistbot at llvm.org
Fri Apr 24 02:48:18 PDT 2026
https://github.com/grypp updated https://github.com/llvm/llvm-project/pull/193775
>From 0a3bb1e9acd2a79e855f4da0dfad52410b64d74a Mon Sep 17 00:00:00 2001
From: Guray Ozen <gozen at nvidia.com>
Date: Thu, 23 Apr 2026 17:22:13 +0200
Subject: [PATCH 1/2] [MLIR][NVVM] Add `nvvm.sin` OP
Implement `nvvm.sin` with ftz flag
---
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 29 +++++++++++++++++++
.../Dialect/LLVMIR/nvvm-transcendentals.mlir | 15 ++++++++++
.../Target/LLVMIR/nvvm/transcendentals.mlir | 15 ++++++++++
3 files changed, 59 insertions(+)
create mode 100644 mlir/test/Dialect/LLVMIR/nvvm-transcendentals.mlir
create mode 100644 mlir/test/Target/LLVMIR/nvvm/transcendentals.mlir
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index c892ee18166f2..cf6523a0fc4fb 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -518,6 +518,35 @@ def NVVM_RcpApproxFtzF32Op : NVVM_IntrOp<"rcp.approx.ftz.f", [Pure], 1> {
let assemblyFormat = "$arg attr-dict `:` type($res)";
}
+//===----------------------------------------------------------------------===//
+// NVVM transcendental op definitions
+//===----------------------------------------------------------------------===//
+// PTX provides only fast approximations for sin/cos/lg2/ex2/tanh. Exposing
+// these as plain `nvvm.{sin,cos,lg2,ex2,tanh}` (without a `.approx` suffix)
+// keeps the NVVM dialect self-contained -- users don't need to reach for
+// `math.*` ops for something that already has a PTX instruction.
+
+def NVVM_SinOp : NVVM_Op<"sin", [Pure, SameOperandsAndResultType]> {
+ let summary = "Sine (fast approximation)";
+ let description = [{
+ Computes a fast approximation of the sine of the input value (in radians).
+ Lowers to PTX `sin.approx{.ftz}.f32`. The `ftz` attribute, when set,
+ flushes subnormal inputs and results to sign-preserving zero.
+
+ For more information, see PTX ISA:
+ [sin](https://docs.nvidia.com/cuda/parallel-thread-execution/#floating-point-instructions-sin)
+ }];
+ let arguments = (ins F32:$src,
+ DefaultValuedAttr<BoolAttr, "false">:$ftz);
+ let results = (outs F32:$res);
+ let assemblyFormat = "$src attr-dict `:` type($src)";
+ string llvmBuilder = [{
+ unsigned IID = $ftz ? llvm::Intrinsic::nvvm_sin_approx_ftz_f
+ : llvm::Intrinsic::nvvm_sin_approx_f;
+ $res = createIntrinsicCall(builder, IID, {$src});
+ }];
+}
+
//===----------------------------------------------------------------------===//
// NVVM redux op definitions
//===----------------------------------------------------------------------===//
diff --git a/mlir/test/Dialect/LLVMIR/nvvm-transcendentals.mlir b/mlir/test/Dialect/LLVMIR/nvvm-transcendentals.mlir
new file mode 100644
index 0000000000000..0a58c1892bded
--- /dev/null
+++ b/mlir/test/Dialect/LLVMIR/nvvm-transcendentals.mlir
@@ -0,0 +1,15 @@
+// RUN: mlir-opt %s -split-input-file -verify-diagnostics | FileCheck %s
+
+// CHECK-LABEL: @nvvm_sin_f32
+func.func @nvvm_sin_f32(%arg0: f32) -> f32 {
+ // CHECK: nvvm.sin {{.*}} : f32
+ %0 = nvvm.sin %arg0 : f32
+ return %0 : f32
+}
+
+// CHECK-LABEL: @nvvm_sin_ftz_f32
+func.func @nvvm_sin_ftz_f32(%arg0: f32) -> f32 {
+ // CHECK: nvvm.sin {{.*}} {ftz = true} : f32
+ %0 = nvvm.sin %arg0 {ftz = true} : f32
+ return %0 : f32
+}
diff --git a/mlir/test/Target/LLVMIR/nvvm/transcendentals.mlir b/mlir/test/Target/LLVMIR/nvvm/transcendentals.mlir
new file mode 100644
index 0000000000000..65374c9bd5e87
--- /dev/null
+++ b/mlir/test/Target/LLVMIR/nvvm/transcendentals.mlir
@@ -0,0 +1,15 @@
+// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
+
+// CHECK-LABEL: @nvvm_sin
+llvm.func @nvvm_sin(%arg0: f32) -> f32 {
+ // CHECK: call float @llvm.nvvm.sin.approx.f(float %{{.*}})
+ %0 = nvvm.sin %arg0 : f32
+ llvm.return %0 : f32
+}
+
+// CHECK-LABEL: @nvvm_sin_ftz
+llvm.func @nvvm_sin_ftz(%arg0: f32) -> f32 {
+ // CHECK: call float @llvm.nvvm.sin.approx.ftz.f(float %{{.*}})
+ %0 = nvvm.sin %arg0 {ftz = true} : f32
+ llvm.return %0 : f32
+}
>From 681c00e775a9c8eca521fcc61fe2ef2dce570d06 Mon Sep 17 00:00:00 2001
From: Guray Ozen <gozen at nvidia.com>
Date: Fri, 24 Apr 2026 11:48:05 +0200
Subject: [PATCH 2/2] [MLIR][NVVM] Address review comments on nvvm.sin
- Switch NVVM_SinOp to the NVVM_SingleResultIntrinsicOp wrapper; move
intrinsic-ID selection into getIntrinsicIDAndArgs.
- Drop -verify-diagnostics from the dialect test RUN line (no
expected-error directives in the file).
---
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 8 ++------
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp | 10 ++++++++++
mlir/test/Dialect/LLVMIR/nvvm-transcendentals.mlir | 2 +-
3 files changed, 13 insertions(+), 7 deletions(-)
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index cf6523a0fc4fb..9e94477c3a60a 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -526,7 +526,8 @@ def NVVM_RcpApproxFtzF32Op : NVVM_IntrOp<"rcp.approx.ftz.f", [Pure], 1> {
// keeps the NVVM dialect self-contained -- users don't need to reach for
// `math.*` ops for something that already has a PTX instruction.
-def NVVM_SinOp : NVVM_Op<"sin", [Pure, SameOperandsAndResultType]> {
+def NVVM_SinOp : NVVM_SingleResultIntrinsicOp<"sin",
+ [Pure, SameOperandsAndResultType]> {
let summary = "Sine (fast approximation)";
let description = [{
Computes a fast approximation of the sine of the input value (in radians).
@@ -540,11 +541,6 @@ def NVVM_SinOp : NVVM_Op<"sin", [Pure, SameOperandsAndResultType]> {
DefaultValuedAttr<BoolAttr, "false">:$ftz);
let results = (outs F32:$res);
let assemblyFormat = "$src attr-dict `:` type($src)";
- string llvmBuilder = [{
- unsigned IID = $ftz ? llvm::Intrinsic::nvvm_sin_approx_ftz_f
- : llvm::Intrinsic::nvvm_sin_approx_f;
- $res = createIntrinsicCall(builder, IID, {$src});
- }];
}
//===----------------------------------------------------------------------===//
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index 31e7ff209db5c..b0cebd45624a9 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -3464,6 +3464,16 @@ mlir::NVVM::IDArgPair NVVM::BarrierOp::getIntrinsicIDAndArgs(
return {id, std::move(args)};
}
+mlir::NVVM::IDArgPair
+SinOp::getIntrinsicIDAndArgs(Operation &op, LLVM::ModuleTranslation &mt,
+ llvm::IRBuilderBase &builder) {
+ auto thisOp = cast<NVVM::SinOp>(op);
+ llvm::Intrinsic::ID id = thisOp.getFtz()
+ ? llvm::Intrinsic::nvvm_sin_approx_ftz_f
+ : llvm::Intrinsic::nvvm_sin_approx_f;
+ return {id, {mt.lookupValue(thisOp.getSrc())}};
+}
+
mlir::NVVM::IDArgPair
PMEventOp::getIntrinsicIDAndArgs(Operation &op, LLVM::ModuleTranslation &mt,
llvm::IRBuilderBase &builder) {
diff --git a/mlir/test/Dialect/LLVMIR/nvvm-transcendentals.mlir b/mlir/test/Dialect/LLVMIR/nvvm-transcendentals.mlir
index 0a58c1892bded..c773e3f06a612 100644
--- a/mlir/test/Dialect/LLVMIR/nvvm-transcendentals.mlir
+++ b/mlir/test/Dialect/LLVMIR/nvvm-transcendentals.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s -split-input-file -verify-diagnostics | FileCheck %s
+// RUN: mlir-opt %s -split-input-file | FileCheck %s
// CHECK-LABEL: @nvvm_sin_f32
func.func @nvvm_sin_f32(%arg0: f32) -> f32 {
More information about the Mlir-commits
mailing list