From c377dd4e18c7321db6f12cc4187ec1a2fbc7cb47 Mon Sep 17 00:00:00 2001 From: Denis Steckelmacher Date: Wed, 19 Oct 2011 12:31:50 +0200 Subject: 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). --- CMakeLists.txt | 1 + src/core/cpu/builtins.cpp | 3 + src/runtime/CMakeLists.txt | 3 +- src/runtime/builtins.def | 174 ++++++++++++++++++++++++++++++++++++++++++++- src/runtime/stdlib.h | 32 ++++++--- tests/test_builtins.cpp | 11 +++ 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 #include #include +#include #include @@ -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 /* 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); } -- cgit v1.2.3