add support and tests for __nvvm_sin/cos_approx intrinsics

Author: bratpiorka

llvm/lib/Target/NVPTX/NVPTXInstrInfo.td

@@ -1223,11 +1223,24 @@ def SIN_APPROX_f32 :
   BasicFlagsNVPTXInst<(outs B32:$dst), (ins B32:$src), (ins FTZFlag:$ftz),
                       "sin.approx$ftz.f32",
                       [(set f32:$dst, (UnaryOpAllowsApproxFn<fsin> f32:$src))]>;
+
+// Patterns for NVVM sin intrinsics
+def : Pat<(f32 (int_nvvm_sin_approx_f f32:$a)),
+          (SIN_APPROX_f32 f32:$a, 0)>;
+def : Pat<(f32 (int_nvvm_sin_approx_ftz_f f32:$a)),
+          (SIN_APPROX_f32 f32:$a, 1)>;
+
 def COS_APPROX_f32 :
   BasicFlagsNVPTXInst<(outs B32:$dst), (ins B32:$src), (ins FTZFlag:$ftz),
                       "cos.approx$ftz.f32",
                       [(set f32:$dst, (UnaryOpAllowsApproxFn<fcos> f32:$src))]>;
 
+// Patterns for NVVM cos intrinsics
+def : Pat<(f32 (int_nvvm_cos_approx_f f32:$a)),
+          (COS_APPROX_f32 f32:$a, 0)>;
+def : Pat<(f32 (int_nvvm_cos_approx_ftz_f f32:$a)),
+          (COS_APPROX_f32 f32:$a, 1)>;
+
 // NOTE: tanh.approx doesn't support the FTZ flag for f16/f16x2
 def TANH_APPROX_f16 :
   BasicNVPTXInst<(outs B16:$dst), (ins B16:$src), "tanh.approx.f16",
@@ -1238,7 +1251,7 @@ def TANH_APPROX_f16x2 :
   BasicNVPTXInst<(outs B32:$dst), (ins B32:$src), "tanh.approx.f16x2",
                  [(set v2f16:$dst, (UnaryOpAllowsApproxFn<ftanh> v2f16:$src))]>,
                  Requires<[hasPTX<70>, hasSM<75>]>;
-
+          
 def TANH_APPROX_f32 :
   BasicFlagsNVPTXInst<(outs B32:$dst), (ins B32:$src), (ins FTZFlag:$ftz),
                       "tanh.approx$ftz.f32",

sycl/test-e2e/DeviceLib/built-ins/nvvm_sin_cos_intrinsics.cpp

@@ -0,0 +1,77 @@
+// RUN: %{build} -fsycl-device-code-split=per_kernel -o %t.out
+// RUN: %{run} %t.out
+
+// Tests NVVM intrinsics for sin and cos approximations
+
+#include <cassert>
+#include <cmath>
+#include <sycl/detail/core.hpp>
+
+// Forward declarations of NVVM intrinsics
+extern "C" {
+float __nvvm_sin_approx_f(float);
+float __nvvm_sin_approx_ftz_f(float);
+float __nvvm_cos_approx_f(float);
+float __nvvm_cos_approx_ftz_f(float);
+}
+
+constexpr float TOLERANCE = 0.01f; // 1% tolerance for approximations
+
+template <typename Func>
+void test_approx(sycl::queue &q, Func intrinsic_func, const char *name,
+                 float input, float expected) {
+  float result = 0.0f;
+
+  {
+    sycl::buffer<float, 1> buf_result(&result, sycl::range<1>(1));
+    q.submit([&](sycl::handler &cgh) {
+       auto acc_result =
+           buf_result.template get_access<sycl::access::mode::write>(cgh);
+       cgh.single_task([=]() { acc_result[0] = intrinsic_func(input); });
+     }).wait();
+  }
+
+  float error = std::abs(result - expected);
+  assert(error < TOLERANCE && name && " approximation out of tolerance");
+}
+
+int main() {
+  sycl::queue q;
+
+  // Test values
+  const float pi = 3.14159265f;
+  const float test_values[] = {0.0f,      pi / 6.0f, pi / 4.0f, pi / 3.0f,
+                               pi / 2.0f, pi,        2.0f * pi};
+
+  // Expected sin values
+  const float expected_sin[] = {0.0f, 0.5f, 0.707107f, 0.866025f,
+                                1.0f, 0.0f, 0.0f};
+
+  // Expected cos values
+  const float expected_cos[] = {1.0f, 0.866025f, 0.707107f, 0.5f,
+                                0.0f, -1.0f,     1.0f};
+
+  // Test __nvvm_sin_approx_f
+  for (size_t i = 0; i < sizeof(test_values) / sizeof(test_values[0]); ++i) {
+    test_approx(q, __nvvm_sin_approx_f, "sin", test_values[i], expected_sin[i]);
+  }
+
+  // Test __nvvm_sin_approx_ftz_f
+  for (size_t i = 0; i < sizeof(test_values) / sizeof(test_values[0]); ++i) {
+    test_approx(q, __nvvm_sin_approx_ftz_f, "sin_ftz", test_values[i],
+                expected_sin[i]);
+  }
+
+  // Test __nvvm_cos_approx_f
+  for (size_t i = 0; i < sizeof(test_values) / sizeof(test_values[0]); ++i) {
+    test_approx(q, __nvvm_cos_approx_f, "cos", test_values[i], expected_cos[i]);
+  }
+
+  // Test __nvvm_cos_approx_ftz_f
+  for (size_t i = 0; i < sizeof(test_values) / sizeof(test_values[0]); ++i) {
+    test_approx(q, __nvvm_cos_approx_ftz_f, "cos_ftz", test_values[i],
+                expected_cos[i]);
+  }
+
+  return 0;
+}