OSX

.travis.yml

@@ -257,24 +257,12 @@ matrix:
     ############################################################################
     # OSX
     ############################################################################
-    # OSX build with running tests
-    - os: osx
-      compiler: clang
-      env:
-        - ENV_CXX_FLAGS="-Wno-c99-extensions"
-        - RUN_TESTS=true
-    # OSX build without running tests
+
+    # OSX build
     - os: osx
       compiler: clang
       env:
         - ENV_CXX_FLAGS="-Wno-c99-extensions"
-        - RUN_TESTS=false
-
-  allow_failures:
-    - os: osx # OSX build with running tests is allowed to fail
-      env:
-        - ENV_CXX_FLAGS="-Wno-c99-extensions"
-        - RUN_TESTS=true
 
 before_install:
     # Install recent cmake

example/matrix_transpose.cpp

@@ -251,6 +251,14 @@ int main(int argc, char *argv[])
     std::cout << "Local Size: " << TILE_DIM << "x" << BLOCK_ROWS << " threads" << std::endl;
     std::cout << std::endl;
 
+    // On OSX this example does not work on CPU devices
+    #if defined(__APPLE__)
+    if(device.type() & compute::device::cpu) {
+        std::cout << "On OSX this example does not work on CPU devices" << std::endl;
+        return 0;
+    }
+    #endif
+
     const size_t global_work_size[2] = {rows, cols*BLOCK_ROWS/TILE_DIM};
     const size_t local_work_size[2] = {TILE_DIM, BLOCK_ROWS};
 

include/boost/compute/random/bernoulli_distribution.hpp

@@ -11,6 +11,9 @@
 #ifndef BOOST_COMPUTE_RANDOM_BERNOULLI_DISTRIBUTION_HPP
 #define BOOST_COMPUTE_RANDOM_BERNOULLI_DISTRIBUTION_HPP
 
+#include <boost/assert.hpp>
+#include <boost/type_traits.hpp>
+
 #include <boost/compute/command_queue.hpp>
 #include <boost/compute/function.hpp>
 #include <boost/compute/types/fundamental.hpp>
@@ -84,6 +87,11 @@ class bernoulli_distribution
 
 private:
     RealType m_p;
+
+    BOOST_STATIC_ASSERT_MSG(
+        boost::is_floating_point<RealType>::value,
+        "Template argument must be a floating point type"
+    );
 };
 
 } // end compute namespace

include/boost/compute/random/discrete_distribution.hpp

@@ -11,6 +11,9 @@
 #ifndef BOOST_COMPUTE_RANDOM_DISCRETE_DISTRIBUTION_HPP
 #define BOOST_COMPUTE_RANDOM_DISCRETE_DISTRIBUTION_HPP
 
+#include <boost/type_traits.hpp>
+#include <boost/static_assert.hpp>
+
 #include <boost/compute/command_queue.hpp>
 #include <boost/compute/function.hpp>
 #include <boost/compute/algorithm/accumulate.hpp>
@@ -42,8 +45,8 @@ class discrete_distribution
     /// the range [\p first, \p last)
     template<class InputIterator>
     discrete_distribution(InputIterator first, InputIterator last)
-        : m_n(std::distance(first, last)),
-          m_probabilities(std::distance(first, last))
+        : m_n((std::max)(size_t(1), static_cast<size_t>(std::distance(first, last)))),
+          m_probabilities((std::max)(size_t(1), static_cast<size_t>(std::distance(first, last))))
     {
         double sum = 0;
 
@@ -52,9 +55,14 @@ class discrete_distribution
             sum += *iter;
         }
 
-        for(size_t i=0; i<m_n; i++)
+        InputIterator iter = first;
+        m_probabilities[0] = (*iter)/sum;
+        iter++;
+
+        for(size_t i = 1; i < m_n; i++)
         {
-            m_probabilities[i] = m_probabilities[i-1] + first[i]/sum;
+            m_probabilities[i] = m_probabilities[i-1] + (*iter)/sum;
+            iter++;
         }
     }
 
@@ -109,6 +117,11 @@ class discrete_distribution
 private:
     size_t m_n;
     ::std::vector<double> m_probabilities;
