[Mlir-commits] [mlir] 9df51a9 - [MLIR][NVVM] Add `nvvm.sin` OP (#193775)

[llvmlistbot at llvm.org]

Author: Guray Ozen
Date: 2026-04-24T12:34:36+02:00
New Revision: 9df51a964ddbe4094f7e5a5b4a43f2776dd0da47

URL: https://github.com/llvm/llvm-project/commit/9df51a964ddbe4094f7e5a5b4a43f2776dd0da47
DIFF: https://github.com/llvm/llvm-project/commit/9df51a964ddbe4094f7e5a5b4a43f2776dd0da47.diff

LOG: [MLIR][NVVM] Add `nvvm.sin` OP (#193775)

Implement `nvvm.sin` with ftz flag

Added: 
    mlir/test/Dialect/LLVMIR/nvvm-transcendentals.mlir
    mlir/test/Target/LLVMIR/nvvm/transcendentals.mlir

Modified: 
    mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
    mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp

Removed: 
    


################################################################################
diff  --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index c892ee18166f2..9e94477c3a60a 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -518,6 +518,31 @@ 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_SingleResultIntrinsicOp<"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)";
+}
+
 //===----------------------------------------------------------------------===//
 // NVVM redux op definitions
 //===----------------------------------------------------------------------===//

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
new file mode 100644
index 0000000000000..c773e3f06a612
--- /dev/null
+++ b/mlir/test/Dialect/LLVMIR/nvvm-transcendentals.mlir
@@ -0,0 +1,15 @@
+// RUN: mlir-opt %s -split-input-file | 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
+}