Fixes for Mac OS X

example/black_scholes.cpp

@@ -145,8 +145,18 @@ int main()
     std::cout << "option 0 call price: " << call0 << std::endl;
     std::cout << "option 0 put price: " << put0 << std::endl;
 
+    // due to the differences in the random-number generators between linux and
+    // mac os x, we will get different "expected" results for this example
+#ifdef __APPLE__
+    double expected_call0 = 0.000249461;
+    double expected_put0 = 26.2798;
+#else
+    double expected_call0 = 0.0999f;
+    double expected_put0 = 43.0524f;
+#endif
+
     // check option prices
-    if(std::abs(call0 - 0.0999f) > 1e-4 || std::abs(put0 - 43.0524f) > 1e-4){
+    if(std::abs(call0 - expected_call0) > 1e-4 || std::abs(put0 - expected_put0) > 1e-4){
         std::cerr << "error: option prices are wrong" << std::endl;
         return -1;
     }

example/simple_moving_average.cpp

@@ -32,7 +32,7 @@ compute::program make_sma_program(const compute::context& context)
         {
             const int gid = get_global_id(0);
 
-            float cumValues = 0.;
+            float cumValues = 0.f;
             int endIdx = gid + wSize/2;
             int startIdx = gid -1 - wSize/2;
 

include/boost/compute/algorithm/lexicographical_compare.hpp

@@ -21,15 +21,14 @@ namespace compute {
 namespace detail {
 
 const char lexicographical_compare_source[] =
-"__kernel void lexicographical_compare(uint size1,\n"
-"                                      uint size2,\n"
+"__kernel void lexicographical_compare(const uint size1,\n"
+"                                      const uint size2,\n"
 "                                      __global const T1 *range1,\n"
 "                                      __global const T2 *range2,\n"
 "                                      __global bool *result_buf)\n"
 "{\n"
-"   const int i = get_global_id(0);\n"
-"   if((i != size1) && (i != size2))\n"
-"   {\n"
+"   const uint i = get_global_id(0);\n"
+"   if((i != size1) && (i != size2)){\n"
         //Individual elements are compared and results are stored in parallel.
         //0 is true
 "       if(range1[i] < range2[i])\n"

include/boost/compute/container/dynamic_bitset.hpp

@@ -160,10 +160,10 @@ class dynamic_bitset
 
         // update block value
         if(value){
-            block_value |= (1 << bit);
+            block_value |= (size_type(1) << bit);
         }
         else {
-            block_value &= ~(1 << bit);
+            block_value &= ~(size_type(1) << bit);
         }
 
         // store new block
@@ -179,7 +179,7 @@ class dynamic_bitset
         block_type block_value;
         copy_n(m_bits.begin() + block, 1, &block_value, queue);
 
-        return block_value & (1 << bit);
+        return block_value & (size_type(1) << bit);
     }
 
     /// Flips the value of the bit at position \p n.

include/boost/compute/context.hpp

@@ -251,6 +251,13 @@ class context
                           size_t cb,
                           void *user_data)
     {
+    #ifdef __APPLE__
+        // on apple, every single opencl failure is reported through the
+        // context error handler. in order to let failures propogate
+        // via opencl_error, we don't throw context_errors from here.
+        return;
+    #endif
+
         context *this_ = static_cast<context *>(user_data);
 
         BOOST_THROW_EXCEPTION(

include/boost/compute/detail/literal.hpp

@@ -0,0 +1,45 @@
+//---------------------------------------------------------------------------//
+// Copyright (c) 2013-2014 Kyle Lutz <kyle.r.lutz@gmail.com>
+//
+// Distributed under the Boost Software License, Version 1.0
+// See accompanying file LICENSE_1_0.txt or copy at
+// http://www.boost.org/LICENSE_1_0.txt
+//
+// See http://kylelutz.github.com/compute for more information.
+//---------------------------------------------------------------------------//
+
+#ifndef BOOST_COMPUTE_DETAIL_LITERAL_HPP
+#define BOOST_COMPUTE_DETAIL_LITERAL_HPP
+
+#include <iomanip>
+#include <limits>
+#include <sstream>
+
+#include <boost/type_traits/is_same.hpp>
+
+#include <boost/compute/types/fundamental.hpp>
+
+namespace boost {
+namespace compute {
+namespace detail {
+
+template<class T>
+std::string make_literal(T x)
+{
+    std::stringstream s;
+    s << std::setprecision(std::numeric_limits<T>::digits10)
+      << std::scientific
+      << x;
+
+    if(boost::is_same<T, float>::value || boost::is_same<T, float_>::value){
+        s << "f";
+    }
+
+    return s.str();
+}
+
+} // end detail namespace
+} // end compute namespace
+} // end boost namespace
+
+#endif // BOOST_COMPUTE_DETAIL_LITERAL_HPP

include/boost/compute/random/bernoulli_distribution.hpp

@@ -15,6 +15,7 @@
 #include <boost/compute/function.hpp>
 #include <boost/compute/types/fundamental.hpp>
 #include <boost/compute/detail/iterator_range_size.hpp>
+#include <boost/compute/detail/literal.hpp>
 
 namespace boost {
 namespace compute {
@@ -70,7 +71,7 @@ class bernoulli_distribution
             return (convert_RealType(x) / MAX_RANDOM) < PARAM;
         });
 
-        scale_random.define("PARAM", boost::lexical_cast<std::string>(m_p));
+        scale_random.define("PARAM", detail::make_literal(m_p));
         scale_random.define("MAX_RANDOM", "UINT_MAX");
         scale_random.define(
             "convert_RealType", std::string("convert_") + type_name<RealType>()

include/boost/compute/random/discrete_distribution.hpp

@@ -13,10 +13,11 @@
 
 #include <boost/compute/command_queue.hpp>
 #include <boost/compute/function.hpp>
-#include <boost/compute/types/fundamental.hpp>
 #include <boost/compute/algorithm/accumulate.hpp>
 #include <boost/compute/algorithm/copy.hpp>
 #include <boost/compute/algorithm/transform.hpp>
+#include <boost/compute/detail/literal.hpp>
+#include <boost/compute/types/fundamental.hpp>
 
 namespace boost {
 namespace compute {
@@ -90,15 +91,12 @@ class discrete_distribution
         for(size_t i=0; i<m_n; i++)
         {
             source = source +
-                "if(rno <= " +
-                    boost::lexical_cast<std::string>(m_probabilities[i]) +
-                    ")\n" +
-                "   return " + boost::lexical_cast<std::string>(i) +
-                        ";\n";
+                "if(rno <= " + detail::make_literal<float>(m_probabilities[i]) + ")\n" +
+                "   return " + detail::make_literal(i) + ";\n";
         }
 
         source = source +
-            "return " + boost::lexical_cast<std::string>(m_n-1) + ";\n" +
+            "return " + detail::make_literal(m_n - 1) + ";\n" +
             "}\n";
 
         BOOST_COMPUTE_FUNCTION(IntType, scale_random, (const uint_ x), {});

include/boost/compute/random/uniform_real_distribution.hpp

@@ -13,6 +13,7 @@
 
 #include <boost/compute/command_queue.hpp>
 #include <boost/compute/function.hpp>
+#include <boost/compute/detail/literal.hpp>
 #include <boost/compute/types/fundamental.hpp>
 
 namespace boost {
@@ -71,8 +72,8 @@ class uniform_real_distribution
             return LO + (convert_RealType(x) / MAX_RANDOM) * (HI - LO);
         });
 
-        scale_random.define("LO", boost::lexical_cast<std::string>(m_a));
-        scale_random.define("HI", boost::lexical_cast<std::string>(m_b));
+        scale_random.define("LO", detail::make_literal(m_a));
+        scale_random.define("HI", detail::make_literal(m_b));
         scale_random.define("MAX_RANDOM", "UINT_MAX");
         scale_random.define(
             "convert_RealType", std::string("convert_") + type_name<RealType>()

test/CMakeLists.txt

@@ -2,6 +2,11 @@ include_directories(../include)
 
 set(BOOST_COMPONENTS unit_test_framework)
 
+if(${BOOST_COMPUTE_USE_CPP11})
+  # allow tests to use C++11 features
+  add_definitions(-DBOOST_COMPUTE_USE_CPP11)
+endif()
+
 if (${BOOST_COMPUTE_USE_OFFLINE_CACHE})
   set(BOOST_COMPONENTS ${BOOST_COMPONENTS} system filesystem)
   add_definitions(-DBOOST_COMPUTE_USE_OFFLINE_CACHE)

test/test_device.cpp

@@ -100,9 +100,6 @@ BOOST_AUTO_TEST_CASE(partition_device_equally)
         return;
     }
 
-    // ensure device is not a sub-device
-    BOOST_CHECK(device.is_subdevice() == false);
-
     // check that the device supports partitioning equally
     if(!supports_partition_type(device, CL_DEVICE_PARTITION_EQUALLY)){
         std::cout << "skipping test: "
@@ -111,6 +108,9 @@ BOOST_AUTO_TEST_CASE(partition_device_equally)
         return;
     }
 
+    // ensure device is not a sub-device
+    BOOST_CHECK(device.is_subdevice() == false);
+
     // partition default device into sub-devices with two compute units each
     std::vector<boost::compute::device>
         sub_devices = device.partition_equally(2);
@@ -199,9 +199,6 @@ BOOST_AUTO_TEST_CASE(partition_by_affinity_domain)
         return;
     }
 
-    // ensure device is not a sub-device
-    BOOST_CHECK(device.is_subdevice() == false);
-
     // check that the device supports splitting by affinity domains
     if(!supports_partition_type(device, CL_DEVICE_AFFINITY_DOMAIN_NEXT_PARTITIONABLE)){
         std::cout << "skipping test: "
@@ -210,6 +207,9 @@ BOOST_AUTO_TEST_CASE(partition_by_affinity_domain)
         return;
     }
 
+    // ensure device is not a sub-device
+    BOOST_CHECK(device.is_subdevice() == false);
+
     std::vector<boost::compute::device> sub_devices =
         device.partition_by_affinity_domain(
             CL_DEVICE_AFFINITY_DOMAIN_NEXT_PARTITIONABLE);

test/test_event.cpp

@@ -11,11 +11,12 @@
 #define BOOST_TEST_MODULE TestEvent
 #include <boost/test/unit_test.hpp>
 
-#include <boost/config.hpp>
+#include <vector>
 
-#if !defined(BOOST_NO_CXX11_HDR_FUTURE) && !defined(BOOST_NO_0X_HDR_FUTURE)
+#ifdef BOOST_COMPUTE_USE_CPP11
+#include <mutex>
 #include <future>
-#endif // BOOST_NO_CXX11_HDR_FUTURE
+#endif // BOOST_COMPUTE_USE_CPP11
 
 #include <boost/compute/event.hpp>
 
@@ -27,43 +28,63 @@ BOOST_AUTO_TEST_CASE(null_event)
     BOOST_CHECK(null.get() == cl_event());
 }
 
-#ifdef CL_VERSION_1_1
+#if defined(CL_VERSION_1_1) && defined(BOOST_COMPUTE_USE_CPP11)
+std::mutex callback_mutex;
+std::condition_variable callback_condition_variable;
 static bool callback_invoked = false;
 
 static void BOOST_COMPUTE_CL_CALLBACK
 callback(cl_event event, cl_int status, void *user_data)
 {
+    std::lock_guard<std::mutex> lock(callback_mutex);
     callback_invoked = true;
+    callback_condition_variable.notify_one();
 }
 
 BOOST_AUTO_TEST_CASE(event_callback)
 {
     REQUIRES_OPENCL_VERSION(1,2);
 
+    // ensure callback has not yet been executed
     BOOST_CHECK_EQUAL(callback_invoked, false);
-    {
-        boost::compute::event marker = queue.enqueue_marker();
-        marker.set_callback(callback);
-        queue.finish();
-    }
+
+    // enqueue marker and set callback to be invoked
+    boost::compute::event marker = queue.enqueue_marker();
+    marker.set_callback(callback);
+    marker.wait();
+
+    // wait up to one second for the callback to be executed
+    std::unique_lock<std::mutex> lock(callback_mutex);
+    callback_condition_variable.wait_for(
+        lock, std::chrono::seconds(1), [&](){ return callback_invoked; }
+    );
+
+    // ensure callback has been executed
     BOOST_CHECK_EQUAL(callback_invoked, true);
 }
 
-#if !defined(BOOST_NO_CXX11_LAMBDAS) && !defined(BOOST_NO_LAMBDAS)
 BOOST_AUTO_TEST_CASE(lambda_callback)
 {
     REQUIRES_OPENCL_VERSION(1,2);
+
     bool lambda_invoked = false;
-    {
-        boost::compute::event marker = queue.enqueue_marker();
-        marker.set_callback([&lambda_invoked](){ lambda_invoked = true; });
-        queue.finish();
-    }
+
+    boost::compute::event marker = queue.enqueue_marker();
+    marker.set_callback([&](){
+        std::lock_guard<std::mutex> lock(callback_mutex);
+        lambda_invoked = true;
+        callback_condition_variable.notify_one();
+    });
+    marker.wait();
+
+    // wait up to one second for the callback to be executed
+    std::unique_lock<std::mutex> lock(callback_mutex);
+    callback_condition_variable.wait_for(
+        lock, std::chrono::seconds(1), [&](){ return lambda_invoked; }
+    );
     BOOST_CHECK_EQUAL(lambda_invoked, true);
 }
-#endif // BOOST_NO_CXX11_LAMBDAS
 
-#if !defined(BOOST_NO_CXX11_HDR_FUTURE) && !defined(BOOST_NO_0X_HDR_FUTURE)
 void BOOST_COMPUTE_CL_CALLBACK
 event_promise_fulfiller_callback(cl_event event, cl_int status, void *user_data)
 {
@@ -75,21 +96,25 @@ event_promise_fulfiller_callback(cl_event event, cl_int status, void *user_data)
 BOOST_AUTO_TEST_CASE(event_to_std_future)
 {
     REQUIRES_OPENCL_VERSION(1,2);
+
+    // enqueue an asynchronous copy to the device
     std::vector<float> vector(1000, 3.14f);
     boost::compute::buffer buffer(context, 1000 * sizeof(float));
     auto event = queue.enqueue_write_buffer_async(
         buffer, 0, 1000 * sizeof(float), vector.data()
     );
+
+    // create a promise and future to be set by the callback
     auto *promise = new std::promise<void>;
     std::future<void> future = promise->get_future();
     event.set_callback(event_promise_fulfiller_callback, CL_COMPLETE, promise);
 
-    // reset the event object (neccessary for intel gpus to fire the callback)
-    event = boost::compute::event();
+    // ensure commands are submitted to the device before waiting
+    queue.flush();
 
+    // wait for future to become ready
     future.wait();
 }
-#endif // BOOST_NO_CXX11_HDR_FUTURE
 #endif // CL_VERSION_1_1
 
 BOOST_AUTO_TEST_SUITE_END()

test/test_function.cpp

@@ -71,8 +71,7 @@ BOOST_AUTO_TEST_CASE(sort_pairs)
     data.push_back(std::make_pair(0, 4.2f));
     data.push_back(std::make_pair(2, 1.0f));
 
-    compute::vector<std::pair<int, float> > vector(data.size());
-    compute::copy(data.begin(), data.end(), vector.begin(), queue);
+    compute::vector<std::pair<int, float> > vector(data.begin(), data.end(), queue);
 
     // sort by first component
     BOOST_COMPUTE_FUNCTION(bool, compare_first, (std::pair<int, float> a, std::pair<int, float> b),

test/test_functional_bind.cpp

@@ -68,7 +68,12 @@ BOOST_AUTO_TEST_CASE(transform_pow_two)
         compute::bind(compute::pow<float>(), 2.0f, _1),
         queue
     );
-    CHECK_RANGE_EQUAL(float, 4, vector, (4.0f, 8.0f, 16.0f, 32.0f));
+
+    compute::copy(vector.begin(), vector.end(), data, queue);
+    BOOST_CHECK_CLOSE(data[0], 4.0f, 1e-4);
+    BOOST_CHECK_CLOSE(data[1], 8.0f, 1e-4);
+    BOOST_CHECK_CLOSE(data[2], 16.0f, 1e-4);
+    BOOST_CHECK_CLOSE(data[3], 32.0f, 1e-4);
 }
 
 BOOST_AUTO_TEST_CASE(find_if_equal)