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"
@@ -83,8 +82,8 @@ inline bool dispatch_lexicographical_compare(InputIterator1 first1,
     kernel lexicographical_compare_kernel(lexicographical_compare_program,
                                           "lexicographical_compare");
 
-    lexicographical_compare_kernel.set_arg(0, (uint)iterator_size1);
-    lexicographical_compare_kernel.set_arg(1, (uint)iterator_size2);
+    lexicographical_compare_kernel.set_arg<uint_>(0, iterator_size1);
+    lexicographical_compare_kernel.set_arg<uint_>(1, iterator_size2);
     lexicographical_compare_kernel.set_arg(2, first1.get_buffer());
     lexicographical_compare_kernel.set_arg(3, first2.get_buffer());
     lexicographical_compare_kernel.set_arg(4, result_vector.get_buffer());

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_int_distribution.hpp

@@ -74,7 +74,7 @@ class uniform_int_distribution
         vector<g_result_type> tmp(size, queue.get_context());
         vector<g_result_type> tmp2(size, queue.get_context());
 
-        uint_ bound = ((uint(-1))/(m_b-m_a+1))*(m_b-m_a+1);
+        uint_ bound = ((uint_(-1))/(m_b-m_a+1))*(m_b-m_a+1);
 
         buffer_iterator<g_result_type> tmp2_iter;
 

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>()

perf/perf_discrete_distribution.cpp

@@ -30,12 +30,12 @@ int main(int argc, char *argv[])
     compute::context context(device);
     compute::command_queue queue(context, device);
 
-    compute::vector<unsigned int> vector(PERF_N, context);
+    compute::vector<compute::uint_> vector(PERF_N, context);
 
     int weights[] = {1, 1};
 
     compute::default_random_engine rng(queue);
-    compute::discrete_distribution<uint> dist(weights, weights+2);
+    compute::discrete_distribution<compute::uint_> dist(weights, weights+2);
 
     perf_timer t;
     t.start();

perf/perf_linear_congruential_engine.cpp

@@ -31,10 +31,10 @@ int main(int argc, char *argv[])
     compute::command_queue queue(context, device);
 
     // create vector on the device
-    compute::vector<unsigned int> vector(PERF_N, context);
+    compute::vector<compute::uint_> vector(PERF_N, context);
 
     // create mersenne twister engine
-    compute::linear_congruential_engine<uint> rng(queue);
+    compute::linear_congruential_engine<compute::uint_> rng(queue);
 
     // generate random numbers
     perf_timer t;

perf/perf_uniform_int_distribution.cpp

@@ -30,10 +30,10 @@ int main(int argc, char *argv[])
     compute::context context(device);
     compute::command_queue queue(context, device);
 
-    compute::vector<unsigned int> vector(PERF_N, context);
+    compute::vector<compute::uint_> vector(PERF_N, context);
 
     compute::default_random_engine rng(queue);
-    compute::uniform_int_distribution<uint> dist(0, 1);
+    compute::uniform_int_distribution<compute::uint_> dist(0, 1);
 
     perf_timer t;
     t.start();

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);