diff --git a/CMakeLists.txt b/CMakeLists.txt index 7657d08a..231c0b85 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -12,6 +12,12 @@ include_directories(${OPENCL_INCLUDE_DIRS}) find_package(Boost 1.48 REQUIRED) include_directories(${Boost_INCLUDE_DIRS}) +# optional third-party libraries +option(BOOST_COMPUTE_HAVE_EIGEN "Have Eigen" OFF) +option(BOOST_COMPUTE_HAVE_OPENCV "Have OpenCV" OFF) +option(BOOST_COMPUTE_HAVE_QT "Have Qt" OFF) +option(BOOST_COMPUTE_HAVE_VTK "Have VTK" OFF) + include_directories(include) if(${OpenCL_HEADER_CL_EXT_FOUND}) diff --git a/cmake/FindEigen.cmake b/cmake/FindEigen.cmake new file mode 100644 index 00000000..2cd3e12c --- /dev/null +++ b/cmake/FindEigen.cmake @@ -0,0 +1,160 @@ +# Ceres Solver - A fast non-linear least squares minimizer +# Copyright 2013 Google Inc. All rights reserved. +# http://code.google.com/p/ceres-solver/ +# +# Redistribution and use in source and binary forms, with or without +# modification, are permitted provided that the following conditions are met: +# +# * Redistributions of source code must retain the above copyright notice, +# this list of conditions and the following disclaimer. +# * Redistributions in binary form must reproduce the above copyright notice, +# this list of conditions and the following disclaimer in the documentation +# and/or other materials provided with the distribution. +# * Neither the name of Google Inc. nor the names of its contributors may be +# used to endorse or promote products derived from this software without +# specific prior written permission. +# +# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE +# ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE +# LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR +# CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF +# SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS +# INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN +# CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) +# ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE +# POSSIBILITY OF SUCH DAMAGE. +# +# Author: alexs.mac@gmail.com (Alex Stewart) +# + +# FindEigen.cmake - Find Eigen library, version >= 3. +# +# This module defines the following variables: +# +# EIGEN_FOUND: TRUE iff Eigen is found. +# EIGEN_INCLUDE_DIRS: Include directories for Eigen. +# +# EIGEN_VERSION: Extracted from Eigen/src/Core/util/Macros.h +# EIGEN_WORLD_VERSION: Equal to 3 if EIGEN_VERSION = 3.2.0 +# EIGEN_MAJOR_VERSION: Equal to 2 if EIGEN_VERSION = 3.2.0 +# EIGEN_MINOR_VERSION: Equal to 0 if EIGEN_VERSION = 3.2.0 +# +# The following variables control the behaviour of this module: +# +# EIGEN_INCLUDE_DIR_HINTS: List of additional directories in which to +# search for eigen includes, e.g: /timbuktu/eigen3. +# +# The following variables are also defined by this module, but in line with +# CMake recommended FindPackage() module style should NOT be referenced directly +# by callers (use the plural variables detailed above instead). These variables +# do however affect the behaviour of the module via FIND_[PATH/LIBRARY]() which +# are NOT re-called (i.e. search for library is not repeated) if these variables +# are set with valid values _in the CMake cache_. This means that if these +# variables are set directly in the cache, either by the user in the CMake GUI, +# or by the user passing -DVAR=VALUE directives to CMake when called (which +# explicitly defines a cache variable), then they will be used verbatim, +# bypassing the HINTS variables and other hard-coded search locations. +# +# EIGEN_INCLUDE_DIR: Include directory for CXSparse, not including the +# include directory of any dependencies. + +# Called if we failed to find Eigen or any of it's required dependencies, +# unsets all public (designed to be used externally) variables and reports +# error message at priority depending upon [REQUIRED/QUIET/] argument. +MACRO(EIGEN_REPORT_NOT_FOUND REASON_MSG) + UNSET(EIGEN_FOUND) + UNSET(EIGEN_INCLUDE_DIRS) + # Make results of search visible in the CMake GUI if Eigen has not + # been found so that user does not have to toggle to advanced view. + MARK_AS_ADVANCED(CLEAR EIGEN_INCLUDE_DIR) + # Note _FIND_[REQUIRED/QUIETLY] variables defined by FindPackage() + # use the camelcase library name, not uppercase. + IF (Eigen_FIND_QUIETLY) + MESSAGE(STATUS "Failed to find Eigen - " ${REASON_MSG} ${ARGN}) + ELSEIF (Eigen_FIND_REQUIRED) + MESSAGE(FATAL_ERROR "Failed to find Eigen - " ${REASON_MSG} ${ARGN}) + ELSE() + # Neither QUIETLY nor REQUIRED, use no priority which emits a message + # but continues configuration and allows generation. + MESSAGE("-- Failed to find Eigen - " ${REASON_MSG} ${ARGN}) + ENDIF () +ENDMACRO(EIGEN_REPORT_NOT_FOUND) + +# Search user-installed locations first, so that we prefer user installs +# to system installs where both exist. +# +# TODO: Add standard Windows search locations for Eigen. +LIST(APPEND EIGEN_CHECK_INCLUDE_DIRS + /usr/local/include/eigen3 + /usr/local/homebrew/include/eigen3 # Mac OS X + /opt/local/var/macports/software/eigen3 # Mac OS X. + /opt/local/include/eigen3 + /usr/include/eigen3) + +# Search supplied hint directories first if supplied. +FIND_PATH(EIGEN_INCLUDE_DIR + NAMES Eigen/Core + PATHS ${EIGEN_INCLUDE_DIR_HINTS} + ${EIGEN_CHECK_INCLUDE_DIRS}) +IF (NOT EIGEN_INCLUDE_DIR OR + NOT EXISTS ${EIGEN_INCLUDE_DIR}) + EIGEN_REPORT_NOT_FOUND( + "Could not find eigen3 include directory, set EIGEN_INCLUDE_DIR to " + "path to eigen3 include directory, e.g. /usr/local/include/eigen3.") +ENDIF (NOT EIGEN_INCLUDE_DIR OR + NOT EXISTS ${EIGEN_INCLUDE_DIR}) + +# Mark internally as found, then verify. EIGEN_REPORT_NOT_FOUND() unsets +# if called. +SET(EIGEN_FOUND TRUE) + +# Extract Eigen version from Eigen/src/Core/util/Macros.h +IF (EIGEN_INCLUDE_DIR) + SET(EIGEN_VERSION_FILE ${EIGEN_INCLUDE_DIR}/Eigen/src/Core/util/Macros.h) + IF (NOT EXISTS ${EIGEN_VERSION_FILE}) + EIGEN_REPORT_NOT_FOUND( + "Could not find file: ${EIGEN_VERSION_FILE} " + "containing version information in Eigen install located at: " + "${EIGEN_INCLUDE_DIR}.") + ELSE (NOT EXISTS ${EIGEN_VERSION_FILE}) + FILE(READ ${EIGEN_VERSION_FILE} EIGEN_VERSION_FILE_CONTENTS) + + STRING(REGEX MATCH "#define EIGEN_WORLD_VERSION [0-9]+" + EIGEN_WORLD_VERSION "${EIGEN_VERSION_FILE_CONTENTS}") + STRING(REGEX REPLACE "#define EIGEN_WORLD_VERSION ([0-9]+)" "\\1" + EIGEN_WORLD_VERSION "${EIGEN_WORLD_VERSION}") + + STRING(REGEX MATCH "#define EIGEN_MAJOR_VERSION [0-9]+" + EIGEN_MAJOR_VERSION "${EIGEN_VERSION_FILE_CONTENTS}") + STRING(REGEX REPLACE "#define EIGEN_MAJOR_VERSION ([0-9]+)" "\\1" + EIGEN_MAJOR_VERSION "${EIGEN_MAJOR_VERSION}") + + STRING(REGEX MATCH "#define EIGEN_MINOR_VERSION [0-9]+" + EIGEN_MINOR_VERSION "${EIGEN_VERSION_FILE_CONTENTS}") + STRING(REGEX REPLACE "#define EIGEN_MINOR_VERSION ([0-9]+)" "\\1" + EIGEN_MINOR_VERSION "${EIGEN_MINOR_VERSION}") + + # This is on a single line s/t CMake does not interpret it as a list of + # elements and insert ';' separators which would result in 3.;2.;0 nonsense. + SET(EIGEN_VERSION "${EIGEN_WORLD_VERSION}.${EIGEN_MAJOR_VERSION}.${EIGEN_MINOR_VERSION}") + ENDIF (NOT EXISTS ${EIGEN_VERSION_FILE}) +ENDIF (EIGEN_INCLUDE_DIR) + +# Set standard CMake FindPackage variables if found. +IF (EIGEN_FOUND) + SET(EIGEN_INCLUDE_DIRS ${EIGEN_INCLUDE_DIR}) +ENDIF (EIGEN_FOUND) + +# Handle REQUIRED / QUIET optional arguments and version. +INCLUDE(FindPackageHandleStandardArgs) +FIND_PACKAGE_HANDLE_STANDARD_ARGS(Eigen + REQUIRED_VARS EIGEN_INCLUDE_DIRS + VERSION_VAR EIGEN_VERSION) + +# Only mark internal variables as advanced if we found Eigen, otherwise +# leave it visible in the standard GUI for the user to set manually. +IF (EIGEN_FOUND) + MARK_AS_ADVANCED(FORCE EIGEN_INCLUDE_DIR) +ENDIF (EIGEN_FOUND) diff --git a/example/CMakeLists.txt b/example/CMakeLists.txt index c0095af6..beebb664 100644 --- a/example/CMakeLists.txt +++ b/example/CMakeLists.txt @@ -25,3 +25,27 @@ foreach(EXAMPLE ${EXAMPLES}) add_executable(${EXAMPLE} ${EXAMPLE}.cpp) target_link_libraries(${EXAMPLE} ${OPENCL_LIBRARIES} ${Boost_LIBRARIES}) endforeach() + +# opencv examples +if(${BOOST_COMPUTE_HAVE_OPENCV}) + find_package(OpenCV REQUIRED) + include_directories(${OpenCV_INCLUDE_DIRS}) + add_executable(opencv_filter opencv_filter.cpp) + target_link_libraries(opencv_filter ${OPENCL_LIBRARIES} ${OpenCV_LIBRARIES}) +endif() + +# opengl/vtk examples +if(${BOOST_COMPUTE_HAVE_VTK}) + find_package(VTK REQUIRED) + include(${VTK_USE_FILE}) + add_executable(opengl_sphere opengl_sphere.cpp) + target_link_libraries(opengl_sphere ${OPENCL_LIBRARIES} ${VTK_LIBRARIES} GL) +endif() + +# qt examples +if(${BOOST_COMPUTE_HAVE_QT}) + find_package(Qt REQUIRED COMPONENTS QtCore QtGui) + include(${QT_USE_FILE}) + add_executable(qimage_blur qimage_blur.cpp) + target_link_libraries(qimage_blur ${OPENCL_LIBRARIES} ${QT_LIBRARIES}) +endif() diff --git a/example/opencv_filter.cpp b/example/opencv_filter.cpp new file mode 100644 index 00000000..314a840c --- /dev/null +++ b/example/opencv_filter.cpp @@ -0,0 +1,101 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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. +//---------------------------------------------------------------------------// + +#include + +#include +#include +#include + +#include +#include +#include +#include + +namespace compute = boost::compute; + +// this example shows how to read an image with OpenCV, transfer the +// image to the GPU, and apply a simple flip filter written in OpenCL +int main(int argc, char *argv[]) +{ + // check command line + if(argc < 2){ + std::cerr << "usage: " << argv[0] << " FILENAME" << std::endl; + return -1; + } + + // read image with opencv + cv::Mat cv_image = cv::imread(argv[1], CV_LOAD_IMAGE_COLOR); + if(!cv_image.data){ + std::cerr << "failed to load image" << std::endl; + return -1; + } + + // get default device and setup context + compute::device gpu = compute::system::default_device(); + compute::context context(gpu); + compute::command_queue queue(context, gpu); + + // convert image to BGRA (OpenCL requires 16-byte aligned data) + cv::cvtColor(cv_image, cv_image, CV_BGR2BGRA); + + // transfer image to gpu + compute::image2d input_image = + compute::opencv_create_image2d_with_mat( + cv_image, compute::image2d::read_write, queue + ); + + // create output image + compute::image2d output_image( + context, + compute::image2d::write_only, + input_image.get_format(), + input_image.width(), + input_image.height() + ); + + // create flip program + const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE( + __kernel void flip_kernel(__read_only image2d_t input, + __write_only image2d_t output) + { + const sampler_t sampler = CLK_ADDRESS_NONE | CLK_FILTER_NEAREST; + int height = get_image_height(input); + int2 input_coord = { get_global_id(0), get_global_id(1) }; + int2 output_coord = { input_coord.x, height - input_coord.y - 1 }; + float4 value = read_imagef(input, sampler, input_coord); + write_imagef(output, output_coord, value); + } + ); + + compute::program flip_program = + compute::program::create_with_source(source, context); + flip_program.build(); + + // create flip kernel and set arguments + compute::kernel flip_kernel(flip_program, "flip_kernel"); + flip_kernel.set_arg(0, input_image); + flip_kernel.set_arg(1, output_image); + + // run flip kernel + size_t origin[2] = { 0, 0 }; + size_t region[2] = { input_image.width(), input_image.height() }; + queue.enqueue_nd_range_kernel(flip_kernel, 2, origin, region, 0); + + // show host image + cv::imshow("opencv image", cv_image); + + // show gpu image + compute::opencv_imshow("filtered image", output_image, queue); + + // wait and return + cv::waitKey(0); + return 0; +} diff --git a/example/opengl_sphere.cpp b/example/opengl_sphere.cpp new file mode 100644 index 00000000..a6794b76 --- /dev/null +++ b/example/opengl_sphere.cpp @@ -0,0 +1,275 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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. +//---------------------------------------------------------------------------// + +#include +#include + +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include + +namespace compute = boost::compute; + +// tesselates a sphere with radius, phi_slices, and theta_slices. returns +// a shared opencl/opengl buffer containing the vertex data. +compute::opengl_buffer tesselate_sphere(float radius, + size_t phi_slices, + size_t theta_slices, + compute::command_queue &queue) +{ + const compute::context &context = queue.get_context(); + + const size_t vertex_count = phi_slices * theta_slices; + + // create opengl buffer + GLuint vbo; + vtkgl::GenBuffersARB(1, &vbo); + vtkgl::BindBufferARB(vtkgl::ARRAY_BUFFER, vbo); + vtkgl::BufferDataARB(vtkgl::ARRAY_BUFFER, + sizeof(float) * 4 * vertex_count, + NULL, + vtkgl::STREAM_DRAW); + vtkgl::BindBufferARB(vtkgl::ARRAY_BUFFER, 0); + + // create shared opengl/opencl buffer + compute::opengl_buffer vertex_buffer(context, vbo); + + // tesselate_sphere kernel source + const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE( + __kernel void tesselate_sphere(float radius, + uint phi_slices, + uint theta_slices, + __global float4 *vertex_buffer) + { + const uint phi_i = get_global_id(0); + const uint theta_i = get_global_id(1); + + const float phi = phi_i * 2.f * M_PI_F / phi_slices; + const float theta = theta_i * 2.f * M_PI_F / theta_slices; + + float4 v; + v.x = radius * cos(theta) * cos(phi); + v.y = radius * cos(theta) * sin(phi); + v.z = radius * sin(theta); + v.w = 1.f; + + vertex_buffer[phi_i*phi_slices+theta_i] = v; + } + ); + + // build tesselate_sphere program + compute::program program = + compute::program::create_with_source(source, context); + program.build(); + + // setup tesselate_sphere kernel + compute::kernel kernel(program, "tesselate_sphere"); + kernel.set_arg(0, radius); + kernel.set_arg(1, phi_slices); + kernel.set_arg(2, theta_slices); + kernel.set_arg(3, vertex_buffer); + + // acqurire buffer so that it is accessible to OpenCL + compute::opengl_enqueue_acquire_buffer(vertex_buffer, queue); + + // execute tesselate_sphere kernel + size_t offset[2] = { 0, 0 }; + size_t work_size[2] = { phi_slices, theta_slices }; + size_t group_size[2] = { 1, 1 }; + queue.enqueue_nd_range_kernel(kernel, 2, offset, work_size, group_size); + + // release buffer so that it is accessible to OpenGL + compute::opengl_enqueue_release_buffer(vertex_buffer, queue); + + return vertex_buffer; +} + +// simple vtkMapper subclass to render the tesselated sphere on the gpu. +class gpu_sphere_mapper : public vtkMapper +{ +public: + vtkTypeMacro(gpu_sphere_mapper, vtkMapper) + + static gpu_sphere_mapper* New() + { + return new gpu_sphere_mapper; + } + + void Render(vtkRenderer *renderer, vtkActor *actor) + { + if(!m_initialized){ + Initialize(renderer, actor); + m_initialized = true; + } + + if(!m_tesselated){ + m_vertex_count = m_phi_slices * m_theta_slices; + + // tesselate sphere + m_vertex_buffer = tesselate_sphere( + m_radius, m_phi_slices, m_theta_slices, m_command_queue + ); + + // create index array buffer + GLuint index_vbo; + vtkgl::GenBuffersARB(1, &index_vbo); + vtkgl::BindBufferARB(GL_ELEMENT_ARRAY_BUFFER, index_vbo); + vtkgl::BufferDataARB(GL_ELEMENT_ARRAY_BUFFER, + sizeof(unsigned short) * m_vertex_count, + NULL, + GL_STATIC_DRAW); + vtkgl::BindBufferARB(GL_ELEMENT_ARRAY_BUFFER, 0); + + // fill index buffer with (0, 1, 2, ..., vertex_count-1) + m_index_buffer = compute::opengl_buffer(m_context, index_vbo); + compute::opengl_enqueue_acquire_buffer(m_index_buffer, m_command_queue); + compute::iota( + compute::make_buffer_iterator(m_index_buffer, 0), + compute::make_buffer_iterator(m_index_buffer, m_vertex_count), + 0, + m_command_queue + ); + compute::opengl_enqueue_release_buffer(m_index_buffer, m_command_queue); + + // set tesselated flag to true + m_tesselated = true; + } + + // draw sphere + glEnableClientState(GL_VERTEX_ARRAY); + vtkgl::BindBufferARB(vtkgl::ARRAY_BUFFER, m_vertex_buffer.get_opengl_object()); + glVertexPointer(4, GL_FLOAT, sizeof(float)*4, 0); + + vtkgl::BindBufferARB(GL_ELEMENT_ARRAY_BUFFER, m_index_buffer.get_opengl_object()); + glDrawElements(GL_POINTS, m_vertex_count, GL_UNSIGNED_SHORT, 0); + } + + void Initialize(vtkRenderer *renderer, vtkActor *actor) + { + // initialize opengl extensions + vtkOpenGLExtensionManager *extensions = + static_cast(renderer->GetRenderWindow()) + ->GetExtensionManager(); + extensions->LoadExtension("GL_ARB_vertex_buffer_object"); + + // initialize opencl/opengl shared context + compute::device device = compute::system::default_device(); + std::cout << "device: " << device.name() << std::endl; + if(!device.supports_extension("cl_khr_gl_sharing")){ + std::cerr << "error: " + << "gpu device: " << device.name() + << " does not support OpenGL sharing" + << std::endl; + return; + } + + // create context for the gpu device + cl_context_properties properties[] = { + CL_GL_CONTEXT_KHR, (cl_context_properties) glXGetCurrentContext(), + CL_GLX_DISPLAY_KHR, (cl_context_properties) glXGetCurrentDisplay(), + 0 + }; + m_context = compute::context(device, properties); + + // create command queue for the gpu device + m_command_queue = compute::command_queue(m_context, device); + } + + double* GetBounds() + { + static double bounds[6]; + bounds[0] = -m_radius; bounds[1] = m_radius; + bounds[2] = -m_radius; bounds[3] = m_radius; + bounds[4] = -m_radius; bounds[5] = m_radius; + return bounds; + } + +protected: + gpu_sphere_mapper() + { + m_radius = 5.0f; + m_phi_slices = 100; + m_theta_slices = 100; + m_initialized = false; + m_tesselated = false; + } + +private: + float m_radius; + int m_phi_slices; + int m_theta_slices; + int m_vertex_count; + bool m_initialized; + bool m_tesselated; + compute::context m_context; + compute::command_queue m_command_queue; + compute::opengl_buffer m_vertex_buffer; + compute::opengl_buffer m_index_buffer; +}; + +int main(int argc, char *argv[]) +{ + // create gpu sphere mapper + vtkSmartPointer mapper = + vtkSmartPointer::New(); + + // create actor for gpu sphere mapper + vtkSmartPointer actor = + vtkSmartPointer::New(); + actor->GetProperty()->LightingOff(); + actor->GetProperty()->SetInterpolationToFlat(); + actor->SetMapper(mapper); + + // create render window + vtkSmartPointer renderer = + vtkSmartPointer::New(); + renderer->SetBackground(.1, .2, .31); + vtkSmartPointer renderWindow = + vtkSmartPointer::New(); + renderWindow->SetSize(800, 600); + renderWindow->AddRenderer(renderer); + vtkSmartPointer renderWindowInteractor = + vtkSmartPointer::New(); + vtkInteractorStyleSwitch *interactorStyle = + vtkInteractorStyleSwitch::SafeDownCast( + renderWindowInteractor->GetInteractorStyle() + ); + interactorStyle->SetCurrentStyleToTrackballCamera(); + renderWindowInteractor->SetRenderWindow(renderWindow); + renderer->AddActor(actor); + + // render + renderer->ResetCamera(); + vtkCamera *camera = renderer->GetActiveCamera(); + camera->Elevation(-90.0); + renderWindowInteractor->Initialize(); + renderWindow->Render(); + renderWindowInteractor->Start(); + + return 0; +} diff --git a/example/qimage_blur.cpp b/example/qimage_blur.cpp new file mode 100644 index 00000000..2a8c7e05 --- /dev/null +++ b/example/qimage_blur.cpp @@ -0,0 +1,133 @@ +#include +#include + +#include + +#include +#include +#include +#include + +namespace compute = boost::compute; + +inline void box_filter_image(const compute::image2d &input, + compute::image2d &output, + compute::uint_ box_height, + compute::uint_ box_width, + compute::command_queue &queue) +{ + const compute::context &context = queue.get_context(); + + // simple box filter kernel source + const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE( + __kernel void box_filter(__read_only image2d_t input, + __write_only image2d_t output, + uint box_height, + uint box_width) + { + int x = get_global_id(0); + int y = get_global_id(1); + int h = get_image_height(input); + int w = get_image_width(input); + int k = box_width; + int l = box_height; + + if(x < k/2 || y < l/2 || x >= w-(k/2) || y >= h-(l/2)){ + write_imagef(output, (int2)(x, y), (float4)(0, 0, 0, 1)); + } + else { + const sampler_t sampler = CLK_ADDRESS_NONE | CLK_FILTER_NEAREST; + + float4 sum = { 0, 0, 0, 0 }; + for(int i = 0; i < k; i++){ + for(int j = 0; j < l; j++){ + sum += read_imagef(input, sampler, (int2)(x+i-k, y+j-l)); + } + } + sum /= (float) k * l; + float4 value = (float4)( sum.x, sum.y, sum.z, 1.f ); + write_imagef(output, (int2)(x, y), value); + } + } + ); + + // build box filter program + compute::program program = + compute::program::create_with_source(source, context); + program.build(); + + // setup box filter kernel + compute::kernel kernel(program, "box_filter"); + kernel.set_arg(0, input); + kernel.set_arg(1, output); + kernel.set_arg(2, box_height); + kernel.set_arg(3, box_width); + + // execute the box filter kernel + size_t origin[2] = { 0, 0 }; + size_t region[2] = { input.width(), input.height() }; + size_t work_size[2] = { 1, 1 }; + + queue.enqueue_nd_range_kernel(kernel, 2, origin, region, work_size); +} + +// this example shows how to load an image using Qt, apply a simple +// box blur filter, and then display it in a Qt window. +int main(int argc, char *argv[]) +{ + QApplication app(argc, argv); + + // check command line + if(argc < 2){ + std::cout << "usage: qimage_blur [FILENAME]" << std::endl; + return -1; + } + + // load image using Qt + QString fileName = argv[1]; + QImage qimage(fileName); + + size_t height = qimage.height(); + size_t width = qimage.width(); + size_t bytes_per_line = qimage.bytesPerLine(); + + qDebug() << "height:" << height + << "width:" << width + << "bytes per line:" << bytes_per_line + << "depth:" << qimage.depth() + << "format:" << qimage.format(); + + // create compute context + compute::device gpu = compute::system::default_device(); + compute::context context(gpu); + compute::command_queue queue(context, gpu); + std::cout << "device: " << gpu.name() << std::endl; + + // get the opencl image format for the qimage + compute::image_format format = + compute::qt_qimage_format_to_image_format(qimage.format()); + + // create input and output images on the gpu + compute::image2d input_image( + context, compute::image2d::read_write, format, width, height + ); + compute::image2d output_image( + context, compute::image2d::read_write, format, width, height + ); + + // copy host qimage to gpu image + compute::qt_copy_qimage_to_image2d(qimage, input_image, queue); + + // apply box filter + box_filter_image(input_image, output_image, 7, 7, queue); + + // copy gpu blurred image from to host qimage + compute::qt_copy_image2d_to_qimage(output_image, qimage, queue); + + // show image as a pixmap + QLabel label; + label.setPixmap(QPixmap::fromImage(qimage)); + label.show(); + + return app.exec(); +} diff --git a/include/boost/compute.hpp b/include/boost/compute.hpp index 9ae9d7c7..35f8b386 100644 --- a/include/boost/compute.hpp +++ b/include/boost/compute.hpp @@ -40,8 +40,4 @@ #include #endif -#ifdef BOOST_COMPUTE_HAVE_GL -#include -#endif - #endif // BOOST_COMPUTE_HPP diff --git a/include/boost/compute/buffer.hpp b/include/boost/compute/buffer.hpp index 7d9e238c..ead99d0c 100644 --- a/include/boost/compute/buffer.hpp +++ b/include/boost/compute/buffer.hpp @@ -19,10 +19,6 @@ #include #include -#ifdef BOOST_COMPUTE_HAVE_GL -#include -#endif - namespace boost { namespace compute { @@ -123,23 +119,6 @@ public: return get_memory_info(info); } - #ifdef BOOST_COMPUTE_HAVE_GL - static buffer from_gl_buffer(const context &context, - GLuint bufobj, - cl_mem_flags flags = read_write) - { - cl_int error = 0; - cl_mem mem = clCreateFromGLBuffer(context, flags, bufobj, &error); - if(!mem){ - BOOST_THROW_EXCEPTION(runtime_exception(error)); - } - - buffer buf(mem); - clReleaseMemObject(mem); - return buf; - } - #endif // BOOST_COMPUTE_HAVE_GL - private: BOOST_COPYABLE_AND_MOVABLE(buffer) }; diff --git a/include/boost/compute/command_queue.hpp b/include/boost/compute/command_queue.hpp index 7bd5fd8d..eeea084c 100644 --- a/include/boost/compute/command_queue.hpp +++ b/include/boost/compute/command_queue.hpp @@ -28,10 +28,6 @@ #include #include -#ifdef BOOST_COMPUTE_HAVE_GL -#include -#endif - namespace boost { namespace compute { @@ -1117,54 +1113,6 @@ public: return m_queue; } - #ifdef BOOST_COMPUTE_HAVE_GL - void enqueue_acquire_gl_objects(size_t num_objects, - const cl_mem *mem_objects) - { - BOOST_ASSERT(m_queue != 0); - - cl_int ret = clEnqueueAcquireGLObjects(m_queue, - num_objects, - mem_objects, - 0, - 0, - 0); - if(ret != CL_SUCCESS){ - BOOST_THROW_EXCEPTION(runtime_exception(ret)); - } - } - - void enqueue_acquire_gl_buffer(const buffer &buffer) - { - BOOST_ASSERT(buffer.get_context() == this->get_context()); - - enqueue_acquire_gl_objects(1, &buffer.get()); - } - - void enqueue_release_gl_objects(size_t num_objects, - const cl_mem *mem_objects) - { - BOOST_ASSERT(m_queue != 0); - - cl_int ret = clEnqueueReleaseGLObjects(m_queue, - num_objects, - mem_objects, - 0, - 0, - 0); - if(ret != CL_SUCCESS){ - BOOST_THROW_EXCEPTION(runtime_exception(ret)); - } - } - - void enqueue_release_gl_buffer(const buffer &buffer) - { - BOOST_ASSERT(buffer.get_context() == this->get_context()); - - enqueue_release_gl_objects(1, &buffer.get()); - } - #endif // BOOST_COMPUTE_HAVE_GL - private: BOOST_COPYABLE_AND_MOVABLE(command_queue) diff --git a/include/boost/compute/interop/eigen.hpp b/include/boost/compute/interop/eigen.hpp new file mode 100644 index 00000000..8039395f --- /dev/null +++ b/include/boost/compute/interop/eigen.hpp @@ -0,0 +1,16 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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_INTEROP_EIGEN_HPP +#define BOOST_COMPUTE_INTEROP_EIGEN_HPP + +#include + +#endif // BOOST_COMPUTE_INTEROP_EIGEN_HPP diff --git a/include/boost/compute/interop/eigen/core.hpp b/include/boost/compute/interop/eigen/core.hpp new file mode 100644 index 00000000..f483472e --- /dev/null +++ b/include/boost/compute/interop/eigen/core.hpp @@ -0,0 +1,72 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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_INTEROP_EIGEN_EIGEN_HPP +#define BOOST_COMPUTE_INTEROP_EIGEN_EIGEN_HPP + +#include + +#include +#include +#include +#include + +namespace boost { +namespace compute { + +/// Copies \p matrix to \p buffer. +template +inline void eigen_copy_matrix_to_buffer(const Eigen::PlainObjectBase &matrix, + buffer_iterator buffer, + command_queue &queue = system::default_queue()) +{ + ::boost::compute::copy_n(matrix.data(), matrix.size(), buffer, queue); +} + +/// Copies \p buffer to \p matrix. +template +inline void eigen_copy_buffer_to_matrix(const buffer_iterator buffer, + Eigen::PlainObjectBase &matrix, + command_queue &queue = system::default_queue()) +{ + ::boost::compute::copy_n(buffer, matrix.size(), matrix.data(), queue); +} + +/// Converts an \c Eigen::Matrix4f to a \c float16_. +inline float16_ eigen_matrix4f_to_float16(const Eigen::Matrix4f &matrix) +{ + float16_ result; + std::memcpy(&result, matrix.data(), 16 * sizeof(float)); + return result; +} + +/// Converts an \c Eigen::Matrix4d to a \c double16_. +inline double16_ eigen_matrix4d_to_double16(const Eigen::Matrix4d &matrix) +{ + double16_ result; + std::memcpy(&result, matrix.data(), 16 * sizeof(double)); + return result; +} + +} // end compute namespace +} // end boost namespace + +BOOST_COMPUTE_TYPE_NAME(Eigen::Vector2i, int2); +BOOST_COMPUTE_TYPE_NAME(Eigen::Vector4i, int4); +BOOST_COMPUTE_TYPE_NAME(Eigen::Vector2f, float2); +BOOST_COMPUTE_TYPE_NAME(Eigen::Vector4f, float4); +BOOST_COMPUTE_TYPE_NAME(Eigen::Matrix2f, float8); +BOOST_COMPUTE_TYPE_NAME(Eigen::Matrix4f, float16); +BOOST_COMPUTE_TYPE_NAME(Eigen::Vector2d, double2); +BOOST_COMPUTE_TYPE_NAME(Eigen::Vector4d, double4); +BOOST_COMPUTE_TYPE_NAME(Eigen::Matrix2d, double8); +BOOST_COMPUTE_TYPE_NAME(Eigen::Matrix4d, double16); + +#endif // BOOST_COMPUTE_INTEROP_EIGEN_EIGEN_HPP diff --git a/include/boost/compute/interop/opencv.hpp b/include/boost/compute/interop/opencv.hpp new file mode 100644 index 00000000..2294f4ba --- /dev/null +++ b/include/boost/compute/interop/opencv.hpp @@ -0,0 +1,17 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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_INTEROP_OPENCV_HPP +#define BOOST_COMPUTE_INTEROP_OPENCV_HPP + +#include +#include + +#endif // BOOST_COMPUTE_INTEROP_OPENCV_HPP diff --git a/include/boost/compute/interop/opencv/core.hpp b/include/boost/compute/interop/opencv/core.hpp new file mode 100644 index 00000000..7b9c5fa7 --- /dev/null +++ b/include/boost/compute/interop/opencv/core.hpp @@ -0,0 +1,119 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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_INTEROP_OPENCV_CORE_HPP +#define BOOST_COMPUTE_INTEROP_OPENCV_CORE_HPP + +#include + +#include +#include +#include + +namespace boost { +namespace compute { + +template +inline void opencv_copy_mat_to_buffer(const cv::Mat &mat, + buffer_iterator buffer, + command_queue &queue = system::default_queue()) +{ + BOOST_ASSERT(mat.isContinuous()); + + ::boost::compute::copy_n( + reinterpret_cast(mat.data), mat.rows * mat.cols, buffer, queue + ); +} + +template +inline void opencv_copy_buffer_to_mat(const buffer_iterator buffer, + cv::Mat &mat, + command_queue &queue = system::default_queue()) +{ + BOOST_ASSERT(mat.isContinuous()); + + ::boost::compute::copy_n( + buffer, mat.cols * mat.rows, reinterpret_cast(mat.data), queue + ); +} + +inline void opencv_copy_mat_to_image(const cv::Mat &mat, + image2d &image, + command_queue &queue = system::default_queue()) +{ + BOOST_ASSERT(mat.data != 0); + BOOST_ASSERT(mat.isContinuous()); + BOOST_ASSERT(image.get_context() == queue.get_context()); + + size_t origin[2] = { 0, 0 }; + size_t region[2] = { image.width(), image.height() }; + + queue.enqueue_write_image(image, origin, region, 0, mat.data); +} + +inline void opencv_copy_image_to_mat(const image2d &image, + cv::Mat &mat, + command_queue &queue = system::default_queue()) +{ + BOOST_ASSERT(mat.isContinuous()); + BOOST_ASSERT(image.get_context() == queue.get_context()); + + size_t origin[2] = { 0, 0 }; + size_t region[2] = { image.width(), image.height() }; + + queue.enqueue_read_image(image, origin, region, 0, mat.data); +} + +inline image_format opencv_get_mat_image_format(const cv::Mat &mat) +{ + switch(mat.type()){ + case CV_8UC4: + return image_format(CL_BGRA, CL_UNORM_INT8); + case CV_16UC4: + return image_format(CL_BGRA, CL_UNORM_INT16); + case CV_32F: + return image_format(CL_INTENSITY, CL_FLOAT); + default: + return image_format(); + } +} + +inline cv::Mat opencv_create_mat_with_image2d(const image2d &image, + command_queue &queue = system::default_queue()) +{ + BOOST_ASSERT(image.get_context() == queue.get_context()); + + cv::Mat mat(image.height(), image.width(), CV_8UC4); + + opencv_copy_image_to_mat(image, mat, queue); + + return mat; +} + +inline image2d opencv_create_image2d_with_mat(const cv::Mat &mat, + cl_mem_flags flags, + command_queue &queue = system::default_queue()) +{ + const context &context = queue.get_context(); + const image_format format = opencv_get_mat_image_format(mat); + + image2d image( + context, flags, format, mat.cols, mat.rows + ); + + opencv_copy_mat_to_image(mat, image, queue); + + return image; +} + +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_INTEROP_OPENCV_CORE_HPP diff --git a/include/boost/compute/interop/opencv/highgui.hpp b/include/boost/compute/interop/opencv/highgui.hpp new file mode 100644 index 00000000..2acae626 --- /dev/null +++ b/include/boost/compute/interop/opencv/highgui.hpp @@ -0,0 +1,34 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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_INTEROP_OPENCV_HIGHGUI_HPP +#define BOOST_COMPUTE_INTEROP_OPENCV_HIGHGUI_HPP + +#include + +#include +#include + +namespace boost { +namespace compute { + +inline void opencv_imshow(const std::string &winname, + const image2d &image, + command_queue &queue = system::default_queue()) +{ + const cv::Mat mat = opencv_create_mat_with_image2d(image, queue); + + cv::imshow(winname, mat); +} + +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_INTEROP_OPENCV_HIGHGUI_HPP diff --git a/include/boost/compute/interop/opengl.hpp b/include/boost/compute/interop/opengl.hpp new file mode 100644 index 00000000..f8ac248f --- /dev/null +++ b/include/boost/compute/interop/opengl.hpp @@ -0,0 +1,18 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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_INTEROP_OPENGL_HPP +#define BOOST_COMPUTE_INTEROP_OPENGL_HPP + +#include +#include +#include + +#endif // BOOST_COMPUTE_INTEROP_OPENGL_HPP diff --git a/include/boost/compute/interop/opengl/acquire.hpp b/include/boost/compute/interop/opengl/acquire.hpp new file mode 100644 index 00000000..10cbca9e --- /dev/null +++ b/include/boost/compute/interop/opengl/acquire.hpp @@ -0,0 +1,76 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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_INTEROP_OPENGL_ACQUIRE_HPP +#define BOOST_COMPUTE_INTEROP_OPENGL_ACQUIRE_HPP + +#include +#include +#include + +namespace boost { +namespace compute { + +/// Enqueues a command to acquire the specified OpenGL memory objects. +/// +/// \see_opencl_ref{clEnqueueAcquireGLObjects} +inline void opengl_enqueue_acquire_gl_objects(size_t num_objects, + const cl_mem *mem_objects, + command_queue &queue) +{ + cl_int ret = clEnqueueAcquireGLObjects(queue.get(), + num_objects, + mem_objects, + 0, + 0, + 0); + if(ret != CL_SUCCESS){ + BOOST_THROW_EXCEPTION(runtime_exception(ret)); + } +} + +/// Enqueues a command to release the specified OpenGL memory objects. +/// +/// \see_opencl_ref{clEnqueueReleaseGLObjects} +void opengl_enqueue_release_gl_objects(size_t num_objects, + const cl_mem *mem_objects, + command_queue &queue) +{ + cl_int ret = clEnqueueReleaseGLObjects(queue.get(), + num_objects, + mem_objects, + 0, + 0, + 0); + if(ret != CL_SUCCESS){ + BOOST_THROW_EXCEPTION(runtime_exception(ret)); + } +} + +inline void opengl_enqueue_acquire_buffer(const opengl_buffer &buffer, + command_queue &queue) +{ + BOOST_ASSERT(buffer.get_context() == queue.get_context()); + + opengl_enqueue_acquire_gl_objects(1, &buffer.get(), queue); +} + +inline void opengl_enqueue_release_buffer(const opengl_buffer &buffer, + command_queue &queue) +{ + BOOST_ASSERT(buffer.get_context() == queue.get_context()); + + opengl_enqueue_release_gl_objects(1, &buffer.get(), queue); +} + +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_INTEROP_OPENGL_ACQUIRE_HPP diff --git a/include/boost/compute/cl_gl.hpp b/include/boost/compute/interop/opengl/cl_gl.hpp similarity index 69% rename from include/boost/compute/cl_gl.hpp rename to include/boost/compute/interop/opengl/cl_gl.hpp index 512cf5ef..ef1c5390 100644 --- a/include/boost/compute/cl_gl.hpp +++ b/include/boost/compute/interop/opengl/cl_gl.hpp @@ -1,5 +1,5 @@ //---------------------------------------------------------------------------// -// Copyright (c) 2013 Kyle Lutz +// Copyright (c) 2013-2014 Kyle Lutz // // Distributed under the Boost Software License, Version 1.0 // See accompanying file LICENSE_1_0.txt or copy at @@ -8,15 +8,13 @@ // See http://kylelutz.github.com/compute for more information. //---------------------------------------------------------------------------// -#ifndef BOOST_COMPUTE_CL_GL_HPP -#define BOOST_COMPUTE_CL_GL_HPP +#ifndef BOOST_COMPUTE_INTEROP_OPENGL_CL_GL_HPP +#define BOOST_COMPUTE_INTEROP_OPENGL_CL_GL_HPP #if defined(__APPLE__) -#include #include #else -#include #include #endif -#endif // BOOST_COMPUTE_CL_GL_HPP +#endif // BOOST_COMPUTE_INTEROP_OPENGL_CL_GL_HPP diff --git a/include/boost/compute/interop/opengl/gl.hpp b/include/boost/compute/interop/opengl/gl.hpp new file mode 100644 index 00000000..c00fb7a8 --- /dev/null +++ b/include/boost/compute/interop/opengl/gl.hpp @@ -0,0 +1,20 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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_INTEROP_OPENGL_GL_HPP +#define BOOST_COMPUTE_INTEROP_OPENGL_GL_HPP + +#if defined(__APPLE__) +#include +#else +#include +#endif + +#endif // BOOST_COMPUTE_INTEROP_OPENGL_GL_HPP diff --git a/include/boost/compute/interop/opengl/opengl_buffer.hpp b/include/boost/compute/interop/opengl/opengl_buffer.hpp new file mode 100644 index 00000000..a93ed342 --- /dev/null +++ b/include/boost/compute/interop/opengl/opengl_buffer.hpp @@ -0,0 +1,93 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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_INTEROP_OPENGL_OPENGL_BUFFER_HPP +#define BOOST_COMPUTE_INTEROP_OPENGL_OPENGL_BUFFER_HPP + +#include +#include +#include + +namespace boost { +namespace compute { + +/// \class opengl_buffer +/// +/// A OpenCL buffer for accessing an OpenGL memory object. +class opengl_buffer : public buffer +{ +public: + /// Creates a null OpenGL buffer object. + opengl_buffer() + : buffer() + { + } + + /// Creates a new OpenGL buffer object for \p mem. + explicit opengl_buffer(cl_mem mem, bool retain = true) + : buffer(mem, retain) + { + } + + /// Creates a new OpenGL buffer object in \p context for \p bufobj + /// with \p flags. + opengl_buffer(const context &context, + GLuint bufobj, + cl_mem_flags flags = read_write) + { + cl_int error = 0; + m_mem = clCreateFromGLBuffer(context, flags, bufobj, &error); + if(!m_mem){ + BOOST_THROW_EXCEPTION(runtime_exception(error)); + } + } + + /// Creates a new OpenGL buffer object as a copy of \p other. + opengl_buffer(const opengl_buffer &other) + : buffer(other) + { + } + + /// Copies the OpenGL buffer object from \p other. + opengl_buffer& operator=(const opengl_buffer &other) + { + if(this != &other){ + buffer::operator=(other); + } + + return *this; + } + + /// Destroys the OpenGL buffer object. + ~opengl_buffer() + { + } + + /// Returns the OpenGL memory object ID. + GLuint get_opengl_object() const + { + GLuint object = 0; + clGetGLObjectInfo(m_mem, 0, &object); + return object; + } + + /// Returns the OpenGL memory object type. + cl_gl_object_type get_opengl_type() const + { + cl_gl_object_type type; + clGetGLObjectInfo(m_mem, &type, 0); + return type; + } +}; + +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_INTEROP_OPENGL_OPENGL_BUFFER_HPP diff --git a/include/boost/compute/interop/opengl/opengl_texture.hpp b/include/boost/compute/interop/opengl/opengl_texture.hpp new file mode 100644 index 00000000..c5325412 --- /dev/null +++ b/include/boost/compute/interop/opengl/opengl_texture.hpp @@ -0,0 +1,92 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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_INTEROP_OPENGL_OPENGL_TEXTURE_HPP +#define BOOST_COMPUTE_INTEROP_OPENGL_OPENGL_TEXTURE_HPP + +#include +#include +#include +#include + +namespace boost { +namespace compute { + +class opengl_texture : public memory_object +{ +public: + opengl_texture() + : memory_object() + { + } + + explicit opengl_texture(cl_mem mem, bool retain = true) + : memory_object(mem, retain) + { + } + + opengl_texture(const context &context, + GLenum texture_target, + GLint miplevel, + GLuint texture, + cl_mem_flags flags = read_write) + { + cl_int error = 0; + + #ifdef CL_VERSION_1_2 + m_mem = clCreateFromGLTexture(context, + flags, + texture_target, + miplevel, + texture, + &error); + #else + m_mem = clCreateFromGLTexture2D(context, + flags, + texture_target, + miplevel, + texture, + &error); + #endif + + if(!m_mem){ + BOOST_THROW_EXCEPTION(runtime_exception(error)); + } + } + + opengl_texture(const opengl_texture &other) + : memory_object(other) + { + } + + opengl_texture& operator=(const opengl_texture &other) + { + if(this != &other){ + memory_object::operator=(other); + } + + return *this; + } + + ~opengl_texture() + { + } + + template + T get_texture_info(cl_gl_texture_info info) const + { + return detail::get_object_info(clGetGLTextureInfo, m_mem, info); + } +}; + +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_INTEROP_OPENGL_OPENGL_TEXTURE_HPP diff --git a/include/boost/compute/interop/qt.hpp b/include/boost/compute/interop/qt.hpp new file mode 100644 index 00000000..0dd40e7a --- /dev/null +++ b/include/boost/compute/interop/qt.hpp @@ -0,0 +1,17 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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_INTEROP_QT_HPP +#define BOOST_COMPUTE_INTEROP_QT_HPP + +#include +#include + +#endif // BOOST_COMPUTE_INTEROP_QT_HPP diff --git a/include/boost/compute/interop/qt/qimage.hpp b/include/boost/compute/interop/qt/qimage.hpp new file mode 100644 index 00000000..034fb8e3 --- /dev/null +++ b/include/boost/compute/interop/qt/qimage.hpp @@ -0,0 +1,70 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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_INTEROP_QT_QIMAGE_HPP +#define BOOST_COMPUTE_INTEROP_QT_QIMAGE_HPP + +#include +#include +#include + +#include + +namespace boost { +namespace compute { + +inline image_format qt_qimage_format_to_image_format(const QImage::Format &format) +{ + if(format == QImage::Format_RGB32){ + return image_format(image_format::bgra, image_format::unorm_int8); + } + + return image_format(); +} + +inline QImage::Format qt_image_format_to_qimage_format(const image_format &format) +{ + if(format == image_format(image_format::bgra, image_format::unorm_int8)){ + return QImage::Format_RGB32; + } + + return QImage::Format_Invalid; +} + +inline image_format qt_qimage_get_format(const QImage &image) +{ + return qt_qimage_format_to_image_format(image.format()); +} + +inline void qt_copy_qimage_to_image2d(const QImage &qimage, + image2d &image, + command_queue &queue) +{ + size_t origin[] = { 0, 0 }; + size_t region[] = { image.width(), image.height() }; + + queue.enqueue_write_image(image, origin, region, 0, qimage.constBits()); +} + +inline void qt_copy_image2d_to_qimage(const image2d &image, + QImage &qimage, + command_queue &queue) +{ + size_t origin[] = { 0, 0 }; + size_t region[] = { static_cast(qimage.width()), + static_cast(qimage.height()) }; + + queue.enqueue_read_image(image, origin, region, 0, qimage.bits()); +} + +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_INTEROP_QT_QIMAGE_HPP diff --git a/include/boost/compute/interop/qt/qpoint.hpp b/include/boost/compute/interop/qt/qpoint.hpp new file mode 100644 index 00000000..a852c967 --- /dev/null +++ b/include/boost/compute/interop/qt/qpoint.hpp @@ -0,0 +1,20 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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_INTEROP_QT_QPOINT_HPP +#define BOOST_COMPUTE_INTEROP_QT_QPOINT_HPP + +#include + +#include + +BOOST_COMPUTE_TYPE_NAME(QPoint, "int2") + +#endif // BOOST_COMPUTE_INTEROP_QT_QPOINT_HPP diff --git a/include/boost/compute/interop/qt/qpointf.hpp b/include/boost/compute/interop/qt/qpointf.hpp new file mode 100644 index 00000000..76fe5b05 --- /dev/null +++ b/include/boost/compute/interop/qt/qpointf.hpp @@ -0,0 +1,20 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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_INTEROP_QT_QPOINTF_HPP +#define BOOST_COMPUTE_INTEROP_QT_QPOINTF_HPP + +#include + +#include + +BOOST_COMPUTE_TYPE_NAME(QPointF, "float2") + +#endif // BOOST_COMPUTE_INTEROP_QT_QPOINTF_HPP diff --git a/include/boost/compute/interop/qt/qtcore.hpp b/include/boost/compute/interop/qt/qtcore.hpp new file mode 100644 index 00000000..e29ada5c --- /dev/null +++ b/include/boost/compute/interop/qt/qtcore.hpp @@ -0,0 +1,18 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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_INTEROP_QT_QTCORE_HPP +#define BOOST_COMPUTE_INTEROP_QT_QTCORE_HPP + +#include +#include +#include + +#endif // BOOST_COMPUTE_INTEROP_QT_QTCORE_HPP diff --git a/include/boost/compute/interop/qt/qtgui.hpp b/include/boost/compute/interop/qt/qtgui.hpp new file mode 100644 index 00000000..9d7ec0c2 --- /dev/null +++ b/include/boost/compute/interop/qt/qtgui.hpp @@ -0,0 +1,16 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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_INTEROP_QT_QTGUI_HPP +#define BOOST_COMPUTE_INTEROP_QT_QTGUI_HPP + +#include + +#endif // BOOST_COMPUTE_INTEROP_QT_QTGUI_HPP diff --git a/include/boost/compute/interop/qt/qvector.hpp b/include/boost/compute/interop/qt/qvector.hpp new file mode 100644 index 00000000..1733f143 --- /dev/null +++ b/include/boost/compute/interop/qt/qvector.hpp @@ -0,0 +1,48 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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_INTEROP_QT_QVECTOR_HPP +#define BOOST_COMPUTE_INTEROP_QT_QVECTOR_HPP + +#include + +#include + +namespace boost { +namespace compute { +namespace detail { + +template +struct _is_contiguous_iterator< + Iterator, + typename boost::enable_if< + typename boost::is_same< + Iterator, + typename QVector::iterator + >::type + >::type +> : public boost::true_type {}; + +template +struct _is_contiguous_iterator< + Iterator, + typename boost::enable_if< + typename boost::is_same< + Iterator, + typename QVector::const_iterator + >::type + >::type +> : public boost::true_type {}; + +} // end detail namespace +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_INTEROP_QT_QVECTOR_HPP diff --git a/include/boost/compute/interop/vtk.hpp b/include/boost/compute/interop/vtk.hpp new file mode 100644 index 00000000..c782b9fb --- /dev/null +++ b/include/boost/compute/interop/vtk.hpp @@ -0,0 +1,19 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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_INTEROP_VTK_HPP +#define BOOST_COMPUTE_INTEROP_VTK_HPP + +#include +#include +#include +#include + +#endif // BOOST_COMPUTE_INTEROP_VTK_HPP diff --git a/include/boost/compute/interop/vtk/bounds.hpp b/include/boost/compute/interop/vtk/bounds.hpp new file mode 100644 index 00000000..d962ae6b --- /dev/null +++ b/include/boost/compute/interop/vtk/bounds.hpp @@ -0,0 +1,59 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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_INTEROP_VTK_BOUNDS_HPP +#define BOOST_COMPUTE_INTEROP_VTK_BOUNDS_HPP + +#include +#include + +#include +#include +#include +#include +#include + +namespace boost { +namespace compute { + +/// Calculates the bounds for the points in the range [\p first, \p last) and +/// stores the result in \p bounds. +/// +/// For example, this can be used to implement the GetBounds() method for a +/// vtkMapper subclass. +template +inline void vtk_compute_bounds(PointIterator first, + PointIterator last, + double bounds[6], + command_queue &queue = system::default_queue()) +{ + typedef typename std::iterator_traits::value_type T; + + const context &context = queue.get_context(); + + // compute min and max point + array extrema(context); + reduce(first, last, extrema.begin() + 0, min(), queue); + reduce(first, last, extrema.begin() + 1, max(), queue); + + // copy results to host buffer + std::vector buffer(2); + copy_n(extrema.begin(), 2, buffer.begin(), queue); + + // copy to vtk-style bounds + bounds[0] = buffer[0][0]; bounds[1] = buffer[1][0]; + bounds[2] = buffer[0][1]; bounds[3] = buffer[1][1]; + bounds[4] = buffer[0][2]; bounds[5] = buffer[1][2]; +} + +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_INTEROP_VTK_BOUNDS_HPP diff --git a/include/boost/compute/interop/vtk/data_array.hpp b/include/boost/compute/interop/vtk/data_array.hpp new file mode 100644 index 00000000..315bcfa3 --- /dev/null +++ b/include/boost/compute/interop/vtk/data_array.hpp @@ -0,0 +1,65 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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_INTEROP_VTK_DATA_ARRAY_HPP +#define BOOST_COMPUTE_INTEROP_VTK_DATA_ARRAY_HPP + +#include +#include + +#include +#include +#include +#include +#include + +namespace boost { +namespace compute { + +/// Copies the values in \p data to \p buffer. +template +inline void vtk_copy_data_array_to_buffer(const vtkDataArray *data, + buffer_iterator buffer, + command_queue &queue = system::default_queue()); + +/// \internal_ +template +inline void vtk_copy_data_array_to_buffer(const vtkDataArrayTemplate *data, + buffer_iterator buffer, + command_queue &queue = system::default_queue()) +{ + vtkDataArrayTemplate *data_ = const_cast *>(data); + const T *data_ptr = static_cast(data_->GetVoidPointer(0)); + size_t data_size = data_->GetNumberOfComponents() * data_->GetNumberOfTuples(); + ::boost::compute::copy_n(data_ptr, data_size, buffer, queue); +} + +/// Copies the values in the range [\p first, \p last) to \p data. +template +inline void vtk_copy_buffer_to_data_array(buffer_iterator first, + buffer_iterator last, + vtkDataArray *data, + command_queue &queue = system::default_queue()); + +/// \internal_ +template +inline void vtk_copy_buffer_to_data_array(buffer_iterator first, + buffer_iterator last, + vtkDataArrayTemplate *data, + command_queue &queue = system::default_queue()) +{ + T *data_ptr = static_cast(data->GetVoidPointer(0)); + ::boost::compute::copy(first, last, data_ptr, queue); +} + +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_INTEROP_VTK_DATA_ARRAY_HPP diff --git a/include/boost/compute/interop/vtk/matrix4x4.hpp b/include/boost/compute/interop/vtk/matrix4x4.hpp new file mode 100644 index 00000000..55162604 --- /dev/null +++ b/include/boost/compute/interop/vtk/matrix4x4.hpp @@ -0,0 +1,46 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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_INTEROP_VTK_MATRIX4X4_HPP +#define BOOST_COMPUTE_INTEROP_VTK_MATRIX4X4_HPP + +#include + +#include + +namespace boost { +namespace compute { + +/// Converts a \c vtkMatrix4x4 to a \c float16_. +inline float16_ vtk_matrix4x4_to_float16(const vtkMatrix4x4 *matrix) +{ + float16_ result; + + for(int i = 0; i < 4; i++){ + for(int j = 0; j < 4; j++){ + result[i*4+j] = matrix->GetElement(i, j); + } + } + + return result; +} + +/// Converts a \c vtkMatrix4x4 to a \c double16_; +inline double16_ vtk_matrix4x4_to_double16(const vtkMatrix4x4 *matrix) +{ + double16_ result; + std::memcpy(&result, matrix->Element, 16 * sizeof(double)); + return result; +} + +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_INTEROP_VTK_MATRIX4X4_HPP diff --git a/include/boost/compute/interop/vtk/points.hpp b/include/boost/compute/interop/vtk/points.hpp new file mode 100644 index 00000000..b2fdd427 --- /dev/null +++ b/include/boost/compute/interop/vtk/points.hpp @@ -0,0 +1,55 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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_INTEROP_VTK_POINTS_HPP +#define BOOST_COMPUTE_INTEROP_VTK_POINTS_HPP + +#include + +#include + +#include +#include +#include +#include + +namespace boost { +namespace compute { + +/// Copies \p points to \p buffer. +/// +/// For example, to copy from a \c vtkPoints object to a \c vector: +/// \code +/// vtkPoints *points = ... +/// vector vector(points->GetNumberOfPoints(), context); +/// vtk_copy_points_to_buffer(points, vector.begin(), queue); +/// \endcode +template +inline void vtk_copy_points_to_buffer(const vtkPoints *points, + buffer_iterator buffer, + command_queue &queue = system::default_queue()) +{ + vtkPoints *points_ = const_cast(points); + + // copy points to aligned buffer + std::vector tmp(points_->GetNumberOfPoints()); + for(vtkIdType i = 0; i < points_->GetNumberOfPoints(); i++){ + double *p = points_->GetPoint(i); + tmp[i] = PointType(p[0], p[1], p[2], 1); + } + + // copy data to device + copy(tmp.begin(), tmp.end(), buffer, queue); +} + +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_INTEROP_VTK_POINTS_HPP diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 9716f036..af3fb3cf 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -113,3 +113,38 @@ add_compute_test("ext.program_cache" test_program_cache.cpp) add_compute_test("ext.tuple" test_tuple.cpp) add_compute_test("issue.11" test_issue_11.cpp) + +# eigen interop tests +if(${BOOST_COMPUTE_HAVE_EIGEN}) + find_package(Eigen REQUIRED) + include_directories(${EIGEN_INCLUDE_DIRS}) + add_compute_test("interop.eigen" test_interop_eigen.cpp) +endif() + +# opencv interop tests +if(${BOOST_COMPUTE_HAVE_OPENCV}) + find_package(OpenCV REQUIRED) + include_directories(${OpenCV_INCLUDE_DIRS}) + add_compute_test("interop.opencv" test_interop_opencv.cpp) + target_link_libraries(test_interop_opencv ${OpenCV_LIBRARIES}) +endif() + +# qt interop tests +if(${BOOST_COMPUTE_HAVE_QT}) + find_package(Qt REQUIRED COMPONENTS QtCore QtGui QtOpenGL) + include(${QT_USE_FILE}) + add_compute_test("interop.qt" test_interop_qt.cpp) + target_link_libraries(test_interop_qt ${QT_LIBRARIES}) + + # the opengl interop test depends on qt to create the opengl context + add_compute_test("interop.opengl" test_interop_opengl.cpp) + target_link_libraries(test_interop_opengl ${QT_LIBRARIES}) +endif() + +# vtk interop tests +if(${BOOST_COMPUTE_HAVE_VTK}) + find_package(VTK REQUIRED) + include(${VTK_USE_FILE}) + add_compute_test("interop.vtk" test_interop_vtk.cpp) + target_link_libraries(test_interop_vtk ${VTK_LIBRARIES}) +endif() diff --git a/test/test_interop_eigen.cpp b/test/test_interop_eigen.cpp new file mode 100644 index 00000000..e9cec277 --- /dev/null +++ b/test/test_interop_eigen.cpp @@ -0,0 +1,102 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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. +//---------------------------------------------------------------------------// + +#define BOOST_TEST_MODULE TestInteropEigen +#include + +#include +#include +#include +#include + +#include "check_macros.hpp" +#include "context_setup.hpp" + +namespace bcl = boost::compute; + +BOOST_AUTO_TEST_CASE(eigen) +{ + Eigen::MatrixXf mat(3, 3); + mat << 1, 2, 3, + 6, 5, 4, + 7, 8, 9; + + // copy matrix to gpu buffer + bcl::vector vec(9, context); + bcl::eigen_copy_matrix_to_buffer(mat, vec.begin(), queue); + CHECK_RANGE_EQUAL(float, 9, vec, (1, 6, 7, 2, 5, 8, 3, 4, 9)); + + // transpose matrix and then copy to gpu buffer + mat = mat.transpose().eval(); + bcl::eigen_copy_matrix_to_buffer(mat, vec.begin(), queue); + CHECK_RANGE_EQUAL(float, 9, vec, (1, 2, 3, 6, 5, 4, 7, 8, 9)); + + // set matrix to zero and copy data back from gpu buffer + mat.setZero(); + bcl::eigen_copy_buffer_to_matrix(vec.begin(), mat, queue); + BOOST_CHECK(mat.isZero() == false); + BOOST_CHECK_EQUAL(mat.sum(), 45); +} + +BOOST_AUTO_TEST_CASE(eigen_types) +{ + BOOST_CHECK(std::strcmp(bcl::type_name(), "int2") == 0); + BOOST_CHECK(std::strcmp(bcl::type_name(), "float2") == 0); + BOOST_CHECK(std::strcmp(bcl::type_name(), "float4") == 0); + BOOST_CHECK(std::strcmp(bcl::type_name(), "double4") == 0); +} + +BOOST_AUTO_TEST_CASE(multiply_matrix4) +{ + std::vector host_vectors; + std::vector host_matrices; + + host_matrices.push_back(Eigen::Matrix4f::Identity()); + host_matrices.push_back(Eigen::Matrix4f::Zero()); + host_matrices.push_back(Eigen::Matrix4f::Ones()); + host_matrices.push_back(Eigen::Matrix4f::Ones() * 2); + + host_vectors.push_back(Eigen::Vector4f(1, 2, 3, 4)); + host_vectors.push_back(Eigen::Vector4f(4, 3, 2, 1)); + host_vectors.push_back(Eigen::Vector4f(1, 2, 3, 4)); + host_vectors.push_back(Eigen::Vector4f(4, 3, 2, 1)); + + // returns the result of M*x + BOOST_COMPUTE_FUNCTION(Eigen::Vector4f, transform4x4, (const Eigen::Matrix4f, const Eigen::Vector4f), + { + float4 r; + r.x = dot(_1.s048c, _2); + r.y = dot(_1.s159d, _2); + r.z = dot(_1.s26ae, _2); + r.w = dot(_1.s37bf, _2); + return r; + }); + + bcl::vector vectors(4, context); + bcl::vector matrices(4, context); + bcl::vector results(4, context); + + bcl::copy(host_vectors.begin(), host_vectors.end(), vectors.begin(), queue); + bcl::copy(host_matrices.begin(), host_matrices.end(), matrices.begin(), queue); + + bcl::transform( + matrices.begin(), matrices.end(), vectors.begin(), results.begin(), transform4x4, queue + ); + + std::vector host_results(4); + bcl::copy(results.begin(), results.end(), host_results.begin(), queue); + + BOOST_CHECK((host_matrices[0] * host_vectors[0]) == host_results[0]); + BOOST_CHECK((host_matrices[1] * host_vectors[1]) == host_results[1]); + BOOST_CHECK((host_matrices[2] * host_vectors[2]) == host_results[2]); + BOOST_CHECK((host_matrices[3] * host_vectors[3]) == host_results[3]); +} + +BOOST_AUTO_TEST_SUITE_END() diff --git a/test/test_interop_opencv.cpp b/test/test_interop_opencv.cpp new file mode 100644 index 00000000..5838a5a5 --- /dev/null +++ b/test/test_interop_opencv.cpp @@ -0,0 +1,62 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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. +//---------------------------------------------------------------------------// + +#define BOOST_TEST_MODULE TestInteropOpenCV +#include + +#include +#include +#include +#include + +#include "check_macros.hpp" +#include "context_setup.hpp" + +namespace bcl = boost::compute; + +BOOST_AUTO_TEST_CASE(opencv_mat_to_buffer) +{ + // create opencv mat + cv::Mat mat(1, 4, CV_32F); + mat.at(0, 0) = 0.0f; + mat.at(0, 1) = 2.5f; + mat.at(0, 2) = 4.1f; + mat.at(0, 3) = 5.6f; + + // copy mat to gpu vector + bcl::vector vector(4, context); + bcl::opencv_copy_mat_to_buffer(mat, vector.begin(), queue); + CHECK_RANGE_EQUAL(float, 4, vector, (0.0f, 2.5f, 4.1f, 5.6f)); + + // reverse gpu vector and copy back to mat + bcl::reverse(vector.begin(), vector.end(), queue); + bcl::opencv_copy_buffer_to_mat(vector.begin(), mat, queue); + BOOST_CHECK_EQUAL(mat.at(0), 5.6f); + BOOST_CHECK_EQUAL(mat.at(1), 4.1f); + BOOST_CHECK_EQUAL(mat.at(2), 2.5f); + BOOST_CHECK_EQUAL(mat.at(3), 0.0f); +} + +BOOST_AUTO_TEST_CASE(opencv_image_format) +{ + // 32-bit BGRA + BOOST_CHECK( + bcl::opencv_get_mat_image_format(cv::Mat(32, 32, CV_8UC4)) == + bcl::image_format(CL_BGRA, CL_UNORM_INT8) + ); + + // 32-bit float + BOOST_CHECK( + bcl::opencv_get_mat_image_format(cv::Mat(32, 32, CV_32F)) == + bcl::image_format(CL_INTENSITY, CL_FLOAT) + ); +} + +BOOST_AUTO_TEST_SUITE_END() diff --git a/test/test_interop_opengl.cpp b/test/test_interop_opengl.cpp new file mode 100644 index 00000000..684f3696 --- /dev/null +++ b/test/test_interop_opengl.cpp @@ -0,0 +1,18 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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. +//---------------------------------------------------------------------------// + +#define BOOST_TEST_MODULE TestInteropOpenGL +#include + +#include + +BOOST_AUTO_TEST_CASE(opengl_buffer) +{ +} diff --git a/test/test_interop_qt.cpp b/test/test_interop_qt.cpp new file mode 100644 index 00000000..88f30ab8 --- /dev/null +++ b/test/test_interop_qt.cpp @@ -0,0 +1,95 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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. +//---------------------------------------------------------------------------// + +#define BOOST_TEST_MODULE TestInteropQt +#include + +#include +#include +#include +#include + +#include +#include + +#include "check_macros.hpp" +#include "context_setup.hpp" + +namespace bcl = boost::compute; + +BOOST_AUTO_TEST_CASE(qimage_format) +{ + BOOST_CHECK( + bcl::qt_qimage_format_to_image_format(QImage::Format_RGB32) == + bcl::image_format(CL_BGRA, CL_UNORM_INT8) + ); +} + +BOOST_AUTO_TEST_CASE(copy_qvector_to_device) +{ + QList qvector; + qvector.append(0); + qvector.append(2); + qvector.append(4); + qvector.append(6); + + bcl::vector vector(4, context); + bcl::copy(qvector.begin(), qvector.end(), vector.begin(), queue); + CHECK_RANGE_EQUAL(int, 4, vector, (0, 2, 4, 6)); +} + +BOOST_AUTO_TEST_CASE(copy_qlist_to_device) +{ + QList list; + list.append(1); + list.append(3); + list.append(5); + list.append(7); + + bcl::vector vector(4, context); + bcl::copy(list.begin(), list.end(), vector.begin(), queue); + CHECK_RANGE_EQUAL(int, 4, vector, (1, 3, 5, 7)); +} + +BOOST_AUTO_TEST_CASE(qvector_of_qpoint) +{ + QVector qt_points; + qt_points.append(QPoint(0, 1)); + qt_points.append(QPoint(2, 3)); + qt_points.append(QPoint(4, 5)); + qt_points.append(QPoint(6, 7)); + + bcl::vector bcl_points(qt_points.size(), context); + bcl::copy(qt_points.begin(), qt_points.end(), bcl_points.begin(), queue); +} + +BOOST_AUTO_TEST_CASE(qvector_of_qpointf) +{ + QVector qt_points; + qt_points.append(QPointF(0.3f, 1.7f)); + qt_points.append(QPointF(2.3f, 3.7f)); + qt_points.append(QPointF(4.3f, 5.7f)); + qt_points.append(QPointF(6.3f, 7.7f)); + + bcl::vector bcl_points(qt_points.size(), context); + bcl::copy(qt_points.begin(), qt_points.end(), bcl_points.begin(), queue); +} + +BOOST_AUTO_TEST_CASE(qvector_iterator) +{ + using boost::compute::detail::is_contiguous_iterator; + + BOOST_STATIC_ASSERT(is_contiguous_iterator::iterator>::value == true); + BOOST_STATIC_ASSERT(is_contiguous_iterator::const_iterator>::value == true); + BOOST_STATIC_ASSERT(is_contiguous_iterator::iterator>::value == false); + BOOST_STATIC_ASSERT(is_contiguous_iterator::const_iterator>::value == false); +} + +BOOST_AUTO_TEST_SUITE_END() diff --git a/test/test_interop_vtk.cpp b/test/test_interop_vtk.cpp new file mode 100644 index 00000000..c2c72ecd --- /dev/null +++ b/test/test_interop_vtk.cpp @@ -0,0 +1,118 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2013-2014 Kyle Lutz +// +// 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. +//---------------------------------------------------------------------------// + +#define BOOST_TEST_MODULE TestInteropVTK +#include + +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include + +#include "check_macros.hpp" +#include "context_setup.hpp" + +namespace compute = boost::compute; + +BOOST_AUTO_TEST_CASE(bounds) +{ + using compute::float4_; + + // create vtk points + vtkNew points; + points->InsertNextPoint(0.0, 0.0, 0.0); + points->InsertNextPoint(1.0, 2.0, 1.0); + points->InsertNextPoint(-1.0, -3.0, -1.0); + points->InsertNextPoint(0.5, 2.5, 1.5); + + // copy points to vector on gpu + compute::vector vector(points->GetNumberOfPoints(), context); + compute::vtk_copy_points_to_buffer(points.GetPointer(), vector.begin(), queue); + + // compute bounds + double bounds[6]; + compute::vtk_compute_bounds(vector.begin(), vector.end(), bounds, queue); + + // check bounds + BOOST_CHECK_CLOSE(bounds[0], -1.0, 1e-8); + BOOST_CHECK_CLOSE(bounds[1], 1.0, 1e-8); + BOOST_CHECK_CLOSE(bounds[2], -3.0, 1e-8); + BOOST_CHECK_CLOSE(bounds[3], 2.5, 1e-8); + BOOST_CHECK_CLOSE(bounds[4], -1.0, 1e-8); + BOOST_CHECK_CLOSE(bounds[5], 1.5, 1e-8); +} + +BOOST_AUTO_TEST_CASE(copy_uchar_array) +{ + // create vtk uchar vector containing 3 RGBA colors + vtkNew array; + array->SetNumberOfComponents(4); + + unsigned char red[4] = { 255, 0, 0, 255 }; + array->InsertNextTupleValue(red); + unsigned char green[4] = { 0, 255, 0, 255 }; + array->InsertNextTupleValue(green); + unsigned char blue[4] = { 0, 0, 255, 255 }; + array->InsertNextTupleValue(blue); + + // create vector on device and copy values from vtk array + compute::vector vector(3, context); + compute::vtk_copy_data_array_to_buffer( + array.GetPointer(), + compute::make_buffer_iterator(vector.get_buffer(), 0), + queue + ); + + // check values + std::vector host_vector(3); + compute::copy( + vector.begin(), vector.end(), host_vector.begin(), queue + ); + BOOST_CHECK(host_vector[0] == compute::uchar4_(255, 0, 0, 255)); + BOOST_CHECK(host_vector[1] == compute::uchar4_(0, 255, 0, 255)); + BOOST_CHECK(host_vector[2] == compute::uchar4_(0, 0, 255, 255)); +} + +BOOST_AUTO_TEST_CASE(sort_float_array) +{ + // create vtk float array + vtkNew array; + array->InsertNextValue(2.5f); + array->InsertNextValue(1.0f); + array->InsertNextValue(6.5f); + array->InsertNextValue(4.0f); + + // create vector on device and copy values from vtk array + compute::vector vector(4, context); + compute::vtk_copy_data_array_to_buffer(array.GetPointer(), vector.begin(), queue); + + // sort values on the gpu + compute::sort(vector.begin(), vector.end(), queue); + CHECK_RANGE_EQUAL(float, 4, vector, (1.0f, 2.5f, 4.0f, 6.5f)); + + // copy sorted values back to the vtk array + compute::vtk_copy_buffer_to_data_array( + vector.begin(), vector.end(), array.GetPointer(), queue + ); + BOOST_CHECK_EQUAL(array->GetValue(0), 1.0f); + BOOST_CHECK_EQUAL(array->GetValue(1), 2.5f); + BOOST_CHECK_EQUAL(array->GetValue(2), 4.0f); + BOOST_CHECK_EQUAL(array->GetValue(3), 6.5f); +} + +BOOST_AUTO_TEST_SUITE_END()