+
+    BOOST_STATIC_ASSERT_MSG(
+        boost::is_integral<IntType>::value,
+        "Template argument must be integral"
+    );
 };
 
 } // end compute namespace

include/boost/compute/random/normal_distribution.hpp

@@ -13,6 +13,9 @@
 
 #include <limits>
 
+#include <boost/assert.hpp>
+#include <boost/type_traits.hpp>
+
 #include <boost/compute/command_queue.hpp>
 #include <boost/compute/function.hpp>
 #include <boost/compute/types/fundamental.hpp>
@@ -124,6 +127,11 @@ class normal_distribution
 private:
     RealType m_mean;
     RealType m_stddev;
+
+    BOOST_STATIC_ASSERT_MSG(
+        boost::is_floating_point<RealType>::value,
+        "Template argument must be a floating point type"
+    );
 };
 
 } // end compute namespace

include/boost/compute/random/uniform_int_distribution.hpp

@@ -13,6 +13,9 @@
 
 #include <limits>
 
+#include <boost/type_traits.hpp>
+#include <boost/static_assert.hpp>
+
 #include <boost/compute/command_queue.hpp>
 #include <boost/compute/container/vector.hpp>
 #include <boost/compute/function.hpp>
@@ -103,6 +106,11 @@ class uniform_int_distribution
 private:
     IntType m_a;
     IntType m_b;
+
+    BOOST_STATIC_ASSERT_MSG(
+        boost::is_integral<IntType>::value,
+        "Template argument must be integral"
+    );
 };
 
 } // end compute namespace

include/boost/compute/random/uniform_real_distribution.hpp

@@ -12,6 +12,7 @@
 #define BOOST_COMPUTE_RANDOM_UNIFORM_REAL_DISTRIBUTION_HPP
 
 #include <boost/assert.hpp>
+#include <boost/type_traits.hpp>
 
 #include <boost/compute/command_queue.hpp>
 #include <boost/compute/function.hpp>
@@ -102,6 +103,11 @@ class uniform_real_distribution
 private:
     RealType m_a;
     RealType m_b;
+
+    BOOST_STATIC_ASSERT_MSG(
+        boost::is_floating_point<RealType>::value,
+        "Template argument must be a floating point type"
+    );
 };
 
 } // end compute namespace

test/quirks.hpp

@@ -25,6 +25,12 @@ inline bool is_pocl_device(const boost::compute::device &device)
     return device.platform().name() == "Portable Computing Language";
 }
 
+// returns true if the device is from Apple OpenCL platform
+inline bool is_apple_device(const boost::compute::device &device)
+{
+    return device.platform().name() == "Apple";
+}
+
 // AMD platforms have a bug when using struct assignment. this affects
 // algorithms like fill() when used with pairs/tuples.
 //
@@ -43,6 +49,25 @@ inline bool bug_in_svmmemcpy(const boost::compute::device &device)
     return boost::compute::detail::is_amd_device(device);
 }
 
