Implement some math built-ins.

acos, acosh, acospi, asin, asinh, asinpi, atan, atan2, atanh, atanpi,
atan2pi, cbrt, ceil, copysign, cos, cosh and cospi.

There is not yet tests for these function, except for copysign (the only
one not simply using a STL or Boost math function).

6 files changed, 211 insertions, 13 deletions

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 25cde42..9449564 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -17,6 +17,7 @@ OPTION(BUILD_SHARED_LIBS "Set to OFF to build static libraries" ON)
 
 Find_Package(LLVM REQUIRED)
 Find_Package(Clang REQUIRED)
+Find_Package(Boost REQUIRED)
 
 ADD_DEFINITIONS("-DCOAL_BUILD_DIR=\"${CMAKE_BINARY_DIR}\"")
 ADD_DEFINITIONS("-DCOAL_INSTALL_DIR=\"${CMAKE_INSTALL_PREFIX}\"")
diff --git a/src/core/cpu/builtins.cpp b/src/core/cpu/builtins.cpp
index 04f54a6..de99499 100644
--- a/src/core/cpu/builtins.cpp
+++ b/src/core/cpu/builtins.cpp
@@ -51,6 +51,7 @@
 #include <iostream>
 #include <cstring>
 #include <cmath>
+#include <boost/math/special_functions.hpp>
 
 #include <stdio.h>
 
