Update SDK to ComputeCpp v0.9.1

The change to the Windows debug library has been reverted. It was fixed in the 0.9.1 release. Additionally, there have been some minor API updates from the SYCL working group in Khronos, resulting in some small changes to the samples. For details, see the blogpost on the Codeplay website: https://www.codeplay.com/portal/07-20-18-computecpp-ce-0-9-1-release

Author: cjdb

.travis.yml

@@ -45,6 +45,6 @@ script:
  ###########################
  - mkdir build
  - cd build
- - LD_LIBRARY_PATH=/tmp/OpenCL-ICD-Loader/build/lib cmake ../ -DComputeCpp_DIR=/tmp/computecpp -DCOMPUTECPP_SDK_BUILD_TESTS=1 -DCMAKE_CXX_FLAGS="-D_GLIBCXX_USE_CXX11_ABI=0"
+ - cmake ../ -DComputeCpp_DIR=/tmp/computecpp -DCOMPUTECPP_SDK_BUILD_TESTS=1 -DCMAKE_CXX_FLAGS="-D_GLIBCXX_USE_CXX11_ABI=0"
  - make -j2
  - COMPUTECPP_TARGET="host" ctest -V

.travis/before_install.sh

@@ -6,11 +6,3 @@ sudo apt-get update -q
 sudo apt-get install ocl-icd-libopencl1 ocl-icd-dev opencl-headers -y
 # Use gcc 5 as default
 sudo update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-5 60 --slave /usr/bin/g++ g++ /usr/bin/g++-5
-cd /tmp
-git clone https://github.com/KhronosGroup/OpenCL-ICD-Loader
-cd OpenCL-ICD-Loader
-rm inc/README.txt
-git clone https://github.com/KhronosGroup/OpenCL-Headers inc
-mkdir -p build && cd build
-cmake ../ -DCMAKE_BUILD_TYPE=Release
-make -j2

.travis/build_computecpp.sh

@@ -3,7 +3,7 @@
 set -ev
 
 # Get ComputeCpp
-wget --no-check-certificate https://computecpp.codeplay.com/downloads/computecpp-ce/0.9.0/ubuntu-14.04-64bit.tar.gz
+wget --no-check-certificate https://computecpp.codeplay.com/downloads/computecpp-ce/0.9.1/ubuntu-14.04-64bit.tar.gz
 mkdir -p /tmp/computecpp
 tar -xzf ubuntu-14.04-64bit.tar.gz -C /tmp/computecpp --strip-components=1
 # Workaround for C99 definition conflict

cmake/Modules/FindComputeCpp.cmake

@@ -84,7 +84,7 @@ find_library(COMPUTECPP_RUNTIME_LIBRARY
  DOC "ComputeCpp Runtime Library")
 
 find_library(COMPUTECPP_RUNTIME_LIBRARY_DEBUG
- NAMES ComputeCpp ComputeCpp_vs2015
+ NAMES ComputeCpp ComputeCpp_vs2015_d
  PATHS ${ComputeCpp_ROOT_DIR}
  PATH_SUFFIXES lib
  DOC "ComputeCpp Debug Runtime Library")
@@ -148,13 +148,6 @@ if(NOT ComputeCpp_FOUND)
  return()
 endif()
 
-if(MSVC)
- message(WARNING " The Debug ComputeCpp library is missing! You might
- experience linker errors or crashes when building a Debug
- configuration. Please file an issue on Github if you do.
- This will be fixed in a subsequent release.")
-endif()
-
 if(CMAKE_CROSSCOMPILING)
  if(NOT SDK_DONT_USE_TOOLCHAIN)
  list(APPEND COMPUTECPP_DEVICE_COMPILER_FLAGS --gcc-toolchain=${SDK_TOOLCHAIN_DIR})

samples/gaussian-blur/gaussian-blur.cpp

@@ -162,13 +162,13 @@ int main(int argc, char* argv[]) {
  for (int x = - offset; x < offset; x++) {
  for (int y = - offset; y < offset; y++) {
  auto inputCoords =
- int2(itemID.get_global(0) + x, itemID.get_global(1) + y);
+ int2(itemID.get_global_id(0) + x, itemID.get_global_id(1) + y);
  newPixel += inPtr.read(inputCoords, smpl) *
  globalGaussian[x + offset][y + offset];
  }
  }
 
- auto outputCoords = int2(itemID.get_global(0), itemID.get_global(1));
+ auto outputCoords = int2(itemID.get_global_id(0), itemID.get_global_id(1));
  newPixel.w() = 1.f;
  outPtr.write(outputCoords, newPixel);
  });

samples/matrix-multiply/matrix-multiply.cpp

@@ -171,8 +171,8 @@ bool local_mxm(cl::sycl::queue& q, T* MA, T* MB, T* MC, int matSize) {
  int blockY = it.get_group(1);
 
  // Current local item
- int localX = it.get_local(0);
- int localY = it.get_local(1);
+ int localX = it.get_local_id(0);
+ int localY = it.get_local_id(1);
 
  // Start in the A matrix
  int a_start = matSize * blockSize * blockY;
@@ -203,7 +203,7 @@ bool local_mxm(cl::sycl::queue& q, T* MA, T* MB, T* MC, int matSize) {
  it.barrier(access::fence_space::local_space);
  }
  auto elemIndex =
- it.get_global(1) * it.get_global_range()[0] + it.get_global(0);
+ it.get_global_id(1) * it.get_global_range()[0] + it.get_global_id(0);
  // Each thread updates its position
  pC[elemIndex] = tmp;
  });

samples/parallel-for/parallel-for.cpp

@@ -76,7 +76,7 @@ int main() {
  * item::get_global() to retrieve the global id as an id<1>.
  * This particular kernel will set the ith element to the value
  * of i. */
- ptr[item.get_global()] = item.get_global()[0];
+ ptr[item.get_global_id()] = item.get_global_id()[0];
  });
 
  /* We call the parallel_for() API with two parameters; the nd_range

samples/reduction/reduction.cpp

@@ -91,8 +91,8 @@ T sycl_reduce(const std::vector<T>& v) {
  * parameter, since the code requires barriers for correctness. */
  h.parallel_for<class sycl_reduction>(
  r, [aI, scratch, local, length](cl::sycl::nd_item<1> id) {
- size_t globalid = id.get_global(0);
- size_t localid = id.get_local(0);
+ size_t globalid = id.get_global_id(0);
+ size_t localid = id.get_local_id(0);
 
  /* All threads collectively read from global memory into local.
  * The barrier ensures all threads' IO is resolved before

samples/tiled-convolution/tiled-convolution.cpp

@@ -76,8 +76,8 @@ class conv {
  total_size(total_size_),
  fil_size(fil_size_) {}
  void inline operator()(cl::sycl::nd_item<2> item_id) {
- index_t id_m = item_id.get_global(0); // global id with offset m
- index_t id_n = item_id.get_global(1); // global id with offser n
+ index_t id_m = item_id.get_global_id(0); // global id with offset m
+ index_t id_n = item_id.get_global_id(1); // global id with offser n
  // input index m , filter index fil_m , input index n, filter index fil_n
  index_t m, fil_m, n, fil_n;
  typename write_accessor_t::value_type val = 0.0;