summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--CMakeLists.txt1
-rw-r--r--src/core/cpu/builtins.cpp3
-rw-r--r--src/runtime/CMakeLists.txt3
-rw-r--r--src/runtime/builtins.def174
-rw-r--r--src/runtime/stdlib.h32
-rw-r--r--tests/test_builtins.cpp11
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);
}
generated by cgit v1.2.3 (git 2.39.1) at 2024-11-23 10:14:37 +0000