Commit 95e01a3d authored by Joseph Noir's avatar Joseph Noir

Enable discovery of multiple platforms.

* The metainfo can handle multiple platforms.
* Devices are now organized by platform.
* More consistent naming (opencl_metainfo -> metainfo, ...).
* The metainfo offers the new function get_device_if to find devices by
  their characteristics, (name, maximum work-items, etc.).
parent 7d1ea179
......@@ -11,9 +11,11 @@ if (OpenCL_LIBRARIES)
# list cpp files excluding platform-dependent files
set (LIBCAF_OPENCL_SRCS
src/global.cpp
src/opencl_metainfo.cpp
src/metainfo.cpp
src/program.cpp
src/opencl_err.cpp)
src/opencl_err.cpp
src/platform.cpp
src/device.cpp)
# build shared library if not compiling static only
if(NOT CAF_BUILD_STATIC_ONLY)
add_library(libcaf_opencl_shared SHARED ${LIBCAF_OPENCL_SRCS}
......
......@@ -46,14 +46,14 @@
namespace caf {
namespace opencl {
class opencl_metainfo;
class metainfo;
template <class List>
struct function_sig_from_outputs;
template <class... Ts>
struct function_sig_from_outputs<detail::type_list<Ts...>> {
using type = std::function<message (Ts...)>;
using type = std::function<message (Ts&...)>;
};
template <class T, class List>
......@@ -113,14 +113,14 @@ public:
};
check_vec(config.offsets(), "offsets");
check_vec(config.local_dimensions(), "local dimensions");
auto itr = prog.available_kernels_.find(kernel_name);
if (itr == prog.available_kernels_.end()) {
cl_int err = 0;
kernel_ptr kernel;
kernel.reset(clCreateKernel(prog.program_.get(), kernel_name, &err), false);
if (err != CL_SUCCESS)
return nullptr;
} else {
return new actor_facade(prog, itr->second, config,
std::move(map_args), std::move(map_result),
std::forward_as_tuple(xs...));
}
return new actor_facade(prog, kernel, config,
std::move(map_args), std::move(map_result),
std::forward_as_tuple(xs...));
}
void enqueue(const actor_addr &sender, message_id mid, message content,
......@@ -136,14 +136,14 @@ public:
if (! content.match_elements(input_types{})) {
return;
}
response_promise handle{this->address(), sender, mid.response_id()};
response_promise hdl{this->address(), sender, mid.response_id()};
evnt_vec events;
args_vec input_buffers;
args_vec output_buffers;
size_vec result_sizes;
add_kernel_arguments(events, input_buffers, output_buffers,
result_sizes, content, indices);
auto cmd = make_counted<command_type>(handle, this,
auto cmd = make_counted<command_type>(hdl, this,
std::move(events),
std::move(input_buffers),
std::move(output_buffers),
......@@ -206,7 +206,7 @@ public:
mem_ptr tmp;
tmp.reset(buffer, false);
input_buffers.push_back(tmp);
v1callcl(CAF_CLF(clSetKernelArg), kernel_.get(), I,
v1callcl(CAF_CLF(clSetKernelArg), kernel_.get(), static_cast<unsigned>(I),
sizeof(cl_mem), static_cast<void*>(&input_buffers.back()));
}
......@@ -227,11 +227,11 @@ public:
mem_ptr tmp;
tmp.reset(buffer, false);
output_buffers.push_back(tmp);
v1callcl(CAF_CLF(clSetKernelArg), kernel_.get(), I,
v1callcl(CAF_CLF(clSetKernelArg), kernel_.get(), static_cast<unsigned>(I),
sizeof(cl_mem), static_cast<void*>(&output_buffers.back()));
sizes.push_back(size);
}
template <long I, class T>
void create_buffer(const out<T>& wrapper, evnt_vec&, size_vec& sizes,
args_vec&, args_vec& output_buffers, message& msg) {
......@@ -244,7 +244,7 @@ public:
mem_ptr tmp;
tmp.reset(buffer, false);
output_buffers.push_back(tmp);
v1callcl(CAF_CLF(clSetKernelArg), kernel_.get(), I,
v1callcl(CAF_CLF(clSetKernelArg), kernel_.get(), static_cast<unsigned>(I),
sizeof(cl_mem), static_cast<void*>(&output_buffers.back()));
sizes.push_back(size);
}
......
......@@ -95,8 +95,11 @@ public:
enqueue_read_buffers(event_k, detail::get_indices(result_buffers_));
cl_event marker;
#if defined(__APPLE__)
err = clEnqueueMarkerWithWaitList(queue_.get(), mem_out_events_.size(),
mem_out_events_.data(), &marker);
err = clEnqueueMarkerWithWaitList(
queue_.get(),
static_cast<cl_uint>(mem_out_events_.size()),
mem_out_events_.data(), &marker
);
#else
err = clEnqueueMarker(queue_.get(), &marker);
#endif
......
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2015 *
* Dominik Charousset <dominik.charousset (at) haw-hamburg.de> *
* *
* Distributed under the terms and conditions of the BSD 3-Clause License or *
* (at your option) under the terms and conditions of the Boost Software *
* License 1.0. See accompanying files LICENSE and LICENSE_ALTERNATIVE. *
* *
* If you did not receive a copy of the license files, see *
* http://opensource.org/licenses/BSD-3-Clause and *
* http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/
#ifndef CAF_OPENCL_DEVICE_HPP
#define CAF_OPENCL_DEVICE_HPP
#include <vector>
#include "caf/opencl/global.hpp"
#include "caf/opencl/smart_ptr.hpp"
namespace caf {
namespace opencl {
class program;
class device {
friend class program;
public:
/// Intialize a new device in a context using a sepcific device_id
static device create(context_ptr context, device_ptr device_id, unsigned id);
/// Get the id assigned by caf
inline unsigned get_id() const;
/// Returns device info on CL_DEVICE_ADDRESS_BITS
inline cl_uint get_address_bits() const;
/// Returns device info on CL_DEVICE_ENDIAN_LITTLE
inline cl_bool get_little_endian() const;
/// Returns device info on CL_DEVICE_GLOBAL_MEM_CACHE_SIZE
inline cl_ulong get_global_mem_cache_size() const;
/// Returns device info on CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE
inline cl_uint get_global_mem_cacheline_size() const;
/// Returns device info on CL_DEVICE_GLOBAL_MEM_SIZE
inline cl_ulong get_global_mem_size() const;
/// Returns device info on CL_DEVICE_HOST_UNIFIED_MEMORY
inline cl_bool get_host_unified_memory() const;
/// Returns device info on CL_DEVICE_LOCAL_MEM_SIZE
inline cl_ulong get_local_mem_size() const;
/// Returns device info on CL_DEVICE_LOCAL_MEM_TYPE
inline cl_ulong get_local_mem_type() const;
/// Returns device info on CL_DEVICE_MAX_CLOCK_FREQUENCY
inline cl_uint get_max_clock_frequency() const;
/// Returns device info on CL_DEVICE_MAX_COMPUTE_UNITS
inline cl_uint get_max_compute_units() const;
/// Returns device info on CL_DEVICE_MAX_CONSTANT_ARGS
inline cl_uint get_max_constant_args() const;
/// Returns device info on CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE
inline cl_ulong get_max_constant_buffer_size() const;
/// Returns device info on CL_DEVICE_MAX_MEM_ALLOC_SIZE
inline cl_ulong get_max_mem_alloc_size() const;
/// Returns device info on CL_DEVICE_MAX_PARAMETER_SIZE
inline size_t get_max_parameter_size() const;
/// Returns device info on CL_DEVICE_MAX_WORK_GROUP_SIZE
inline size_t get_max_work_group_size() const;
/// Returns device info on CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS
inline cl_uint get_max_work_item_dimensions() const;
/// Returns device info on CL_DEVICE_PROFILING_TIMER_RESOLUTION
inline size_t get_profiling_timer_resolution() const;
/// Returns device info on CL_DEVICE_MAX_WORK_ITEM_SIZES
inline const dim_vec& get_max_work_item_sizes() const;
/// Returns device info on CL_DEVICE_TYPE
inline device_type get_device_type() const;
/// Returns device info on CL_DEVICE_EXTENSIONS
inline const std::vector<std::string>& get_extensions() const;
/// Returns device info on CL_DEVICE_OPENCL_C_VERSION
inline const std::string& get_opencl_c_version() const;
/// Returns device info on CL_DEVICE_VENDOR
inline const std::string& get_device_vendor() const;
/// Returns device info on CL_DEVICE_VERSION
inline const std::string& get_device_version() const;
/// Returns device info on CL_DRIVER_VERSION
inline const std::string& get_driver_version() const;
/// Returns device info on CL_DEVICE_NAME
inline const std::string& get_name() const;
private:
device(device_ptr device_id, command_queue_ptr queue, context_ptr context,
unsigned id);
template <class T>
static T info(device_ptr device_id, unsigned info_flag);
static std::string info_string(device_ptr device_id, unsigned info_flag);
device_ptr device_id_;
command_queue_ptr command_queue_;
context_ptr context_;
unsigned id_;
bool profiling_enabled_; // CL_DEVICE_QUEUE_PROPERTIES
bool out_of_order_execution_; // CL_DEVICE_QUEUE_PROPERTIES
cl_uint address_bits_; // CL_DEVICE_ADDRESS_BITS
cl_bool little_endian_; // CL_DEVICE_ENDIAN_LITTLE
cl_ulong global_mem_cache_size_; // CL_DEVICE_GLOBAL_MEM_CACHE_SIZE
cl_uint global_mem_cacheline_size_; // CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE
cl_ulong global_mem_size_; // CL_DEVICE_GLOBAL_MEM_SIZE
cl_bool host_unified_memory_; // CL_DEVICE_HOST_UNIFIED_MEMORY
cl_ulong local_mem_size_; // CL_DEVICE_LOCAL_MEM_SIZE
cl_uint local_mem_type_; // CL_DEVICE_LOCAL_MEM_TYPE
cl_uint max_clock_frequency_; // CL_DEVICE_MAX_CLOCK_FREQUENCY
cl_uint max_compute_units_; // CL_DEVICE_MAX_COMPUTE_UNITS
cl_uint max_constant_args_; // CL_DEVICE_MAX_CONSTANT_ARGS
cl_ulong max_constant_buffer_size_; // CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE
cl_ulong max_mem_alloc_size_; // CL_DEVICE_MAX_MEM_ALLOC_SIZE
size_t max_parameter_size_; // CL_DEVICE_MAX_PARAMETER_SIZE
size_t max_work_group_size_; // CL_DEVICE_MAX_WORK_GROUP_SIZE
cl_uint max_work_item_dimensions_; // CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS
size_t profiling_timer_resolution_; // CL_DEVICE_PROFILING_TIMER_RESOLUTION
dim_vec max_work_item_sizes_; // CL_DEVICE_MAX_WORK_ITEM_SIZES
device_type device_type_; // CL_DEVICE_TYPE
std::vector<std::string> extensions_; // CL_DEVICE_EXTENSIONS
std::string opencl_c_version_; // CL_DEVICE_OPENCL_C_VERSION
std::string device_vendor_; // CL_DEVICE_VENDOR
std::string device_version_; // CL_DEVICE_VERSION
std::string driver_version_; // CL_DRIVER_VERSION
std::string name_; // CL_DEVICE_NAME
};
/******************************************************************************\
* implementation of inline member functions *
\******************************************************************************/
inline unsigned device::get_id() const {
return id_;
}
inline cl_uint device::get_address_bits() const {
return address_bits_;
}
inline cl_bool device::get_little_endian() const {
return little_endian_;
}
inline cl_ulong device::get_global_mem_cache_size() const {
return global_mem_cache_size_;
}
inline cl_uint device::get_global_mem_cacheline_size() const {
return global_mem_cacheline_size_;
}
inline cl_ulong device::get_global_mem_size() const {
return global_mem_size_;
}
inline cl_bool device::get_host_unified_memory() const {
return host_unified_memory_;
}
inline cl_ulong device::get_local_mem_size() const {
return local_mem_size_;
}
inline cl_ulong device::get_local_mem_type() const {
return local_mem_size_;
}
inline cl_uint device::get_max_clock_frequency() const {
return max_clock_frequency_;
}
inline cl_uint device::get_max_compute_units() const {
return max_compute_units_;
}
inline cl_uint device::get_max_constant_args() const {
return max_constant_args_;
}
inline cl_ulong device::get_max_constant_buffer_size() const {
return max_constant_buffer_size_;
}
inline cl_ulong device::get_max_mem_alloc_size() const {
return max_mem_alloc_size_;
}
inline size_t device::get_max_parameter_size() const {
return max_parameter_size_;
}
inline size_t device::get_max_work_group_size() const {
return max_work_group_size_;
}
inline cl_uint device::get_max_work_item_dimensions() const {
return max_work_item_dimensions_;
}
inline size_t device::get_profiling_timer_resolution() const {
return profiling_timer_resolution_;
}
inline const dim_vec& device::get_max_work_item_sizes() const {
return max_work_item_sizes_;
}
inline device_type device::get_device_type() const {
return device_type_;
}
inline const std::vector<std::string>& device::get_extensions() const {
return extensions_;
}
inline const std::string& device::get_opencl_c_version() const {
return opencl_c_version_;
}
inline const std::string& device::get_device_vendor() const {
return device_vendor_;
}
inline const std::string& device::get_device_version() const {
return device_version_;
}
inline const std::string& device::get_driver_version() const {
return driver_version_;
}
inline const std::string& device::get_name() const {
return name_;
}
} // namespace opencl
} // namespace caf
#endif // CAF_OPENCL_DEVICE_HPP
......@@ -21,6 +21,7 @@
#define CAF_OPENCL_GLOBAL_HPP
#include <string>
#include <iostream>
#include "caf/config.hpp"
#include "caf/detail/limited_vector.hpp"
......@@ -40,6 +41,18 @@ cl_int clRetainDeviceDummy(cl_device_id);
namespace caf {
namespace opencl {
enum device_type {
def = CL_DEVICE_TYPE_DEFAULT,
cpu = CL_DEVICE_TYPE_CPU,
gpu = CL_DEVICE_TYPE_GPU,
accelerator = CL_DEVICE_TYPE_ACCELERATOR,
custom = CL_DEVICE_TYPE_CUSTOM,
all = CL_DEVICE_TYPE_ALL
};
std::ostream& operator<<(std::ostream& os, device_type dev);
device_type device_type_from_ulong(cl_ulong dev);
/// A vector of up to three elements used for OpenCL dimensions.
using dim_vec = detail::limited_vector<size_t, 3>;
......
......@@ -17,8 +17,8 @@
* http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/
#ifndef CAF_OPENCL_METAINFO_HPP
#define CAF_OPENCL_METAINFO_HPP
#ifndef CAF_METAINFO_HPP
#define CAF_METAINFO_HPP
#include <atomic>
#include <vector>
......@@ -26,44 +26,61 @@
#include <functional>
#include "caf/all.hpp"
#include "caf/config.hpp"
#include "caf/optional.hpp"
#include "caf/opencl/device.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/program.hpp"
#include "caf/opencl/platform.hpp"
#include "caf/opencl/smart_ptr.hpp"
#include "caf/opencl/device_info.hpp"
#include "caf/opencl/actor_facade.hpp"
//#include "caf/detail/singleton_mixin.hpp"
//#include "caf/detail/singleton_manager.hpp"
#include "caf/detail/singletons.hpp"
namespace caf {
namespace opencl {
class opencl_metainfo : public detail::abstract_singleton {
class metainfo : public detail::abstract_singleton {
friend class program;
friend class detail::singletons;
friend command_queue_ptr get_command_queue(uint32_t id);
public:
const std::vector<device_info> get_devices() const;
/// Get a list of all available devices. This is depricated, use the more specific
/// get_deivce and get_deivce_if functions.
/// (Returns only devices of the first discovered platform).
const std::vector<device>& get_devices() const CAF_DEPRECATED;
/// Get the device with id. These ids are assigned sequientally to all available devices.
const optional<const device&> get_device(size_t id = 0) const;
/// Get the first device that satisfies the predicate.
/// The predicate should accept a `const device&` and return a bool;
template <class UnaryPredicate>
const optional<const device&> get_device_if(UnaryPredicate p) const {
for (auto& pl : platforms_) {
for (auto& dev : pl.get_devices()) {
if (p(dev))
return dev;
}
}
return none;
}
/// Get opencl_metainfo instance.
static opencl_metainfo* instance();
/// Get metainfo instance.
static metainfo* instance();
private:
opencl_metainfo() = default;
metainfo() = default;
void stop() override;
void initialize() override;
void dispose() override;
context_ptr context_;
std::vector<device_info> devices_;
std::vector<platform> platforms_;
};
} // namespace opencl
} // namespace caf
#endif // CAF_OPENCL_METAINFO_HPP
#endif // CAF_METAINFO_HPP
......@@ -17,57 +17,61 @@
* http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/
#ifndef CAF_OPENCL_DEVICE_INFO_HPP
#define CAF_OPENCL_DEVICE_INFO_HPP
#ifndef CAF_OPENCL_PLATFORM_HPP
#define CAF_OPENCL_PLATFORM_HPP
#include "caf/opencl/global.hpp"
#include "caf/opencl/program.hpp"
#include "caf/opencl/smart_ptr.hpp"
#include <caf/opencl/device.hpp>
namespace caf {
namespace opencl {
class device_info {
class platform {
friend class program;
public:
device_info(device_ptr device, command_queue_ptr queue,
size_t work_group_size, cl_uint dimensons,
const dim_vec& items_per_dimension)
: max_work_group_size_(work_group_size),
max_dimensions_(dimensons),
max_work_items_per_dim_(items_per_dimension),
device_(device),
cmd_queue_(queue) {}
inline size_t get_max_work_group_size();
inline cl_uint get_max_dimensions();
inline dim_vec get_max_work_items_per_dim();
inline const std::vector<device>& get_devices() const;
inline const std::string& get_name() const;
inline const std::string& get_vendor() const;
inline const std::string& get_version() const;
static platform create(cl_platform_id platform_id, unsigned start_id);
private:
size_t max_work_group_size_;
cl_uint max_dimensions_;
dim_vec max_work_items_per_dim_;
device_ptr device_;
command_queue_ptr cmd_queue_;
platform(cl_platform_id platform_id, context_ptr context,
std::string name, std::string vendor, std::string version,
std::vector<device> devices);
static std::string platform_info(cl_platform_id platform_id,
unsigned info_flag);
cl_platform_id platform_id_;
context_ptr context_;
std::string name_;
std::string vendor_;
std::string version_;
std::vector<device> devices_;
};
/******************************************************************************\
* implementation of inline member functions *
\******************************************************************************/
inline size_t device_info::get_max_work_group_size() {
return max_work_group_size_;
inline const std::vector<device>& platform::get_devices() const {
return devices_;
}
inline cl_uint device_info::get_max_dimensions() { return max_dimensions_; }
inline const std::string& platform::get_name() const {
return name_;
}
inline dim_vec device_info::get_max_work_items_per_dim() {
return max_work_items_per_dim_;
inline const std::string& platform::get_vendor() const {
return vendor_;
}
inline const std::string& platform::get_version() const {
return version_;
}
} // namespace opencl
} // namespace caf
#endif // CAF_OPENCL_DEVICE_INFO_HPP
#endif // CAF_OPENCL_PLTFORM_HPP
......@@ -20,8 +20,10 @@
#ifndef CAF_OPENCL_PROGRAM_HPP
#define CAF_OPENCL_PROGRAM_HPP
#include <map>
#include <memory>
#include "caf/opencl/device.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/smart_ptr.hpp"
......@@ -44,14 +46,18 @@ public:
static program create(const char* kernel_source,
const char* options = nullptr, uint32_t device_id = 0);
/// @brief Factory method, that creates a caf::opencl::program
/// from a given @p kernel_source.
/// @returns A program object.
static program create(const char* kernel_source,
const char* options, const device& dev);
private:
program(context_ptr context, command_queue_ptr queue, program_ptr program,
std::map<std::string,kernel_ptr> available_kernels);
program(context_ptr context, command_queue_ptr queue, program_ptr prog);
context_ptr context_;
program_ptr program_;
command_queue_ptr queue_;
std::map<std::string,kernel_ptr> available_kernels_;
// save CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE
};
} // namespace opencl
......
......@@ -29,10 +29,10 @@
#include "caf/detail/limited_vector.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/metainfo.hpp"
#include "caf/opencl/arguments.hpp"
#include "caf/opencl/actor_facade.hpp"
#include "caf/opencl/spawn_config.hpp"
#include "caf/opencl/opencl_metainfo.hpp"
namespace caf {
......
......@@ -153,7 +153,7 @@ void multiplier(event_based_actor* self) {
);
};
auto box_res = [] (fvec result) -> message {
auto box_res = [] (fvec& result) -> message {
return make_message(matrix_type{move(result)});
};
......
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2015 *
* Dominik Charousset <dominik.charousset (at) haw-hamburg.de> *
* *
* Distributed under the terms and conditions of the BSD 3-Clause License or *
* (at your option) under the terms and conditions of the Boost Software *
* License 1.0. See accompanying files LICENSE and LICENSE_ALTERNATIVE. *
* *
* If you did not receive a copy of the license files, see *
* http://opensource.org/licenses/BSD-3-Clause and *
* http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/
#include <iostream>
#include "caf/string_algorithms.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/device.hpp"
#include "caf/opencl/opencl_err.hpp"
using namespace std;
namespace caf {
namespace opencl {
device device::create(context_ptr context, device_ptr device_id, unsigned id) {
CAF_LOGF_DEBUG("creating device for opencl device id '"
<< device_id.get() << "' with id '" << id << "'");
// look up properties we need to create the command queue
auto supported = info<cl_ulong>(device_id, CL_DEVICE_QUEUE_PROPERTIES);
bool profiling = supported & CL_QUEUE_PROFILING_ENABLE;
bool out_of_order = supported & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE;
unsigned properties = profiling ? CL_QUEUE_PROFILING_ENABLE : 0;
properties |= out_of_order ? CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE : 0;
// create the command queue
command_queue_ptr command_queue;
command_queue.reset(v2get(CAF_CLF(clCreateCommandQueue), context.get(),
device_id.get(), properties),
false);
// create the device
device dev{device_id, command_queue, context, id};
// look up device properties
dev.address_bits_ = info<cl_uint>(device_id, CL_DEVICE_ADDRESS_BITS);
dev.little_endian_ = info<cl_bool>(device_id, CL_DEVICE_ENDIAN_LITTLE);
dev.global_mem_cache_size_ =
info<cl_ulong>(device_id, CL_DEVICE_GLOBAL_MEM_CACHE_SIZE);
dev.global_mem_cacheline_size_ =
info<cl_uint>(device_id, CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE);
dev.global_mem_size_ =
info<cl_ulong >(device_id,CL_DEVICE_GLOBAL_MEM_SIZE);
dev.host_unified_memory_ =
info<cl_bool>(device_id, CL_DEVICE_HOST_UNIFIED_MEMORY);
dev.local_mem_size_ =
info<cl_ulong>(device_id, CL_DEVICE_LOCAL_MEM_SIZE);
dev.local_mem_type_ =
info<cl_uint>(device_id, CL_DEVICE_LOCAL_MEM_TYPE);
dev.max_clock_frequency_ =
info<cl_uint>(device_id, CL_DEVICE_MAX_CLOCK_FREQUENCY);
dev.max_compute_units_ =
info<cl_uint>(device_id, CL_DEVICE_MAX_COMPUTE_UNITS);
dev.max_constant_args_ =
info<cl_uint>(device_id, CL_DEVICE_MAX_CONSTANT_ARGS);
dev.max_constant_buffer_size_ =
info<cl_ulong>(device_id, CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE);
dev.max_mem_alloc_size_ =
info<cl_ulong>(device_id, CL_DEVICE_MAX_MEM_ALLOC_SIZE);
dev.max_parameter_size_ =
info<size_t>(device_id, CL_DEVICE_MAX_PARAMETER_SIZE);
dev.max_work_group_size_ =
info<size_t>(device_id, CL_DEVICE_MAX_WORK_GROUP_SIZE);
dev.max_work_item_dimensions_ =
info<cl_uint>(device_id, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS);
dev.profiling_timer_resolution_ =
info<size_t>(device_id, CL_DEVICE_PROFILING_TIMER_RESOLUTION);
dev.max_work_item_sizes_.resize(dev.max_work_item_dimensions_);
clGetDeviceInfo(device_id.get(), CL_DEVICE_MAX_WORK_ITEM_SIZES,
sizeof(size_t) * dev.max_work_item_dimensions_,
dev.max_work_item_sizes_.data(), nullptr);
dev.device_type_ =
device_type_from_ulong(info<cl_ulong>(device_id, CL_DEVICE_TYPE));
string extensions = info_string(device_id, CL_DEVICE_EXTENSIONS);
split(dev.extensions_, extensions, " ", false);
dev.opencl_c_version_ = info_string(device_id, CL_DEVICE_EXTENSIONS);
dev.device_vendor_ = info_string(device_id, CL_DEVICE_VENDOR);
dev.device_version_ = info_string(device_id, CL_DEVICE_VERSION);
dev.name_ = info_string(device_id, CL_DEVICE_NAME);
return dev;
}
template <class T>
T device::info(device_ptr device_id, unsigned info_flag) {
T value;
clGetDeviceInfo(device_id.get(), info_flag, sizeof(T), &value, nullptr);
return value;
}
string device::info_string(device_ptr device_id, unsigned info_flag) {
size_t size;
clGetDeviceInfo(device_id.get(), info_flag, 0, nullptr, &size);
vector<char> buffer(size);
clGetDeviceInfo(device_id.get(), info_flag, sizeof(char) * size, buffer.data(),
nullptr);
return string(buffer.data());
}
device::device(device_ptr device_id, command_queue_ptr queue,
context_ptr context, unsigned id)
: device_id_(device_id),
command_queue_(queue),
context_(context),
id_(id) { }
} // namespace opencl
} // namespace caf
......@@ -32,6 +32,30 @@ cl_int clRetainDeviceDummy(cl_device_id) {
namespace caf {
namespace opencl {
std::ostream& operator<<(std::ostream& os, device_type dev) {
switch(dev) {
case def : os << "default"; break;
case cpu : os << "CPU"; break;
case gpu : os << "GPU"; break;
case accelerator : os << "accelerator"; break;
case custom : os << "custom"; break;
case all : os << "all"; break;
default : os.setstate(std::ios_base::failbit);
}
return os;
}
device_type device_type_from_ulong(cl_ulong dev) {
switch(dev) {
case CL_DEVICE_TYPE_CPU : return cpu;
case CL_DEVICE_TYPE_GPU : return gpu;
case CL_DEVICE_TYPE_ACCELERATOR : return accelerator;
case CL_DEVICE_TYPE_CUSTOM : return custom;
case CL_DEVICE_TYPE_ALL : return all;
default : return def;
}
}
std::string get_opencl_error(cl_int err) {
switch (err) {
case CL_SUCCESS:
......
......@@ -18,94 +18,61 @@
* http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/
#include "caf/opencl/device.hpp"
#include "caf/opencl/metainfo.hpp"
#include "caf/opencl/platform.hpp"
#include "caf/opencl/opencl_err.hpp"
#include "caf/opencl/opencl_metainfo.hpp"
namespace caf {
namespace opencl {
opencl_metainfo* opencl_metainfo::instance() {
metainfo* metainfo::instance() {
auto sid = detail::singletons::opencl_plugin_id;
auto fac = [] { return new opencl_metainfo; };
auto fac = [] { return new metainfo; };
auto res = detail::singletons::get_plugin_singleton(sid, fac);
return static_cast<opencl_metainfo*>(res);
return static_cast<metainfo*>(res);
}
const std::vector<device_info> opencl_metainfo::get_devices() const {
return devices_;
const std::vector<device>& metainfo::get_devices() const {
return platforms_.front().get_devices();
}
void opencl_metainfo::initialize() {
const optional<const device&> metainfo::get_device(size_t id) const{
if (platforms_.empty())
return none;
size_t from = 0;
size_t to = 0;
for (auto& pl : platforms_) {
from = to;
to += pl.get_devices().size();
if (id >= from && id < to)
return pl.get_devices()[id - from];
}
return none;
}
void metainfo::initialize() {
// get number of available platforms
auto num_platforms = v1get<cl_uint>(CAF_CLF(clGetPlatformIDs));
// get platform ids
std::vector<cl_platform_id> platforms(num_platforms);
v2callcl(CAF_CLF(clGetPlatformIDs), num_platforms, platforms.data());
if (platforms.empty()) {
std::vector<cl_platform_id> platform_ids(num_platforms);
v2callcl(CAF_CLF(clGetPlatformIDs), num_platforms, platform_ids.data());
if (platform_ids.empty())
throw std::runtime_error("no OpenCL platform found");
}
// support multiple platforms -> "for (auto platform : platforms)"?
auto platform = platforms.front();
// detect how many devices we got
cl_uint num_devs = 0;
cl_device_type dev_type = CL_DEVICE_TYPE_GPU;
// try get some GPU devices and try falling back to CPU devices on error
try {
num_devs = v1get<cl_uint>(CAF_CLF(clGetDeviceIDs), platform, dev_type);
}
catch (std::runtime_error&) {
dev_type = CL_DEVICE_TYPE_CPU;
num_devs = v1get<cl_uint>(CAF_CLF(clGetDeviceIDs), platform, dev_type);
}
// get available devices
std::vector<cl_device_id> ds(num_devs);
v2callcl(CAF_CLF(clGetDeviceIDs), platform, dev_type, num_devs, ds.data());
std::vector<device_ptr> devices(num_devs);
// lift raw pointer as returned by OpenCL to C++ smart pointers
auto lift = [](cl_device_id ptr) { return device_ptr{ptr, false}; };
std::transform(ds.begin(), ds.end(), devices.begin(), lift);
// create a context
context_.reset(v2get(CAF_CLF(clCreateContext), nullptr, num_devs, ds.data(),
pfn_notify, nullptr),
false);
for (auto& device : devices) {
CAF_LOG_DEBUG("creating command queue for device(s)");
command_queue_ptr cmd_queue;
try {
cmd_queue.reset(v2get(CAF_CLF(clCreateCommandQueue),
context_.get(), device.get(),
unsigned{CL_QUEUE_PROFILING_ENABLE}),
false);
}
catch (std::runtime_error&) {
CAF_LOG_DEBUG("unable to create command queue for device");
}
if (cmd_queue) {
auto max_wgs = v3get<size_t>(CAF_CLF(clGetDeviceInfo), device.get(),
unsigned{CL_DEVICE_MAX_WORK_GROUP_SIZE});
auto max_wid = v3get<cl_uint>(CAF_CLF(clGetDeviceInfo), device.get(),
unsigned{CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS});
dim_vec max_wi_per_dim(max_wid);
v2callcl(CAF_CLF(clGetDeviceInfo), device.get(),
unsigned{CL_DEVICE_MAX_WORK_ITEM_SIZES},
sizeof(size_t) * max_wid,
max_wi_per_dim.data());
devices_.push_back(device_info{std::move(device), std::move(cmd_queue),
max_wgs, max_wid, max_wi_per_dim});
}
}
if (devices_.empty()) {
std::string errstr = "could not create a command queue for any device";
CAF_LOG_ERROR(errstr);
throw std::runtime_error(std::move(errstr));
// initialize platforms (device discovery)
unsigned current_device_id = 0;
for (auto& id : platform_ids) {
platforms_.push_back(platform::create(id, current_device_id));
current_device_id +=
static_cast<unsigned>(platforms_.back().get_devices().size());
}
}
void opencl_metainfo::dispose() {
void metainfo::dispose() {
delete this;
}
void opencl_metainfo::stop() {
void metainfo::stop() {
// nop
}
......
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2015 *
* Dominik Charousset <dominik.charousset (at) haw-hamburg.de> *
* *
* Distributed under the terms and conditions of the BSD 3-Clause License or *
* (at your option) under the terms and conditions of the Boost Software *
* License 1.0. See accompanying files LICENSE and LICENSE_ALTERNATIVE. *
* *
* If you did not receive a copy of the license files, see *
* http://opensource.org/licenses/BSD-3-Clause and *
* http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/
#include <vector>
#include <iostream>
#include "caf/opencl/platform.hpp"
#include "caf/opencl/opencl_err.hpp"
using namespace std;
namespace caf {
namespace opencl {
platform platform::create(cl_platform_id platform_id,
unsigned start_id) {
vector<unsigned> device_types = {CL_DEVICE_TYPE_GPU,
CL_DEVICE_TYPE_ACCELERATOR,
CL_DEVICE_TYPE_CPU};
vector<cl_device_id> ids;
for (cl_device_type device_type : device_types) {
auto known = ids.size();
cl_uint discoverd;
auto err = clGetDeviceIDs(platform_id, device_type, 0, nullptr, &discoverd);
if (err == CL_DEVICE_NOT_FOUND) {
continue; // no devices of the type found
} else if (err != CL_SUCCESS) {
throwcl("clGetDeviceIDs", err);
}
ids.resize(known + discoverd);
v2callcl(CAF_CLF(clGetDeviceIDs), platform_id, device_type,
discoverd, (ids.data() + known));
}
vector<device_ptr> devices;
devices.resize(ids.size());
auto lift = [](cl_device_id ptr) { return device_ptr{ptr, false}; };
transform(ids.begin(), ids.end(), devices.begin(), lift);
context_ptr context;
context.reset(v2get(CAF_CLF(clCreateContext), nullptr,
static_cast<unsigned>(ids.size()),
ids.data(), pfn_notify, nullptr),
false);
vector<device> device_information;
for (auto& device_id : devices) {
device_information.push_back(device::create(context, device_id,
start_id++));
}
if (device_information.empty()) {
string errstr = "no devices for the platform found";
CAF_LOGF_ERROR(errstr);
throw runtime_error(move(errstr));
}
auto name = platform_info(platform_id, CL_PLATFORM_NAME);
auto vendor = platform_info(platform_id, CL_PLATFORM_VENDOR);
auto version = platform_info(platform_id, CL_PLATFORM_VERSION);
return platform{platform_id, move(context), move(name),
move(vendor), move(version),
move(device_information)};
}
string platform::platform_info(cl_platform_id platform_id,
unsigned info_flag) {
size_t size;
auto err = clGetPlatformInfo(platform_id, info_flag, 0, nullptr,
&size);
throwcl("clGetPlatformInfo", err);
vector<char> buffer(size);
v2callcl(CAF_CLF(clGetPlatformInfo), platform_id, info_flag,
sizeof(char) * size, buffer.data());
return string(buffer.data());
}
platform::platform(cl_platform_id platform_id, context_ptr context,
string name, string vendor, string version,
vector<device> devices)
: platform_id_(platform_id),
context_(context),
name_(move(name)),
vendor_(move(vendor)),
version_(move(version)),
devices_(move(devices)) { }
} // namespace opencl
} // namespace caf
......@@ -26,50 +26,49 @@
#include "caf/detail/singletons.hpp"
#include "caf/opencl/program.hpp"
#include "caf/opencl/metainfo.hpp"
#include "caf/opencl/opencl_err.hpp"
#include "caf/opencl/opencl_metainfo.hpp"
using namespace std;
namespace caf {
namespace opencl {
program::program(context_ptr context, command_queue_ptr queue,
program_ptr program,
map<string, kernel_ptr> available_kernels)
program::program(context_ptr context, command_queue_ptr queue, program_ptr prog)
: context_(move(context)),
program_(move(program)),
queue_(move(queue)),
available_kernels_(move(available_kernels)) {}
program_(move(prog)),
queue_(move(queue)) {}
program program::create(const char* kernel_source, const char* options,
uint32_t device_id) {
auto metainfo = opencl_metainfo::instance();
auto devices = metainfo->get_devices();
auto context = metainfo->context_;
if (devices.size() <= device_id) {
auto info = metainfo::instance();
auto dev = info->get_device(device_id);
if (! dev) {
ostringstream oss;
oss << "Device id " << device_id
<< " is not a vaild device. Maximum id is: " << (devices.size() - 1)
<< ".";
oss << "No device with id '" << device_id << "' found.";
CAF_LOGF_ERROR(oss.str());
throw runtime_error(oss.str());
}
return program::create(kernel_source, options, *dev);
}
program program::create(const char* kernel_source, const char* options,
const device& dev) {
// create program object from kernel source
size_t kernel_source_length = strlen(kernel_source);
program_ptr pptr;
auto rawptr = v2get(CAF_CLF(clCreateProgramWithSource),context.get(),
auto rawptr = v2get(CAF_CLF(clCreateProgramWithSource), dev.context_.get(),
cl_uint{1}, &kernel_source, &kernel_source_length);
pptr.reset(rawptr, false);
// build programm from program object
auto dev_tmp = devices[device_id].device_.get();
auto dev_tmp = dev.device_id_.get();
cl_int err = clBuildProgram(pptr.get(), 1, &dev_tmp,
options,nullptr, nullptr);
options, nullptr, nullptr);
if (err != CL_SUCCESS) {
ostringstream oss;
oss << "clBuildProgram: " << get_opencl_error(err);
// the build log will be printed by the
// pfn_notify (see opencl_metainfo.cpp)
// pfn_notify (see metainfo.cpp)
#ifndef __APPLE__
// seems that just apple implemented the
// pfn_notify callback, but we can get
......@@ -80,49 +79,18 @@ program program::create(const char* kernel_source, const char* options,
clGetProgramBuildInfo(pptr.get(), dev_tmp, CL_PROGRAM_BUILD_LOG,
sizeof(buildlog_buffer_size), nullptr,
&buildlog_buffer_size);
vector<char> buffer(buildlog_buffer_size);
// fill the buffer with buildlog informations
clGetProgramBuildInfo(pptr.get(), dev_tmp, CL_PROGRAM_BUILD_LOG,
sizeof(buffer[0]) * buildlog_buffer_size,
buffer.data(), nullptr);
CAF_LOGF_ERROR("Build log:\n" + string(buffer.data()) +
"\n########################################");
}
#endif
throw runtime_error(oss.str());
}
cl_uint number_of_kernels = 0;
clCreateKernelsInProgram(pptr.get(), 0, nullptr, &number_of_kernels);
vector<cl_kernel> kernels(number_of_kernels);
err = clCreateKernelsInProgram(pptr.get(), number_of_kernels, kernels.data(),
nullptr);
if (err != CL_SUCCESS) {
ostringstream oss;
oss << "clCreateKernelsInProgram: " << get_opencl_error(err);
throw runtime_error(oss.str());
}
map<string, kernel_ptr> available_kernels;
for (cl_uint i = 0; i < number_of_kernels; ++i) {
size_t ret_size;
clGetKernelInfo(kernels[i], CL_KERNEL_FUNCTION_NAME, 0, nullptr, &ret_size);
vector<char> name(ret_size);
err = clGetKernelInfo(kernels[i], CL_KERNEL_FUNCTION_NAME, ret_size,
(void*) name.data(), nullptr);
if (err != CL_SUCCESS) {
ostringstream oss;
oss << "clGetKernelInfo (CL_KERNEL_FUNCTION_NAME): "
<< get_opencl_error(err);
throw runtime_error(oss.str());
}
kernel_ptr kernel;
kernel.reset(move(kernels[i]));
available_kernels.emplace(string(name.data()), move(kernel));
}
return {context, devices[device_id].cmd_queue_, pptr,
move(available_kernels)};
return {dev.context_, dev.command_queue_, pptr};
}
} // namespace opencl
......
......@@ -10,6 +10,7 @@
#include "caf/all.hpp"
#include "caf/opencl/spawn_cl.hpp"
using namespace std;
using namespace caf;
using namespace caf::opencl;
......@@ -17,7 +18,7 @@ using detail::limited_vector;
namespace {
using ivec = std::vector<int>;
using ivec = vector<int>;
using dims = opencl::dim_vec;
constexpr size_t matrix_size = 4;
......@@ -153,7 +154,7 @@ public:
}
void data(ivec new_data) {
data_ = std::move(new_data);
data_ = move(new_data);
}
private:
......@@ -162,17 +163,17 @@ private:
template <class T>
std::vector<T> make_iota_vector(size_t num_elements) {
std::vector<T> result;
vector<T> make_iota_vector(size_t num_elements) {
vector<T> result;
result.resize(num_elements);
std::iota(result.begin(), result.end(), T{0});
iota(result.begin(), result.end(), T{0});
return result;
}
template <size_t Size>
square_matrix<Size> make_iota_matrix() {
square_matrix<Size> result;
std::iota(result.data().begin(), result.data().end(), 0);
iota(result.data().begin(), result.data().end(), 0);
return result;
}
......@@ -191,32 +192,37 @@ bool operator!=(const square_matrix<Size>& lhs,
using matrix_type = square_matrix<matrix_size>;
template <class T>
void check_vector_results(const std::string& description,
const std::vector<T>& expected,
const std::vector<T>& result) {
void check_vector_results(const string& description,
const vector<T>& expected,
const vector<T>& result) {
auto cond = (expected == result);
CAF_CHECK(cond);
if (!cond) {
if (! cond) {
CAF_TEST_INFO(description << " failed.");
std::cout << "Expected: " << std::endl;
cout << "Expected: " << endl;
for (size_t i = 0; i < expected.size(); ++i) {
std::cout << " " << expected[i];
cout << " " << expected[i];
}
std::cout << std::endl << "Received: " << std::endl;
cout << endl << "Received: " << endl;
for (size_t i = 0; i < result.size(); ++i) {
std::cout << " " << result[i];
cout << " " << result[i];
}
std::cout << std::endl;
cout << endl;
}
}
void test_opencl() {
auto info = metainfo::instance();
auto opt = info->get_device_if([](const device&){ return true; });
if (! opt)
CAF_TEST_ERROR("No OpenCL device found.");
auto dev = *opt;
scoped_actor self;
const ivec expected1{ 56, 62, 68, 74,
152, 174, 196, 218,
248, 286, 324, 362,
344, 398, 452, 506};
auto w1 = spawn_cl(program::create(kernel_source), kernel_name,
auto w1 = spawn_cl(program::create(kernel_source, "", dev), kernel_name,
opencl::spawn_config{dims{matrix_size, matrix_size}},
opencl::in<ivec>{}, opencl::out<ivec>{});
self->send(w1, make_iota_vector<int>(matrix_size * matrix_size));
......@@ -245,16 +251,16 @@ void test_opencl() {
<< to_string(self->current_message()));
}
);
const matrix_type expected2(std::move(expected1));
const matrix_type expected2(move(expected1));
auto map_arg = [](message& msg) -> optional<message> {
return msg.apply(
[](matrix_type& mx) {
return make_message(std::move(mx.data()));
return make_message(move(mx.data()));
}
);
};
auto map_res = [=](ivec result) -> message {
return make_message(matrix_type{std::move(result)});
return make_message(matrix_type{move(result)});
};
opencl::spawn_config cfg3{dims{matrix_size, matrix_size}};
auto w3 = spawn_cl(program::create(kernel_source), kernel_name, cfg3,
......@@ -292,11 +298,11 @@ void test_opencl() {
try {
auto create_error = program::create(kernel_source_error);
}
catch (const std::exception& exc) {
catch (const exception& exc) {
auto cond = (strcmp("clBuildProgram: CL_BUILD_PROGRAM_FAILURE",
exc.what()) == 0);
CAF_CHECK(cond);
if (!cond) {
if (! cond) {
CAF_TEST_INFO("Wrong exception cought for program build failure.");
}
}
......@@ -317,45 +323,44 @@ void test_opencl() {
}
);
auto get_max_workgroup_size = [](size_t dev_id, size_t dims) -> size_t {
size_t max_size = 512;
auto devices = opencl_metainfo::instance()->get_devices()[dev_id];
size_t dimsize = devices.get_max_work_items_per_dim()[dims];
return max_size < dimsize ? max_size : dimsize;
};
// test for manuel return size selection (max workgroup size 1d)
const int max_workgroup_size = static_cast<int>(get_max_workgroup_size(0,1));
const size_t reduce_buffer_size = static_cast<size_t>(max_workgroup_size) * 8;
const size_t reduce_local_size = static_cast<size_t>(max_workgroup_size);
const size_t reduce_work_groups = reduce_buffer_size / reduce_local_size;
const size_t reduce_global_size = reduce_buffer_size;
const size_t reduce_result_size = reduce_work_groups;
ivec arr6(reduce_buffer_size);
int n = static_cast<int>(arr6.capacity());
std::generate(arr6.begin(), arr6.end(), [&]{ return --n; });
opencl::spawn_config cfg6{dims{reduce_global_size},
dims{ /* no offsets */ },
dims{reduce_local_size}};
auto get_result_size_6 = [=](const ivec&) {
return reduce_result_size;
};
auto w6 = spawn_cl(kernel_source_reduce, kernel_name_reduce, cfg6,
opencl::in<ivec>{}, opencl::out<ivec>{get_result_size_6});
self->send(w6, move(arr6));
ivec expected4{max_workgroup_size * 7, max_workgroup_size * 6,
max_workgroup_size * 5, max_workgroup_size * 4,
max_workgroup_size * 3, max_workgroup_size * 2,
max_workgroup_size, 0};
self->receive(
[&](const ivec& result) {
check_vector_results("Passing size for the output", expected4, result);
},
others >> [&] {
CAF_TEST_ERROR("Unexpected message "
<< to_string(self->current_message()));
}
);
auto dev6 = info->get_device_if([](const device& d) {
return d.get_device_type() != cpu;
});
if (dev6) {
// test for manuel return size selection (max workgroup size 1d)
const size_t max_wg_size = min(dev6->get_max_work_item_sizes()[0], 512UL);
const size_t reduce_buffer_size = static_cast<size_t>(max_wg_size) * 8;
const size_t reduce_local_size = static_cast<size_t>(max_wg_size);
const size_t reduce_work_groups = reduce_buffer_size / reduce_local_size;
const size_t reduce_global_size = reduce_buffer_size;
const size_t reduce_result_size = reduce_work_groups;
ivec arr6(reduce_buffer_size);
int n = static_cast<int>(arr6.capacity());
generate(arr6.begin(), arr6.end(), [&]{ return --n; });
opencl::spawn_config cfg6{dims{reduce_global_size},
dims{ /* no offsets */ },
dims{reduce_local_size}};
auto get_result_size_6 = [reduce_result_size](const ivec&) {
return reduce_result_size;
};
auto w6 = spawn_cl(program::create(kernel_source_reduce, "", *dev6),
kernel_name_reduce, cfg6,
opencl::in<ivec>{}, opencl::out<ivec>{get_result_size_6});
self->send(w6, move(arr6));
auto wg_size_as_int = static_cast<int>(max_wg_size);
ivec expected4{wg_size_as_int * 7, wg_size_as_int * 6, wg_size_as_int * 5,
wg_size_as_int * 4, wg_size_as_int * 3, wg_size_as_int * 2,
wg_size_as_int , 0};
self->receive(
[&](const ivec& result) {
check_vector_results("Passing size for the output", expected4, result);
},
others >> [&] {
CAF_TEST_ERROR("Unexpected message "
<< to_string(self->current_message()));
}
);
}
// calculator function for getting the size of the output
auto get_result_size_7 = [=](const ivec&) {
return problem_size;
......@@ -381,12 +386,12 @@ void test_opencl() {
// test in_out argument type
ivec input9 = make_iota_vector<int>(problem_size);
ivec expected9{input9};
std::transform(begin(expected9), end(expected9), begin(expected9),
transform(begin(expected9), end(expected9), begin(expected9),
[](const int& val){ return val * 2; });
auto w9 = spawn_cl(kernel_source_inout, kernel_name_inout,
spawn_config{dims{problem_size}},
opencl::in_out<ivec>{});
self->send(w9, std::move(input9));
self->send(w9, move(input9));
self->receive(
[&](const ivec& result) {
check_vector_results("Testing in_out arugment", expected9, result);
......
......@@ -28,7 +28,6 @@ constexpr int magic_number = 23;
constexpr const char* kernel_name = "matrix_square";
constexpr const char* kernel_name_compiler_flag = "compiler_flag";
constexpr const char* kernel_name_reduce = "reduce";
constexpr const char* kernel_name_const = "const_mod";
constexpr const char* compiler_flag = "-D CAF_OPENCL_TEST_FLAG";
......@@ -65,29 +64,6 @@ constexpr const char* kernel_source_compiler_flag = R"__(
}
)__";
// http://developer.amd.com/resources/documentation-articles/articles-whitepapers/
// opencl-optimization-case-study-simple-reductions
constexpr const char* kernel_source_reduce = R"__(
__kernel void reduce(__global int* buffer,
__global int* result) {
__local int scratch[512];
int local_index = get_local_id(0);
scratch[local_index] = buffer[get_global_id(0)];
barrier(CLK_LOCAL_MEM_FENCE);
for(int offset = get_local_size(0) / 2; offset > 0; offset = offset / 2) {
if (local_index < offset) {
int other = scratch[local_index + offset];
int mine = scratch[local_index];
scratch[local_index] = (mine < other) ? mine : other;
}
barrier(CLK_LOCAL_MEM_FENCE);
}
if (local_index == 0) {
result[get_group_id(0)] = scratch[0];
}
}
)__";
constexpr const char* kernel_source_const = R"__(
__kernel void const_mod(__constant int* input,
__global int* output) {
......@@ -190,7 +166,7 @@ void check_vector_results(const std::string& description,
const std::vector<T>& result) {
auto cond = (expected == result);
CAF_CHECK(cond);
if (!cond) {
if (! cond) {
CAF_TEST_INFO(description << " failed.");
std::cout << "Expected: " << std::endl;
for (size_t i = 0; i < expected.size(); ++i) {
......@@ -239,7 +215,7 @@ void test_opencl_deprecated() {
}
);
};
auto map_res = [](ivec result) -> message {
auto map_res = [](ivec& result) -> message {
return make_message(matrix_type{std::move(result)});
};
auto w3 = spawn_cl<ivec (ivec)>(program::create(kernel_source),
......@@ -274,7 +250,7 @@ void test_opencl_deprecated() {
auto cond = (strcmp("clBuildProgram: CL_BUILD_PROGRAM_FAILURE",
exc.what()) == 0);
CAF_CHECK(cond);
if (!cond) {
if (! cond) {
CAF_TEST_INFO("Wrong exception cought for program build failure.");
}
}
......@@ -289,39 +265,6 @@ void test_opencl_deprecated() {
check_vector_results("Passing compiler flags", expected3, result);
}
);
auto get_max_workgroup_size = [](size_t dev_id, size_t dims) -> size_t {
size_t max_size = 512;
auto devices = opencl_metainfo::instance()->get_devices()[dev_id];
size_t dimsize = devices.get_max_work_items_per_dim()[dims];
return max_size < dimsize ? max_size : dimsize;
};
// test for manuel return size selection (max workgroup size 1d)
const int max_workgroup_size = static_cast<int>(get_max_workgroup_size(0,1));
const int reduce_buffer_size = max_workgroup_size * 8;
const int reduce_local_size = max_workgroup_size;
const int reduce_work_groups = reduce_buffer_size / reduce_local_size;
const int reduce_global_size = reduce_buffer_size;
const int reduce_result_size = reduce_work_groups;
ivec arr6(static_cast<size_t>(reduce_buffer_size));
int n = static_cast<int>(arr6.capacity());
std::generate(arr6.begin(), arr6.end(), [&]{ return --n; });
auto w6 = spawn_cl<ivec (ivec)>(kernel_source_reduce, kernel_name_reduce,
limited_vector<size_t, 3>{static_cast<size_t>(reduce_global_size)},
{},
limited_vector<size_t, 3>{static_cast<size_t>(reduce_local_size)},
static_cast<size_t>(reduce_result_size));
self->send(w6, move(arr6));
ivec expected4{max_workgroup_size * 7, max_workgroup_size * 6,
max_workgroup_size * 5, max_workgroup_size * 4,
max_workgroup_size * 3, max_workgroup_size * 2,
max_workgroup_size, 0};
self->receive(
[&](const ivec& result) {
check_vector_results("Passing size for the output", expected4, result);
}
);
// constant memory arguments
const ivec arr7{magic_number};
auto w7 = spawn_cl<ivec (ivec)>(kernel_source_const, kernel_name_const,
......
Markdown is supported
0%
or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment