[arrayfire] 137/248: 3d surface rendering features
Ghislain Vaillant
ghisvail-guest at moszumanska.debian.org
Tue Nov 17 15:54:16 UTC 2015
This is an automated email from the git hooks/post-receive script.
ghisvail-guest pushed a commit to branch dfsg-clean
in repository arrayfire.
commit 2a1d63d6a89c0d16b5094fd29744a85f96d10583
Author: syurkevi <stefan at arrayfire.com>
Date: Mon Oct 19 17:58:07 2015 -0400
3d surface rendering features
---
examples/graphics/plot3.cpp | 4 +-
examples/graphics/{plot3.cpp => surface.cpp} | 35 +++---
include/af/graphics.h | 47 +++++++-
src/api/c/graphics_common.cpp | 19 ++++
src/api/c/graphics_common.hpp | 5 +
src/api/c/surface.cpp | 133 ++++++++++++++++++++++
src/api/cpp/graphics.cpp | 14 +++
src/api/unified/graphics.cpp | 5 +
src/backend/cpu/surface.cpp | 48 ++++++++
src/backend/{opencl/plot.hpp => cpu/surface.hpp} | 5 +-
src/backend/cuda/surface.cu | 59 ++++++++++
src/backend/{opencl/plot.hpp => cuda/surface.hpp} | 5 +-
src/backend/opencl/interopManager.cpp | 12 ++
src/backend/opencl/interopManager.hpp | 1 +
src/backend/opencl/plot.hpp | 1 -
src/backend/opencl/surface.cpp | 73 ++++++++++++
src/backend/opencl/{plot.hpp => surface.hpp} | 2 +-
17 files changed, 438 insertions(+), 30 deletions(-)
diff --git a/examples/graphics/plot3.cpp b/examples/graphics/plot3.cpp
index 40bd6d4..ea2ca8d 100644
--- a/examples/graphics/plot3.cpp
+++ b/examples/graphics/plot3.cpp
@@ -34,7 +34,9 @@ int main(int argc, char *argv[])
Y = max(min(Y, bounds),-bounds);
array Pts = join(1, X, Y, Z);
- myWindow.plot3(flat(Pts));
+ //Pts can be passed in as a matrix in the form n x 3, 3 x n
+ //or in the flattened xyz-triplet array with size 3n x 1
+ myWindow.plot3(Pts);
t+=0.01;
} while(!myWindow.close());
diff --git a/examples/graphics/plot3.cpp b/examples/graphics/surface.cpp
similarity index 57%
copy from examples/graphics/plot3.cpp
copy to examples/graphics/surface.cpp
index 40bd6d4..3517617 100644
--- a/examples/graphics/plot3.cpp
+++ b/examples/graphics/surface.cpp
@@ -13,7 +13,7 @@
using namespace af;
-static const int ITERATIONS = 200;
+static const int ITERATIONS = 30;
static const float PRECISION = 1.0f/ITERATIONS;
int main(int argc, char *argv[])
@@ -21,23 +21,22 @@ int main(int argc, char *argv[])
try {
// Initialize the kernel array just once
af::info();
- af::Window myWindow(800, 800, "3D Line Plot example: ArrayFire");
-
- static float t=0.1;
- array Z = seq( 0.1f, 10.f, PRECISION);
- array bounds = constant(1, Z.dims());
-
- do{
- array Y = sin((Z*t) + t) / Z;
- array X = cos((Z*t) + t) / Z;
- X = max(min(X, bounds),-bounds);
- Y = max(min(Y, bounds),-bounds);
-
- array Pts = join(1, X, Y, Z);
- myWindow.plot3(flat(Pts));
-
- t+=0.01;
- } while(!myWindow.close());
+ af::Window myWindow(800, 800, "3D Surface example: ArrayFire");
+
+ array X = seq(-1, 1, PRECISION);
+ array Y = seq(-1, 1, PRECISION);
+ array Z = randn(X.dims(0), Y.dims(0));
+
+ static float t=0;
+ for (double val=-af::Pi; !myWindow.close(); ) {
+ t+=0.07;
+ //Z = sin(tile(X,1, Y.dims(0))*t + t) + cos(transpose(tile(Y, 1, X.dims(0)))*t + t);
+ array x = tile(X,1, Y.dims(0));
+ array y = transpose(tile(Y, 1, X.dims(0)));
+ Z = 10*x*-abs(y) * cos(x*x*(y+t))+sin(y*(x+t))-1.5;
+
+ myWindow.surface(X, Y, Z, NULL);
+ }
} catch (af::exception& e) {
fprintf(stderr, "%s\n", e.what());
diff --git a/include/af/graphics.h b/include/af/graphics.h
index c6cf597..6c7061f 100644
--- a/include/af/graphics.h
+++ b/include/af/graphics.h
@@ -164,6 +164,28 @@ class AFAPI Window {
void hist(const array& X, const double minval, const double maxval, const char* const title=NULL);
/**
+ Renders the input arrays as a 3D surface plot to the window
+
+ \param[in] S is an \ref array with the z-axis data points
+ \param[in] title parameter is used when this function is called in grid mode
+
+ \note \p S should be a 2D array
+ */
+ void surface(const array& S, const char* const title);
+
+ /**
+ Renders the input arrays as a 3D surface plot to the window
+
+ \param[in] xVals is an \ref array with the x-axis data points
+ \param[in] yVals is an \ref array with the y-axis data points
+ \param[in] S is an \ref array with the z-axis data points
+ \param[in] title parameter is used when this function is called in grid mode
+
+ \note \p X and \p Y should be vectors or 2D arrays \p S should be s 2D array
+ */
+ void surface(const array& xVals, const array& yVals, const array& S, const char* const title);
+
+ /**
Setup grid layout for multiview mode in a window
\param[in] rows is number of rows you want to show in a window
@@ -306,15 +328,14 @@ AFAPI af_err af_draw_plot(const af_window wind, const af_array X, const af_array
C Interface wrapper for drawing an array as a plot
\param[in] wind is the window handle
- \param[in] X is an \ref af_array with the x-axis data points
- \param[in] Y is an \ref af_array with the y-axis data points
+ \param[in] P is an \ref af_array or matrix with the xyz-values of the points
\param[in] props is structure \ref af_cell that has the properties that are used
for the current rendering.
\return \ref AF_SUCCESS if rendering is successful, otherwise an appropriate error code
is returned.
- \note \p X and \p Y should be vectors.
+ \note \p P should be a 3n x 1 vector or one of a 3xn or nx3 matrices.
\ingroup gfx_func_draw
*/
@@ -340,6 +361,26 @@ AFAPI af_err af_draw_plot3(const af_window wind, const af_array P, const af_cell
AFAPI af_err af_draw_hist(const af_window wind, const af_array X, const double minval, const double maxval, const af_cell* const props);
/**
+ C Interface wrapper for drawing arrayis as a surface
+
+ \param[in] wind is the window handle
+ \param[in] xVals is an \ref af_array with the x-axis data points
+ \param[in] yVals is an \ref af_array with the y-axis data points
+ \param[in] S is an \ref af_array with the z-axis data points
+ \param[in] props is structure \ref af_cell that has the properties that are used
+ for the current rendering.
+
+ \return \ref AF_SUCCESS if rendering is successful, otherwise an appropriate error code
+ is returned.
+
+ \note \p X and \p Y should be vectors. \p S should be a 2D array
+
+ \ingroup gfx_func_draw
+*/
+
+af_err af_draw_surface(const af_window wind, const af_array xVals, const af_array yVals, const af_array S, const af_cell* const props);
+
+/**
C Interface wrapper for grid setup in a window
\param[in] wind is the window handle
diff --git a/src/api/c/graphics_common.cpp b/src/api/c/graphics_common.cpp
index eb03f25..df0ecf9 100644
--- a/src/api/c/graphics_common.cpp
+++ b/src/api/c/graphics_common.cpp
@@ -214,6 +214,25 @@ fg::Histogram* ForgeManager::getHistogram(int nBins, fg::dtype type)
return mHstMap[key];
}
+fg::Surface* ForgeManager::getSurface(int nX, int nY, fg::dtype type)
+{
+ /* nX * nY needs to fall in the range of [0, 2^48]
+ * for the ForgeManager to correctly retrieve
+ * the necessary Forge Plot object. So, this implementation
+ * is a limitation on how big of an plot graph can be rendered
+ * using arrayfire graphics funtionality */
+ assert(nX * nY <= 2ll<<48);
+ long long key = (((nX * nY) & _48BIT) << 48) | (type & _16BIT);
+
+ SfcMapIter iter = mSfcMap.find(key);
+ if (iter==mSfcMap.end()) {
+ fg::Surface* temp = new fg::Surface(nX, nY, type);
+ mSfcMap[key] = temp;
+ }
+
+ return mSfcMap[key];
+}
+
void ForgeManager::destroyResources()
{
/* clear all OpenGL resource objects (images, plots, histograms etc) first
diff --git a/src/api/c/graphics_common.hpp b/src/api/c/graphics_common.hpp
index 4f799b5..f649fe9 100644
--- a/src/api/c/graphics_common.hpp
+++ b/src/api/c/graphics_common.hpp
@@ -46,11 +46,13 @@ typedef std::map<long long, fg::Image*> ImageMap_t;
typedef std::map<long long, fg::Plot*> PlotMap_t;
typedef std::map<long long, fg::Histogram*> HistogramMap_t;
typedef std::map<long long, fg::Plot3*> Plot3Map_t;
+typedef std::map<long long, fg::Surface*> SurfaceMap_t;
typedef ImageMap_t::iterator ImgMapIter;
typedef PlotMap_t::iterator PltMapIter;
typedef Plot3Map_t::iterator Plt3MapIter;
typedef HistogramMap_t::iterator HstMapIter;
+typedef SurfaceMap_t::iterator SfcMapIter;
/**
* ForgeManager class follows a single pattern. Any user of this class, has
@@ -62,6 +64,7 @@ typedef HistogramMap_t::iterator HstMapIter;
* fg::Plot
* fg::Plot3
* fg::Histogram
+ * fg::Surface
* */
class ForgeManager
{
@@ -70,6 +73,7 @@ class ForgeManager
PlotMap_t mPltMap;
Plot3Map_t mPlt3Map;
HistogramMap_t mHstMap;
+ SurfaceMap_t mSfcMap;
public:
static ForgeManager& getInstance();
@@ -81,6 +85,7 @@ class ForgeManager
fg::Plot* getPlot(int nPoints, fg::dtype type);
fg::Plot3* getPlot3(int nPoints, fg::dtype type);
fg::Histogram* getHistogram(int nBins, fg::dtype type);
+ fg::Surface* getSurface(int nX, int nY, fg::dtype type);
protected:
ForgeManager() {}
diff --git a/src/api/c/surface.cpp b/src/api/c/surface.cpp
new file mode 100644
index 0000000..d1ae00c
--- /dev/null
+++ b/src/api/c/surface.cpp
@@ -0,0 +1,133 @@
+/*******************************************************
+ * Copyright (c) 2014, ArrayFire
+ * All rights reserved.
+ *
+ * This file is distributed under 3-clause BSD license.
+ * The complete license agreement can be obtained at:
+ * http://arrayfire.com/licenses/BSD-3-Clause
+ ********************************************************/
+
+#include <af/graphics.h>
+#include <af/image.h>
+
+#include <ArrayInfo.hpp>
+#include <graphics_common.hpp>
+#include <err_common.hpp>
+#include <backend.hpp>
+#include <surface.hpp>
+#include <reduce.hpp>
+#include <join.hpp>
+#include <tile.hpp>
+#include <reorder.hpp>
+#include <handle.hpp>
+
+using af::dim4;
+using namespace detail;
+
+#if defined(WITH_GRAPHICS)
+using namespace graphics;
+
+template<typename T>
+fg::Surface* setup_surface(const af_array xVals, const af_array yVals, const af_array zVals)
+{
+ Array<T> xIn = getArray<T>(xVals);
+ Array<T> yIn = getArray<T>(yVals);
+ Array<T> zIn = getArray<T>(zVals);
+
+ T xmax = reduce_all<af_max_t, T, T>(xIn);
+ T xmin = reduce_all<af_min_t, T, T>(xIn);
+ T ymax = reduce_all<af_max_t, T, T>(yIn);
+ T ymin = reduce_all<af_min_t, T, T>(yIn);
+ T zmax = reduce_all<af_max_t, T, T>(zIn);
+ T zmin = reduce_all<af_min_t, T, T>(zIn);
+
+ ArrayInfo Xinfo = getInfo(xVals);
+ ArrayInfo Yinfo = getInfo(yVals);
+ ArrayInfo Zinfo = getInfo(zVals);
+
+ af::dim4 X_dims = Xinfo.dims();
+ af::dim4 Y_dims = Yinfo.dims();
+ af::dim4 Z_dims = Zinfo.dims();
+
+ dim4 rdims(1, 0, 2, 3);
+ dim4 x_tdims(1, Y_dims[0], 1, 1);
+ dim4 y_tdims(1, X_dims[0], 1, 1);
+ if(Xinfo.isVector()){
+ xIn = tile(xIn, x_tdims);
+ yIn = tile(yIn, y_tdims);
+ yIn = reorder(yIn, rdims);
+ }
+
+ xIn.modDims(xIn.elements());
+ yIn.modDims(yIn.elements());
+ zIn.modDims(zIn.elements());
+ Array<T> Z = join(1, join(1, xIn, yIn), zIn);
+ Z = reorder(Z, rdims);
+ Z.modDims(Z.elements());
+
+ ForgeManager& fgMngr = ForgeManager::getInstance();
+ fg::Surface* surface = fgMngr.getSurface(Z_dims[0], Z_dims[1], getGLType<T>());
+ surface->setColor(1.0, 0.0, 0.0);
+ surface->setAxesLimits(xmax, xmin, ymax, ymin, zmax, zmin);
+ surface->setAxesTitles("X Axis", "Y Axis", "Z Axis");
+
+ copy_surface<T>(Z, surface);
+
+ return surface;
+}
+#endif
+
+af_err af_draw_surface(const af_window wind, const af_array xVals, const af_array yVals, const af_array S, const af_cell* const props)
+{
+#if defined(WITH_GRAPHICS)
+ if(wind==0) {
+ std::cerr<<"Not a valid window"<<std::endl;
+ return AF_SUCCESS;
+ }
+
+ try {
+ ArrayInfo Xinfo = getInfo(xVals);
+ af::dim4 X_dims = Xinfo.dims();
+ af_dtype Xtype = Xinfo.getType();
+
+ ArrayInfo Yinfo = getInfo(yVals);
+ af::dim4 Y_dims = Yinfo.dims();
+ af_dtype Ytype = Yinfo.getType();
+
+ ArrayInfo Sinfo = getInfo(S);
+ af::dim4 S_dims = Sinfo.dims();
+ af_dtype Stype = Sinfo.getType();
+
+ TYPE_ASSERT(Xtype == Ytype);
+ TYPE_ASSERT(Ytype == Stype);
+
+ if(!Yinfo.isVector()){
+ DIM_ASSERT(1, X_dims == Y_dims);
+ DIM_ASSERT(3, Y_dims == S_dims);
+ }else{
+ DIM_ASSERT(3, ( X_dims[0] * Y_dims[0] == Sinfo.elements()));
+ }
+
+ fg::Window* window = reinterpret_cast<fg::Window*>(wind);
+ window->makeCurrent();
+ fg::Surface* surface = NULL;
+
+ switch(Xtype) {
+ case f32: surface = setup_surface<float >(xVals, yVals , S); break;
+ case s32: surface = setup_surface<int >(xVals, yVals , S); break;
+ case u32: surface = setup_surface<uint >(xVals, yVals , S); break;
+ case u8 : surface = setup_surface<uchar >(xVals, yVals , S); break;
+ default: TYPE_ERROR(1, Xtype);
+ }
+
+ if (props->col>-1 && props->row>-1)
+ window->draw(props->col, props->row, *surface, props->title);
+ else
+ window->draw(*surface);
+ }
+ CATCHALL;
+ return AF_SUCCESS;
+#else
+ return AF_ERR_NO_GFX;
+#endif
+}
diff --git a/src/api/cpp/graphics.cpp b/src/api/cpp/graphics.cpp
index c9ff571..b748019 100644
--- a/src/api/cpp/graphics.cpp
+++ b/src/api/cpp/graphics.cpp
@@ -92,6 +92,20 @@ void Window::hist(const array& X, const double minval, const double maxval, cons
AF_THROW(af_draw_hist(get(), X.get(), minval, maxval, &temp));
}
+void Window::surface(const array& S, const char* const title){
+ //TODO: fix offset on forge?
+ af::array xVals = seq(0, S.dims(0)-1);
+ af::array yVals = seq(0, S.dims(1)-1);
+ af_cell temp{_r, _c, title, AF_COLORMAP_DEFAULT};
+ AF_THROW(af_draw_surface(get(), xVals.get(), yVals.get(), S.get(), &temp));
+}
+
+void Window::surface(const array& xVals, const array& yVals, const array& S, const char* const title)
+{
+ af_cell temp{_r, _c, title, AF_COLORMAP_DEFAULT};
+ AF_THROW(af_draw_surface(get(), xVals.get(), yVals.get(), S.get(), &temp));
+}
+
void Window::grid(const int rows, const int cols)
{
AF_THROW(af_grid(get(), rows, cols));
diff --git a/src/api/unified/graphics.cpp b/src/api/unified/graphics.cpp
index 27afb81..ca74f02 100644
--- a/src/api/unified/graphics.cpp
+++ b/src/api/unified/graphics.cpp
@@ -52,6 +52,11 @@ af_err af_draw_hist(const af_window wind, const af_array X, const double minval,
return CALL(wind, X, minval, maxval, props);
}
+af_err af_draw_surface(const af_window wind, const af_array xVals, const af_array yVals, const af_array S, const af_cell* const props)
+{
+ return CALL(wind, xVals, yVals, S, props);
+}
+
af_err af_grid(const af_window wind, const int rows, const int cols)
{
return CALL(wind, rows, cols);
diff --git a/src/backend/cpu/surface.cpp b/src/backend/cpu/surface.cpp
new file mode 100644
index 0000000..39f375a
--- /dev/null
+++ b/src/backend/cpu/surface.cpp
@@ -0,0 +1,48 @@
+/*******************************************************
+ * Copyright (c) 2014, ArrayFire
+ * All rights reserved.
+ *
+ * This file is distributed under 3-clause BSD license.
+ * The complete license agreement can be obtained at:
+ * http://arrayfire.com/licenses/BSD-3-Clause
+ ********************************************************/
+
+#if defined(WITH_GRAPHICS)
+
+#include <Array.hpp>
+#include <surface.hpp>
+#include <err_cpu.hpp>
+#include <stdexcept>
+#include <graphics_common.hpp>
+#include <reduce.hpp>
+#include <memory.hpp>
+
+using af::dim4;
+
+namespace cpu
+{
+ template<typename T>
+ void copy_surface(const Array<T> &P, fg::Surface* surface)
+ {
+ CheckGL("Before CopyArrayToVBO");
+
+ glBindBuffer(GL_ARRAY_BUFFER, surface->vbo());
+ glBufferSubData(GL_ARRAY_BUFFER, 0, surface->size(), P.get());
+ glBindBuffer(GL_ARRAY_BUFFER, 0);
+
+ CheckGL("In CopyArrayToVBO");
+ }
+
+ #define INSTANTIATE(T) \
+ template void copy_surface<T>(const Array<T> &P, fg::Surface* surface);
+
+ INSTANTIATE(float)
+ INSTANTIATE(double)
+ INSTANTIATE(int)
+ INSTANTIATE(uint)
+ INSTANTIATE(uchar)
+ INSTANTIATE(short)
+ INSTANTIATE(ushort)
+}
+
+#endif // WITH_GRAPHICS
diff --git a/src/backend/opencl/plot.hpp b/src/backend/cpu/surface.hpp
similarity index 85%
copy from src/backend/opencl/plot.hpp
copy to src/backend/cpu/surface.hpp
index 582d02e..46a4c4b 100644
--- a/src/backend/opencl/plot.hpp
+++ b/src/backend/cpu/surface.hpp
@@ -12,12 +12,11 @@
#include <Array.hpp>
#include <graphics_common.hpp>
-namespace opencl
+namespace cpu
{
template<typename T>
- void copy_plot(const Array<T> &P, fg::Plot* plot);
+ void copy_surface(const Array<T> &P, fg::Surface* surface);
}
#endif
-
diff --git a/src/backend/cuda/surface.cu b/src/backend/cuda/surface.cu
new file mode 100644
index 0000000..cb8bf4e
--- /dev/null
+++ b/src/backend/cuda/surface.cu
@@ -0,0 +1,59 @@
+/*******************************************************
+ * Copyright (c) 2014, ArrayFire
+ * All rights reserved.
+ *
+ * This file is distributed under 3-clause BSD license.
+ * The complete license agreement can be obtained at:
+ * http://arrayfire.com/licenses/BSD-3-Clause
+ ********************************************************/
+
+#if defined (WITH_GRAPHICS)
+
+#include <interopManager.hpp>
+#include <Array.hpp>
+#include <surface.hpp>
+#include <err_cuda.hpp>
+#include <debug_cuda.hpp>
+#include <join.hpp>
+#include <reduce.hpp>
+#include <reorder.hpp>
+
+using af::dim4;
+
+namespace cuda
+{
+
+template<typename T>
+void copy_surface(const Array<T> &P, fg::Surface* surface)
+{
+ const T *d_P = P.get();
+
+ InteropManager& intrpMngr = InteropManager::getInstance();
+
+ cudaGraphicsResource *cudaVBOResource = intrpMngr.getBufferResource(surface);
+ // Map resource. Copy data to VBO. Unmap resource.
+ size_t num_bytes = surface->size();
+ T* d_vbo = NULL;
+ cudaGraphicsMapResources(1, &cudaVBOResource, 0);
+ cudaGraphicsResourceGetMappedPointer((void **)&d_vbo, &num_bytes, cudaVBOResource);
+ cudaMemcpyAsync(d_vbo, d_P, num_bytes, cudaMemcpyDeviceToDevice,
+ cuda::getStream(cuda::getActiveDeviceId()));
+ cudaGraphicsUnmapResources(1, &cudaVBOResource, 0);
+
+ CheckGL("After cuda resource copy");
+
+ POST_LAUNCH_CHECK();
+}
+
+#define INSTANTIATE(T) \
+ template void copy_surface<T>(const Array<T> &P, fg::Surface* surface);
+
+INSTANTIATE(float)
+INSTANTIATE(double)
+INSTANTIATE(int)
+INSTANTIATE(uint)
+INSTANTIATE(uchar)
+
+}
+
+#endif // WITH_GRAPHICS
diff --git a/src/backend/opencl/plot.hpp b/src/backend/cuda/surface.hpp
similarity index 85%
copy from src/backend/opencl/plot.hpp
copy to src/backend/cuda/surface.hpp
index 582d02e..d701983 100644
--- a/src/backend/opencl/plot.hpp
+++ b/src/backend/cuda/surface.hpp
@@ -12,12 +12,11 @@
#include <Array.hpp>
#include <graphics_common.hpp>
-namespace opencl
+namespace cuda
{
template<typename T>
- void copy_plot(const Array<T> &P, fg::Plot* plot);
+ void copy_surface(const Array<T> &P, fg::Surface* surface);
}
#endif
-
diff --git a/src/backend/opencl/interopManager.cpp b/src/backend/opencl/interopManager.cpp
index 099adb1..89487ec 100644
--- a/src/backend/opencl/interopManager.cpp
+++ b/src/backend/opencl/interopManager.cpp
@@ -83,6 +83,18 @@ cl::Buffer* InteropManager::getBufferResource(const fg::Histogram* hist)
return interop_maps[device][key];
}
+cl::Buffer* InteropManager::getBufferResource(const fg::Surface* surface)
+{
+ void * key = (void*)surface;
+ int device = getActiveDeviceId();
+ iter_t iter = interop_maps[device].find(key);
+
+ if (iter == interop_maps[device].end())
+ interop_maps[device][key] = new cl::BufferGL(getContext(), CL_MEM_WRITE_ONLY, surface->vbo(), NULL);
+
+ return interop_maps[device][key];
+}
+
}
#endif
diff --git a/src/backend/opencl/interopManager.hpp b/src/backend/opencl/interopManager.hpp
index 7bd530c..c7a2c25 100644
--- a/src/backend/opencl/interopManager.hpp
+++ b/src/backend/opencl/interopManager.hpp
@@ -32,6 +32,7 @@ class InteropManager
cl::Buffer* getBufferResource(const fg::Plot* plot);
cl::Buffer* getBufferResource(const fg::Plot3* plot3);
cl::Buffer* getBufferResource(const fg::Histogram* hist);
+ cl::Buffer* getBufferResource(const fg::Surface* surface);
protected:
InteropManager() {}
diff --git a/src/backend/opencl/plot.hpp b/src/backend/opencl/plot.hpp
index 582d02e..c195694 100644
--- a/src/backend/opencl/plot.hpp
+++ b/src/backend/opencl/plot.hpp
@@ -20,4 +20,3 @@ namespace opencl
#endif
-
diff --git a/src/backend/opencl/surface.cpp b/src/backend/opencl/surface.cpp
new file mode 100644
index 0000000..587ad38
--- /dev/null
+++ b/src/backend/opencl/surface.cpp
@@ -0,0 +1,73 @@
+/*******************************************************
+ * Copyright (c) 2014, ArrayFire
+ * All rights reserved.
+ *
+ * This file is distributed under 3-clause BSD license.
+ * The complete license agreement can be obtained at:
+ * http://arrayfire.com/licenses/BSD-3-Clause
+ ********************************************************/
+
+#if defined (WITH_GRAPHICS)
+
+#include <interopManager.hpp>
+#include <Array.hpp>
+#include <surface.hpp>
+#include <err_opencl.hpp>
+#include <debug_opencl.hpp>
+#include <join.hpp>
+#include <reduce.hpp>
+#include <reorder.hpp>
+
+using af::dim4;
+
+namespace opencl
+{
+
+template<typename T>
+void copy_surface(const Array<T> &P, fg::Surface* surface)
+{
+ if (isGLSharingSupported()) {
+ CheckGL("Begin OpenCL resource copy");
+ const cl::Buffer *d_P = P.get();
+ size_t bytes = surface->size();
+
+ InteropManager& intrpMngr = InteropManager::getInstance();
+
+ cl::Buffer *clPBOResource = intrpMngr.getBufferResource(surface);
+
+ std::vector<cl::Memory> shared_objects;
+ shared_objects.push_back(*clPBOResource);
+
+ glFinish();
+ getQueue().enqueueAcquireGLObjects(&shared_objects);
+ getQueue().enqueueCopyBuffer(*d_P, *clPBOResource, 0, 0, bytes, NULL, NULL);
+ getQueue().finish();
+ getQueue().enqueueReleaseGLObjects(&shared_objects);
+
+ CL_DEBUG_FINISH(getQueue());
+ CheckGL("End OpenCL resource copy");
+ } else {
+ CheckGL("Begin OpenCL fallback-resource copy");
+ glBindBuffer(GL_ARRAY_BUFFER, surface->vbo());
+ GLubyte* ptr = (GLubyte*)glMapBuffer(GL_ARRAY_BUFFER, GL_WRITE_ONLY);
+ if (ptr) {
+ getQueue().enqueueReadBuffer(*P.get(), CL_TRUE, 0, surface->size(), ptr);
+ glUnmapBuffer(GL_ARRAY_BUFFER);
+ }
+ glBindBuffer(GL_ARRAY_BUFFER, 0);
+ CheckGL("End OpenCL fallback-resource copy");
+ }
+}
+
+#define INSTANTIATE(T) \
+ template void copy_surface<T>(const Array<T> &P, fg::Surface* surface);
+
+INSTANTIATE(float)
+INSTANTIATE(double)
+INSTANTIATE(int)
+INSTANTIATE(uint)
+INSTANTIATE(uchar)
+
+}
+
+#endif // WITH_GRAPHICS
diff --git a/src/backend/opencl/plot.hpp b/src/backend/opencl/surface.hpp
similarity index 88%
copy from src/backend/opencl/plot.hpp
copy to src/backend/opencl/surface.hpp
index 582d02e..15079f0 100644
--- a/src/backend/opencl/plot.hpp
+++ b/src/backend/opencl/surface.hpp
@@ -15,7 +15,7 @@
namespace opencl
{
template<typename T>
- void copy_plot(const Array<T> &P, fg::Plot* plot);
+ void copy_surface(const Array<T> &P, fg::Surface* surface);
}
#endif
--
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/debian-science/packages/arrayfire.git
More information about the debian-science-commits
mailing list