+// For CPU devices on Apple platform local memory can not be used when work
+// group size is not [1;1;1]. If work group size is greater "Invalid Work Group
+// Size" error is thrown. (Apple OpenCL implementation can sometimes reduce
+// max work group size for other reasons.)
+// When local memory is not used max work group size for CPU devices on Apple
+// platform should be [1024;1;1].
+inline bool is_apple_cpu_device(const boost::compute::device &device)
+{
+    return is_apple_device(device) && (device.type() & ::boost::compute::device::cpu);
+}
+
+// On Apple devices clCreateBuffer does not return NULL and does no set error
+// to CL_INVALID_BUFFER_SIZE when size of the buffer memory object is greater
+// than CL_DEVICE_MAX_MEM_ALLOC_SIZE.
+inline bool bug_in_clcreatebuffer(const boost::compute::device &device)
+{
+    return is_apple_device(device);
+}
+
 // returns true if the device supports image samplers.
 inline bool supports_image_samplers(const boost::compute::device &device)
 {

test/test_command_queue.cpp

@@ -283,6 +283,29 @@ BOOST_AUTO_TEST_CASE(enqueue_kernel_with_extents)
     kernel.set_arg(0, output1);
     kernel.set_arg(1, output2);
 
+    queue.enqueue_nd_range_kernel(kernel, dim(0, 0), dim(4, 4), dim(1, 1));
+
+    CHECK_RANGE_EQUAL(int, 4, output1, (0, 0, 0, 0));
+    CHECK_RANGE_EQUAL(int, 4, output2, (0, 0, 0, 0));
+
+    // Maximum number of work-items that can be specified in each
+    // dimension of the work-group to clEnqueueNDRangeKernel.
+    std::vector<size_t> max_work_item_sizes =
+        device.get_info<CL_DEVICE_MAX_WORK_ITEM_SIZES>();
+
+    if(max_work_item_sizes[0] < size_t(2)) {
+        return;
+    }
+
+    queue.enqueue_nd_range_kernel(kernel, dim(0, 0), dim(4, 4), dim(2, 1));
+
+    CHECK_RANGE_EQUAL(int, 4, output1, (0, 1, 0, 1));
+    CHECK_RANGE_EQUAL(int, 4, output2, (0, 0, 0, 0));
+
+    if(max_work_item_sizes[1] < size_t(2)) {
+        return;
+    }
+
     queue.enqueue_nd_range_kernel(kernel, dim(0, 0), dim(4, 4), dim(2, 2));
 
     CHECK_RANGE_EQUAL(int, 4, output1, (0, 1, 0, 1));

test/test_inplace_reduce.cpp

@@ -17,10 +17,19 @@
 #include <boost/compute/algorithm/detail/inplace_reduce.hpp>
 #include <boost/compute/container/vector.hpp>
 
+#include "quirks.hpp"
 #include "context_setup.hpp"
 
 BOOST_AUTO_TEST_CASE(sum_int)
 {
+    if(is_apple_cpu_device(device)) {
+        std::cerr
+            << "skipping all inplace_reduce tests due to Apple platform"
+            << " behavior when local memory is used on a CPU device"
+            << std::endl;
+        return;
+    }
+
     int data[] = { 1, 5, 3, 4, 9, 3, 5, 3 };
     boost::compute::vector<int> vector(data, data + 8, queue);
 
@@ -43,6 +52,10 @@ BOOST_AUTO_TEST_CASE(sum_int)
 
 BOOST_AUTO_TEST_CASE(multiply_int)
 {
+    if(is_apple_cpu_device(device)) {
+        return;
+    }
+
     int data[] = { 1, 5, 3, 4, 9, 3, 5, 3 };
     boost::compute::vector<int> vector(data, data + 8, queue);
 
@@ -65,6 +78,10 @@ BOOST_AUTO_TEST_CASE(multiply_int)
 
 BOOST_AUTO_TEST_CASE(reduce_iota)
 {
+    if(is_apple_cpu_device(device)) {
+        return;
+    }
+
     // 1 value
     boost::compute::vector<int> vector(1, context);
     boost::compute::iota(vector.begin(), vector.end(), int(0), queue);

test/test_lambda.cpp

@@ -307,7 +307,7 @@ BOOST_AUTO_TEST_CASE(lambda_get_tuple)
     vector.push_back(boost::make_tuple(5, 'c', 5.6f), queue);
     vector.push_back(boost::make_tuple(7, 'd', 7.8f), queue);
 
-    // extract first compoenent of each tuple
+    // extract first component of each tuple
     compute::vector<int> first_component(4, context);
     compute::transform(
         vector.begin(),
@@ -318,7 +318,7 @@ BOOST_AUTO_TEST_CASE(lambda_get_tuple)
     );
     CHECK_RANGE_EQUAL(int, 4, first_component, (1, 3, 5, 7));
 
-    // extract second compoenent of each tuple
+    // extract second component of each tuple
     compute::vector<char> second_component(4, context);
     compute::transform(
         vector.begin(),
@@ -329,7 +329,7 @@ BOOST_AUTO_TEST_CASE(lambda_get_tuple)
     );
     CHECK_RANGE_EQUAL(char, 4, second_component, ('a', 'b', 'c', 'd'));
 
-    // extract third compoenent of each tuple
+    // extract third component of each tuple
     compute::vector<float> third_component(4, context);
     compute::transform(
         vector.begin(),