/// Enqueues a command to migrate \p mem_objects.
///
/// \see_opencl_ref{clEnqueueMigrateMemObjects}
///
/// \opencl_version_warning{1,2}
event enqueue_migrate_memory_objects(uint_ num_mem_objects,
const cl_mem *mem_objects,
cl_mem_migration_flags flags,
const wait_list &events = wait_list())
{
BOOST_ASSERT(m_queue != 0);
event event_;
cl_int ret = clEnqueueMigrateMemObjects(
m_queue,
num_mem_objects,
mem_objects,
flags,
events.size(),
events.get_event_ptr(),
&event_.get()
);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
return event_;
}
/// Enqueues a kernel for execution.
///
/// \see_opencl_ref{clEnqueueNDRangeKernel}
event enqueue_nd_range_kernel(const kernel &kernel,
size_t work_dim,
const size_t *global_work_offset,
const size_t *global_work_size,
const size_t *local_work_size,
const wait_list &events = wait_list())
{
BOOST_ASSERT(m_queue != 0);
BOOST_ASSERT(kernel.get_context() == this->get_context());
event event_;
cl_int ret = clEnqueueNDRangeKernel(
m_queue,
kernel,
static_cast<cl_uint>(work_dim),
global_work_offset,
global_work_size,
local_work_size,
events.size(),
events.get_event_ptr(),
&event_.get()
);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
return event_;
}
/// Enqueues a command to fill \p image with \p fill_color.
///
/// \see_opencl_ref{clEnqueueFillImage}
///
/// \opencl_version_warning{1,2}
event enqueue_fill_image(const image3d &image,
const void *fill_color,
const size_t origin[3],
const size_t region[3],
const wait_list &events = wait_list())
{
BOOST_ASSERT(m_queue != 0);
BOOST_ASSERT(image.get_context() == this->get_context());
event event_;
cl_int ret = clEnqueueFillImage(
m_queue,
image.get(),
fill_color,
origin,
region,
events.size(),
events.get_event_ptr(),
&event_.get()
);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
return event_;
}
static VALUE
legacy_wait(VALUE unused, legacy_wait_args *args)
{
wait_list(&args->condvar->waiting);
rb_funcall(args->mutex, rb_intern("lock"), 0);
return Qnil;
}
static VALUE
lock_mutex(Mutex *mutex)
{
VALUE current;
current = rb_thread_current();
rb_thread_critical = 1;
if (!MUTEX_LOCKED_P(mutex)) {
mutex->owner = current;
}
else {
do {
wait_list(&mutex->waiting);
rb_thread_critical = 1;
if (!MUTEX_LOCKED_P(mutex)) {
mutex->owner = current;
break;
}
} while (mutex->owner != current);
}
rb_thread_critical = 0;
return Qnil;
}
/// Enqueues a command to fill \p buffer with \p pattern.
///
/// \see_opencl_ref{clEnqueueFillBuffer}
///
/// \opencl_version_warning{1,2}
///
/// \see fill()
event enqueue_fill_buffer(const buffer &buffer,
const void *pattern,
size_t pattern_size,
size_t offset,
size_t size,
const wait_list &events = wait_list())
{
BOOST_ASSERT(m_queue != 0);
BOOST_ASSERT(offset + size <= buffer.size());
BOOST_ASSERT(buffer.get_context() == this->get_context());
event event_;
cl_int ret = clEnqueueFillBuffer(
m_queue,
buffer.get(),
pattern,
pattern_size,
offset,
size,
events.size(),
events.get_event_ptr(),
&event_.get()
);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
return event_;
}
/// Enqueues a command to map \p buffer into the host address space.
///
/// \see_opencl_ref{clEnqueueMapBuffer}
void* enqueue_map_buffer(const buffer &buffer,
cl_map_flags flags,
size_t offset,
size_t size,
const wait_list &events = wait_list())
{
BOOST_ASSERT(m_queue != 0);
BOOST_ASSERT(offset + size <= buffer.size());
BOOST_ASSERT(buffer.get_context() == this->get_context());
cl_int ret = 0;
void *pointer = clEnqueueMapBuffer(
m_queue,
buffer.get(),
CL_TRUE,
flags,
offset,
size,
events.size(),
events.get_event_ptr(),
0,
&ret
);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
return pointer;
}
/// Enqueues a command to write data from host memory to \p buffer.
/// The copy is performed asynchronously.
///
/// \see_opencl_ref{clEnqueueWriteBuffer}
///
/// \see copy_async()
event enqueue_write_buffer_async(const buffer &buffer,
size_t offset,
size_t size,
const void *host_ptr,
const wait_list &events = wait_list())
{
BOOST_ASSERT(m_queue != 0);
BOOST_ASSERT(size <= buffer.size());
BOOST_ASSERT(buffer.get_context() == this->get_context());
BOOST_ASSERT(host_ptr != 0);
event event_;
cl_int ret = clEnqueueWriteBuffer(
m_queue,
buffer.get(),
CL_FALSE,
offset,
size,
host_ptr,
events.size(),
events.get_event_ptr(),
&event_.get()
);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
return event_;
}
/// Enqueues a kernel to execute using a single work-item.
///
/// \see_opencl_ref{clEnqueueTask}
event enqueue_task(const kernel &kernel, const wait_list &events = wait_list())
{
BOOST_ASSERT(m_queue != 0);
BOOST_ASSERT(kernel.get_context() == this->get_context());
event event_;
// clEnqueueTask() was deprecated in OpenCL 2.0. In that case we
// just forward to the equivalent clEnqueueNDRangeKernel() call.
#ifdef CL_VERSION_2_0
size_t one = 1;
cl_int ret = clEnqueueNDRangeKernel(
m_queue, kernel, 1, 0, &one, &one,
events.size(), events.get_event_ptr(), &event_.get()
);
#else
cl_int ret = clEnqueueTask(
m_queue, kernel, events.size(), events.get_event_ptr(), &event_.get()
);
#endif
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
return event_;
}
/// Enqueues a command to copy data from \p src_buffer to \p dst_image.
///
/// \see_opencl_ref{clEnqueueCopyBufferToImage}
event enqueue_copy_buffer_to_image(const buffer &src_buffer,
const image3d &dst_image,
size_t src_offset,
const size_t dst_origin[3],
const size_t region[3],
const wait_list &events = wait_list())
{
BOOST_ASSERT(m_queue != 0);
BOOST_ASSERT(src_buffer.get_context() == this->get_context());
BOOST_ASSERT(dst_image.get_context() == this->get_context());
event event_;
cl_int ret = clEnqueueCopyBufferToImage(
m_queue,
src_buffer.get(),
dst_image.get(),
src_offset,
dst_origin,
region,
events.size(),
events.get_event_ptr(),
&event_.get()
);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
return event_;
}
/// Enqueues a command to write data from host memory to \p image.
///
/// \see_opencl_ref{clEnqueueWriteImage}
void enqueue_write_image(const image3d &image,
const size_t origin[3],
const size_t region[3],
size_t input_row_pitch,
size_t input_slice_pitch,
const void *host_ptr,
const wait_list &events = wait_list())
{
BOOST_ASSERT(m_queue != 0);
BOOST_ASSERT(image.get_context() == this->get_context());
cl_int ret = clEnqueueWriteImage(
m_queue,
image.get(),
CL_TRUE,
origin,
region,
input_row_pitch,
input_slice_pitch,
host_ptr,
events.size(),
events.get_event_ptr(),
0
);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
}
/// Enqueues a command to read data from \p buffer to host memory.
///
/// \see_opencl_ref{clEnqueueReadBuffer}
///
/// \see copy()
void enqueue_read_buffer(const buffer &buffer,
size_t offset,
size_t size,
void *host_ptr,
const wait_list &events = wait_list())
{
BOOST_ASSERT(m_queue != 0);
BOOST_ASSERT(size <= buffer.size());
BOOST_ASSERT(buffer.get_context() == this->get_context());
BOOST_ASSERT(host_ptr != 0);
cl_int ret = clEnqueueReadBuffer(
m_queue,
buffer.get(),
CL_TRUE,
offset,
size,
host_ptr,
events.size(),
events.get_event_ptr(),
0
);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
}
/// Enqueues a command to fill \p size bytes of data at \p svm_ptr with
/// \p pattern.
///
/// \opencl_version_warning{2,0}
///
/// \see_opencl2_ref{clEnqueueSVMMemFill}
event enqueue_svm_fill(void *svm_ptr,
const void *pattern,
size_t pattern_size,
size_t size,
const wait_list &events = wait_list())
{
event event_;
cl_int ret = clEnqueueSVMMemFill(
m_queue,
svm_ptr,
pattern,
pattern_size,
size,
events.size(),
events.get_event_ptr(),
&event_.get()
);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
return event_;
}
/// Enqueues a command to copy data from \p src_image to \p dst_image.
///
/// \see_opencl_ref{clEnqueueCopyImage}
event enqueue_copy_image(const image3d &src_image,
const image3d &dst_image,
const size_t src_origin[3],
const size_t dst_origin[3],
const size_t region[3],
const wait_list &events = wait_list())
{
BOOST_ASSERT(m_queue != 0);
BOOST_ASSERT(src_image.get_context() == this->get_context());
BOOST_ASSERT(dst_image.get_context() == this->get_context());
BOOST_ASSERT_MSG(src_image.get_format() == dst_image.get_format(),
"Source and destination image formats must match.");
event event_;
cl_int ret = clEnqueueCopyImage(
m_queue,
src_image.get(),
dst_image.get(),
src_origin,
dst_origin,
region,
events.size(),
events.get_event_ptr(),
&event_.get()
);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
return event_;
}
/// Enqueues a command to copy data from \p src_image to \p dst_buffer.
///
/// \see_opencl_ref{clEnqueueCopyImageToBuffer}
event enqueue_copy_image_to_buffer(const image2d &src_image,
const buffer &dst_buffer,
const size_t src_origin[2],
const size_t region[2],
size_t dst_offset,
const wait_list &events = wait_list())
{
BOOST_ASSERT(m_queue != 0);
BOOST_ASSERT(src_image.get_context() == this->get_context());
BOOST_ASSERT(dst_buffer.get_context() == this->get_context());
const size_t src_origin3[3] = { src_origin[0], src_origin[1], size_t(0) };
const size_t region3[3] = { region[0], region[1], size_t(1) };
event event_;
cl_int ret = clEnqueueCopyImageToBuffer(
m_queue,
src_image.get(),
dst_buffer.get(),
src_origin3,
region3,
dst_offset,
events.size(),
events.get_event_ptr(),
&event_.get()
);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
return event_;
}
/// Enqueues a function to execute on the host.
event enqueue_native_kernel(void (BOOST_COMPUTE_CL_CALLBACK *user_func)(void *),
void *args,
size_t cb_args,
uint_ num_mem_objects,
const cl_mem *mem_list,
const void **args_mem_loc,
const wait_list &events = wait_list())
{
BOOST_ASSERT(m_queue != 0);
event event_;
cl_int ret = clEnqueueNativeKernel(
m_queue,
user_func,
args,
cb_args,
num_mem_objects,
mem_list,
args_mem_loc,
events.size(),
events.get_event_ptr(),
&event_.get()
);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
return event_;
}
/// Enqueues a command to read data from \p image to host memory.
///
/// \see_opencl_ref{clEnqueueReadImage}
void enqueue_read_image(const image2d &image,
const size_t origin[2],
const size_t region[2],
size_t row_pitch,
void *host_ptr,
const wait_list &events = wait_list())
{
BOOST_ASSERT(m_queue != 0);
BOOST_ASSERT(image.get_context() == this->get_context());
const size_t origin3[3] = { origin[0], origin[1], size_t(0) };
const size_t region3[3] = { region[0], region[1], size_t(1) };
cl_int ret = clEnqueueReadImage(
m_queue,
image.get(),
CL_TRUE,
origin3,
region3,
row_pitch,
0,
host_ptr,
events.size(),
events.get_event_ptr(),
0
);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
}
/// Enqueues a command to copy data from \p src_buffer to
/// \p dst_buffer.
///
/// \see_opencl_ref{clEnqueueCopyBuffer}
///
/// \see copy()
event enqueue_copy_buffer(const buffer &src_buffer,
const buffer &dst_buffer,
size_t src_offset,
size_t dst_offset,
size_t size,
const wait_list &events = wait_list())
{
BOOST_ASSERT(m_queue != 0);
BOOST_ASSERT(src_offset + size <= src_buffer.size());
BOOST_ASSERT(dst_offset + size <= dst_buffer.size());
BOOST_ASSERT(src_buffer.get_context() == this->get_context());
BOOST_ASSERT(dst_buffer.get_context() == this->get_context());
event event_;
cl_int ret = clEnqueueCopyBuffer(
m_queue,
src_buffer.get(),
dst_buffer.get(),
src_offset,
dst_offset,
size,
events.size(),
events.get_event_ptr(),
&event_.get()
);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
return event_;
}
/// Enqueues a command to unmap \p buffer from the host memory space.
///
/// \see_opencl_ref{clEnqueueUnmapMemObject}
event enqueue_unmap_buffer(const buffer &buffer,
void *mapped_ptr,
const wait_list &events = wait_list())
{
BOOST_ASSERT(buffer.get_context() == this->get_context());
return enqueue_unmap_mem_object(buffer.get(), mapped_ptr, events);
}
/// Enqueues a command to release the specified OpenGL buffer.
///
/// \see_opencl_ref{clEnqueueReleaseGLObjects}
inline event opengl_enqueue_release_buffer(const opengl_buffer &buffer,
command_queue &queue,
const wait_list &events = wait_list())
{
BOOST_ASSERT(buffer.get_context() == queue.get_context());
return opengl_enqueue_release_gl_objects(1, &buffer.get(), queue, events);
}
/// Convenience overload for enqueue_native_kernel() which enqueues a
/// native kernel on the host with a nullary function.
event enqueue_native_kernel(void (BOOST_COMPUTE_CL_CALLBACK *user_func)(void),
const wait_list &events = wait_list())
{
return enqueue_native_kernel(
detail::nullary_native_kernel_trampoline,
reinterpret_cast<void *>(&user_func),
sizeof(user_func),
0,
0,
0,
events
);
}
/// Convenience method which calls enqueue_nd_range_kernel() with a
/// one-dimensional range.
event enqueue_1d_range_kernel(const kernel &kernel,
size_t global_work_offset,
size_t global_work_size,
size_t local_work_size,
const wait_list &events = wait_list())
{
return enqueue_nd_range_kernel(
kernel,
1,
&global_work_offset,
&global_work_size,
local_work_size ? &local_work_size : 0,
events
);
}
static VALUE
lock_mutex(Mutex *mutex)
{
VALUE current;
current = rb_thread_current();
//rb_thread_critical = 1;
while (RTEST(mutex->owner)) {
wait_list(&mutex->waiting);
//rb_thread_critical = 1;
}
mutex->owner = current;
//rb_thread_critical = 0;
return Qnil;
}
/// Enqueues a command to unmap \p svm_ptr from the host memory space.
///
/// \opencl_version_warning{2,0}
///
/// \see_opencl2_ref{clEnqueueSVMUnmap}
event enqueue_svm_unmap(void *svm_ptr,
const wait_list &events = wait_list())
{
event event_;
cl_int ret = clEnqueueSVMUnmap(
m_queue,
svm_ptr,
events.size(),
events.get_event_ptr(),
&event_.get()
);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
return event_;
}
/// Enqueues a command to copy \p size bytes of data from \p src_ptr to
/// \p dst_ptr.
///
/// \opencl_version_warning{2,0}
///
/// \see_opencl2_ref{clEnqueueSVMMemcpy}
void enqueue_svm_memcpy(void *dst_ptr,
const void *src_ptr,
size_t size,
const wait_list &events = wait_list())
{
cl_int ret = clEnqueueSVMMemcpy(
m_queue,
CL_TRUE,
dst_ptr,
src_ptr,
size,
events.size(),
events.get_event_ptr(),
0
);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
}
/// Enqueues a command to map \p svm_ptr to the host memory space.
///
/// \opencl_version_warning{2,0}
///
/// \see_opencl2_ref{clEnqueueSVMMap}
void enqueue_svm_map(void *svm_ptr,
size_t size,
cl_map_flags flags,
const wait_list &events = wait_list())
{
cl_int ret = clEnqueueSVMMap(
m_queue,
CL_TRUE,
flags,
svm_ptr,
size,
events.size(),
events.get_event_ptr(),
0
);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
}
/// Enqueues a command to release the specified OpenGL memory objects.
///
/// \see_opencl_ref{clEnqueueReleaseGLObjects}
inline event opengl_enqueue_release_gl_objects(size_t num_objects,
const cl_mem *mem_objects,
command_queue &queue,
const wait_list &events = wait_list())
{
BOOST_ASSERT(queue != 0);
event event_;
cl_int ret = clEnqueueReleaseGLObjects(queue.get(),
num_objects,
mem_objects,
events.size(),
events.get_event_ptr(),
&event_.get());
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
return event_;
}
/// Enqueues a command to copy a rectangular region from
/// \p src_buffer to \p dst_buffer.
///
/// \see_opencl_ref{clEnqueueCopyBufferRect}
///
/// \opencl_version_warning{1,1}
event enqueue_copy_buffer_rect(const buffer &src_buffer,
const buffer &dst_buffer,
const size_t src_origin[3],
const size_t dst_origin[3],
const size_t region[3],
size_t buffer_row_pitch,
size_t buffer_slice_pitch,
size_t host_row_pitch,
size_t host_slice_pitch,
const wait_list &events = wait_list())
{
BOOST_ASSERT(m_queue != 0);
BOOST_ASSERT(src_buffer.get_context() == this->get_context());
BOOST_ASSERT(dst_buffer.get_context() == this->get_context());
event event_;
cl_int ret = clEnqueueCopyBufferRect(
m_queue,
src_buffer.get(),
dst_buffer.get(),
src_origin,
dst_origin,
region,
buffer_row_pitch,
buffer_slice_pitch,
host_row_pitch,
host_slice_pitch,
events.size(),
events.get_event_ptr(),
&event_.get()
);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
return event_;
}
/// Enqueues a command to unmap \p mem from the host memory space.
///
/// \see_opencl_ref{clEnqueueUnmapMemObject}
event enqueue_unmap_mem_object(cl_mem mem,
void *mapped_ptr,
const wait_list &events = wait_list())
{
BOOST_ASSERT(m_queue != 0);
event event_;
cl_int ret = clEnqueueUnmapMemObject(
m_queue,
mem,
mapped_ptr,
events.size(),
events.get_event_ptr(),
&event_.get()
);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
return event_;
}
/// Enqueues a command to write a rectangular region from host memory
/// to \p buffer.
///
/// \see_opencl_ref{clEnqueueWriteBufferRect}
///
/// \opencl_version_warning{1,1}
void enqueue_write_buffer_rect(const buffer &buffer,
const size_t buffer_origin[3],
const size_t host_origin[3],
const size_t region[3],
size_t buffer_row_pitch,
size_t buffer_slice_pitch,
size_t host_row_pitch,
size_t host_slice_pitch,
void *host_ptr,
const wait_list &events = wait_list())
{
BOOST_ASSERT(m_queue != 0);
BOOST_ASSERT(buffer.get_context() == this->get_context());
BOOST_ASSERT(host_ptr != 0);
cl_int ret = clEnqueueWriteBufferRect(
m_queue,
buffer.get(),
CL_TRUE,
buffer_origin,
host_origin,
region,
buffer_row_pitch,
buffer_slice_pitch,
host_row_pitch,
host_slice_pitch,
host_ptr,
events.size(),
events.get_event_ptr(),
0
);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
}
