Device code test for cuda native fcts.

Signed-off-by: JackAKirk <[email protected]>

Author: JackAKirk

sycl/test/check_device_code/math-builtins/native-math-cuda.cpp

@@ -0,0 +1,69 @@
+// REQUIRES: cuda
+
+// RUN: %clangxx -fsycl -fsycl-targets=nvptx64-nvidia-cuda -S -Xclang -emit-llvm %s -o -| FileCheck %s
+
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+
+int main() {
+
+  queue q;
+
+  float input[2];
+  float res[13];
+  {
+    buffer<float, 1> input_buff(&input[0], range<1>(2));
+    buffer<float, 1> res_buff(&res[0], range<1>(13));
+    q.submit([&](handler &cgh) {
+      accessor<float, 1, access::mode::write, target::device> res_acc(res_buff,
+                                                                      cgh);
+      accessor<float, 1, access::mode::read, target::device> input_acc(
+          input_buff, cgh);
+      cgh.single_task([=]() {
+        // CHECK: tail call float @llvm.nvvm.cos.approx.f
+        res_acc[0] = sycl::native::cos(input_acc[0]);
+        // CHECK: tail call float @llvm.nvvm.sin.approx.f
+        res_acc[1] = sycl::native::sin(input_acc[0]);
+        // CHECK: tail call float @llvm.nvvm.ex2.approx.f
+        res_acc[2] = sycl::native::exp2(input_acc[0]);
+        // CHECK: tail call float @llvm.nvvm.lg2.approx.f
+        res_acc[3] = sycl::native::log2(input_acc[0]);
+        // CHECK: tail call float @llvm.nvvm.rsqrt.approx.f
+        res_acc[4] = sycl::native::rsqrt(input_acc[0]);
+        // CHECK: tail call float @llvm.nvvm.sqrt.approx.f
+        res_acc[5] = sycl::native::sqrt(input_acc[0]);
+        // CHECK: tail call float @llvm.nvvm.rcp.approx.f
+        res_acc[6] = sycl::native::recip(input_acc[0]);
+        // CHECK: tail call float @llvm.nvvm.div.approx.f
+        res_acc[7] = sycl::native::divide(input_acc[0], input_acc[1]);
+
+        // Functions that use the above builtins:
+
+        // CHECK: tail call float @llvm.nvvm.sin.approx.f
+        // CHECK: tail call float @llvm.nvvm.cos.approx.f
+        // CHECK: tail call float @llvm.nvvm.div.approx.f
+        res_acc[8] = sycl::native::tan(input_acc[0]);
+        // CHECK: fmul float {{.*}}, 0x3FF7154760000000
+        // CHECK: tail call float @llvm.nvvm.ex2.approx.f
+        res_acc[9] = sycl::native::exp(input_acc[0]);
+        // CHECK: fmul float {{.*}}, 0x400A934F00000000
+        // CHECK: tail call float @llvm.nvvm.ex2.approx.f
+        res_acc[10] = sycl::native::exp10(input_acc[0]);
+        // CHECK: tail call float @llvm.nvvm.lg2.approx.f
+        // CHECK: fmul float {{.*}}, 0x3FE62E4300000000
+        res_acc[11] = sycl::native::log(input_acc[0]);
+        // CHECK: tail call float @llvm.nvvm.lg2.approx.f
+        // CHECK: fmul float {{.*}}, 0x3FD3441360000000
+        res_acc[12] = sycl::native::log10(input_acc[0]);
+
+        // CHECK: tail call float @llvm.nvvm.lg2.approx.f
+        // CHECK: fmul float {{.*}}, {{.*}}
+        // CHECK: tail call float @llvm.nvvm.ex2.approx.f
+        res_acc[13] = sycl::native::powr(input_acc[0], input_acc[1]);
+      });
+    });
+  }
+
+  return 0;
+};