@@ -401,6 +402,8 @@ static void read_imageuif(uint32_t *result, Image2D *image, float x, float y,
  * Built-in functions generated by src/runtime/builtins.py
  */
 
+#define REPL(x) for (unsigned int i=0; i<x; ++i)
+
 #include <runtime/builtins_impl.h>
 
 /*
diff --git a/src/runtime/CMakeLists.txt b/src/runtime/CMakeLists.txt
index f825825..046049b 100644
--- a/src/runtime/CMakeLists.txt
+++ b/src/runtime/CMakeLists.txt
@@ -17,7 +17,8 @@ add_custom_command(
             ${CMAKE_CURRENT_SOURCE_DIR}/stdlib.c
             -I${CMAKE_CURRENT_BINARY_DIR}
             -o ${CMAKE_CURRENT_BINARY_DIR}/stdlib.c.bc
-    DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/stdlib.c)
+    DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/stdlib.c
+            ${CMAKE_CURRENT_BINARY_DIR}/stdlib_impl.h)
 
 add_custom_command(
     OUTPUT stdlib.c.bc.embed.h
diff --git a/src/runtime/builtins.def b/src/runtime/builtins.def
index 39e705b..b01279b 100644
--- a/src/runtime/builtins.def
+++ b/src/runtime/builtins.def
@@ -4,15 +4,183 @@ def veci : int2 int3 int4 int8 int16
 def vec : $vecf $veci
 def gentype : float $vecf
 
-func $type fmin $gentype : x:$type y:$type
-    return (x < y ? x : y);
+// gentype acos(gentype)
+// REPL is defined in src/core/cpu/builtins.cpp
+native float acos float : x:float
+    return std::acos(x);
+end
+
+native $type acos $vecf : x:$type
+    REPL($vecdim)
+        result[i] = std::acos(x[i]);
+end
+
+// gentype acosh(gentype)
+native float acosh float : x:float
+    return boost::math::acosh(x);
+end
+
+native $type acosh $vecf : x:$type
+    REPL($vecdim)
+        result[i] = boost::math::acosh(x[i]);
+end
+
+// gentype acospi(gentype)
+func float acospi float : x:float
+    return acos(x) / M_PI;
+end
+
+native $type acospi $vecf : x:$type
+    REPL($vecdim)
+        result[i] = std::acos(x[i]) / M_PI;
+end
+
+// gentype asin (gentype)
+native float asin float : x:float
+    return std::asin(x);
+end
+
+native $type asin $vecf : x:$type
+    REPL($vecdim)
+        result[i] = std::asin(x[i]);
+end
+
+// gentype asinh (gentype)
+native float asinh float : x:float
+    return boost::math::asinh(x);
+end
+
+native $type asinh $vecf : x:$type
+    REPL($vecdim)
+        result[i] = boost::math::asinh(x[i]);
+end
+
+// gentype asinpi (gentype x)
+func float asinpi float : x:float
+    return asin(x) / M_PI;
+end
+
+native $type asinpi $vecf : x:$type
+    REPL($vecdim)
+        result[i] = std::asin(x[i]) / M_PI;
+end
+
+// gentype atan (gentype y_over_x)
+native float atan float : y_over_x:float
+    return std::atan(y_over_x);
+end
+
+native $type atan $vecf : y_over_x:$type
+    REPL($vecdim)
+        result[i] = std::atan(y_over_x[i]);
+end
+
+// gentype atan2 (gentype y, gentype x)
+func float atan2 float : x:float y:float
+    return atan(y / x);
+end
+
+native $type atan2 $vecf : x:$type y:$type
+    REPL($vecdim)
+        result[i] = std::atan(y[i] / x[i]);
+end
+
+// gentype atanh (gentype)
+native float atanh float : x:float
+    return boost::math::atanh(x);
 end
 
+native $type atanh $vecf : x:$type
+    REPL($vecdim)
+        result[i] = boost::math::atanh(x[i]);
+end
+
+// gentype atanpi (gentype x)
+func float atanpi float : x:float
+    return atan(x) / M_PI;
+end
+
+native $type atanpi $vecf : x:$type
+    REPL($vecdim)
+        result[i] = std::atan(x[i]) / M_PI;
+end
+
+// gentype atan2pi (gentype y, gentype x)
+func float atan2pi float : x:float y:float
+    return atan2(y, x) / M_PI;
+end
+
+native $type atan2pi $vecf : x:$type y:$type
+    REPL($vecdim)
+        result[i] = std::atan(y[i] / x[i]) / M_PI;
+end
+
+// gentype cbrt (gentype)
+native float cbrt float : x:float
+    return boost::math::cbrt(x);
+end
+
+native $type cbrt $vecf : x:$type
+    REPL($vecdim)
+        result[i] = boost::math::cbrt(x[i]);
+end
+
+// gentype ceil (gentype)
+native float ceil float : x:float
+    return std::ceil(x);
+end
+
+native $type ceil $vecf : x:$type
+    REPL($vecdim)
+        result[i] = std::ceil(x[i]);
+end
+
+// gentype copysign (gentype x, gentype y)
+func float copysign float : x:float y:float
+    if ((x < 0.0f && y > 0.0f) ||
+        (x > 0.0f && y < 0.0f))
+        return -x;
+end
+
+native $type copysign $vecf : x:$type y:$type
+    REPL($vecdim)
+    {
+        if ((x[i] < 0.0f && y[i] > 0.0f) ||
+            (x[i] > 0.0f && y[i] < 0.0f))
+            result[i] = -x[i];
+        else
+            result[i] = x[i];
+    }
+end
+
+//gentype cos (gentype)
 native float cos float : x:float
     return std::cos(x);
 end
 
 native $type cos $vecf : x:$type
-    for (unsigned int i=0; i<$vecdim; ++i)
+    REPL($vecdim)
         result[i] = std::cos(x[i]);
 end
+
+// gentype cosh (gentype)
+native float cosh float : x:float
+    return std::cosh(x);
+end
+
+native $type cosh $vecf : x:$type
+    REPL($vecdim)
+        result[i] = std::cosh(x[i]);
+end
+
+// gentype cospi (gentype x)
+func $type cospi $gentype : x:$type
+    return cos(x * M_PI);
+end
+
+// TODO: gentype erfc (gentype)
+// TODO: gentype erf (gentype)
+
+func $type fmin $gentype : x:$type y:$type
+    return (x < y ? x : y);
+end
diff --git a/src/runtime/stdlib.h b/src/runtime/stdlib.h
index b7208fe..2cfe88e 100644
--- a/src/runtime/stdlib.h
+++ b/src/runtime/stdlib.h
@@ -66,19 +66,19 @@ typedef struct image3d *image3d_t;
    COAL_VECTOR(type, 8);      \
    COAL_VECTOR(type, 16);
 
-COAL_VECTOR_SET(char);
-COAL_VECTOR_SET(uchar);
+COAL_VECTOR_SET(char)
+COAL_VECTOR_SET(uchar)
 
-COAL_VECTOR_SET(short);
-COAL_VECTOR_SET(ushort);
+COAL_VECTOR_SET(short)
+COAL_VECTOR_SET(ushort)
 
-COAL_VECTOR_SET(int);
-COAL_VECTOR_SET(uint);
+COAL_VECTOR_SET(int)
+COAL_VECTOR_SET(uint)
 
-COAL_VECTOR_SET(long);
-COAL_VECTOR_SET(ulong);
+COAL_VECTOR_SET(long)
+COAL_VECTOR_SET(ulong)
 
-COAL_VECTOR_SET(float);
+COAL_VECTOR_SET(float)
 
 #undef COAL_VECTOR_SET
 #undef COAL_VECTOR
@@ -146,6 +146,20 @@ COAL_VECTOR_SET(float);
 #define CLK_HALF_FLOAT               0x10DD
 #define CLK_FLOAT                    0x10DE
 
+#define M_E            2.7182818284590452354   /* e */
+#define M_LOG2E        1.4426950408889634074   /* log_2 e */
+#define M_LOG10E       0.43429448190325182765  /* log_10 e */
+#define M_LN2          0.69314718055994530942  /* log_e 2 */
+#define M_LN10         2.30258509299404568402  /* log_e 10 */
+#define M_PI           3.14159265358979323846  /* pi */
+#define M_PI_2         1.57079632679489661923  /* pi/2 */
+#define M_PI_4         0.78539816339744830962  /* pi/4 */
+#define M_1_PI         0.31830988618379067154  /* 1/pi */
+#define M_2_PI         0.63661977236758134308  /* 2/pi */
+#define M_2_SQRTPI     1.12837916709551257390  /* 2/sqrt(pi) */
+#define M_SQRT2        1.41421356237309504880  /* sqrt(2) */
+#define M_SQRT1_2      0.70710678118654752440  /* 1/sqrt(2) */
+
 /* Typedefs */
 typedef unsigned int cl_mem_fence_flags;
 
diff --git a/tests/test_builtins.cpp b/tests/test_builtins.cpp
index a687b9a..80852a9 100644
--- a/tests/test_builtins.cpp
+++ b/tests/test_builtins.cpp
@@ -98,12 +98,17 @@ const char image_source[] =
 const char builtins_source[] =
     "__kernel void test_case(__global uint *rs) {\n"
     "   float2 f2;\n"
+    "   float2 f2b;\n"
     "\n"
     "   f2.x = 1.0f;\n"
     "   f2.y = 0.0f;\n"
+    "   f2b.x = -0.5f;\n"
+    "   f2b.y = (float)M_PI;\n"
     "\n"
     "   if (cos(f2).y != 1.0f) { *rs = 1; return; }\n"
     "   if (cos(0.0f) != 1.0f) { *rs = 2; return; }\n"
+    "   if (copysign(1.0f, -0.5f) != -1.0f) { *rs = 3; return; }\n"
+    "   if (copysign(f2, f2b).x != -1.0f) { *rs = 4; return; }\n"
     "}\n";
 
 enum TestCaseKind
@@ -381,6 +386,12 @@ START_TEST (test_builtins)
         case 2:
             errstr = "float cos(float) doesn't behave correctly";
             break;
+        case 3:
+            errstr = "float copysign(float) doesn't behave correctly";
+            break;
+        case 4:
+            errstr = "float2 copysign(float2) doesn't behave correctly";
+            break;
         default:
             errstr = default_error(rs);
     }