[arrayfire] 132/248: 3D line plot feature
Ghislain Vaillant
ghisvail-guest at moszumanska.debian.org
Tue Nov 17 15:54:15 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 2b4ed2512629b5baa509d03bec432e3689131950
Author: syurkevi <stefan at arrayfire.com>
Date: Fri Oct 16 19:52:56 2015 -0400
3D line plot feature
Also adds example
---
CMakeModules/build_forge.cmake | 2 +-
examples/graphics/plot3.cpp | 56 +++++++++++++++++
include/af/graphics.h | 28 +++++++++
src/api/c/graphics_common.cpp | 37 +++++++++---
src/api/c/graphics_common.hpp | 13 ++--
src/api/c/hist.cpp | 3 +-
src/api/c/image.cpp | 4 +-
src/api/c/plot.cpp | 3 +-
src/api/c/plot3.cpp | 111 ++++++++++++++++++++++++++++++++++
src/api/cpp/graphics.cpp | 7 +++
src/api/unified/graphics.cpp | 5 ++
src/backend/cpu/plot3.cpp | 48 +++++++++++++++
src/backend/cpu/plot3.hpp | 22 +++++++
src/backend/cuda/interopManager.cu | 17 ++++++
src/backend/cuda/interopManager.hpp | 1 +
src/backend/cuda/plot3.cu | 59 ++++++++++++++++++
src/backend/cuda/plot3.hpp | 22 +++++++
src/backend/opencl/interopManager.cpp | 12 ++++
src/backend/opencl/interopManager.hpp | 1 +
src/backend/opencl/plot3.cpp | 70 +++++++++++++++++++++
src/backend/opencl/plot3.hpp | 22 +++++++
21 files changed, 523 insertions(+), 20 deletions(-)
diff --git a/CMakeModules/build_forge.cmake b/CMakeModules/build_forge.cmake
index 5784b76..62ea5cb 100644
--- a/CMakeModules/build_forge.cmake
+++ b/CMakeModules/build_forge.cmake
@@ -22,7 +22,7 @@ ENDIF()
ExternalProject_Add(
forge-ext
GIT_REPOSITORY https://github.com/arrayfire/forge.git
- GIT_TAG af3.1.2
+ GIT_TAG 50959f2f04592d23d5207623c43e675bc4a648dc
PREFIX "${prefix}"
INSTALL_DIR "${prefix}"
UPDATE_COMMAND ""
diff --git a/examples/graphics/plot3.cpp b/examples/graphics/plot3.cpp
new file mode 100644
index 0000000..40bd6d4
--- /dev/null
+++ b/examples/graphics/plot3.cpp
@@ -0,0 +1,56 @@
+/*******************************************************
+ * 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 <arrayfire.h>
+#include <cstdio>
+#include <math.h>
+
+using namespace af;
+
+static const int ITERATIONS = 200;
+static const float PRECISION = 1.0f/ITERATIONS;
+
+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());
+
+ } catch (af::exception& e) {
+ fprintf(stderr, "%s\n", e.what());
+ throw;
+ }
+
+ #ifdef WIN32 // pause in Windows
+ if (!(argc == 2 && argv[1][0] == '-')) {
+ printf("hit [enter]...");
+ fflush(stdout);
+ getchar();
+ }
+ #endif
+ return 0;
+}
+
diff --git a/include/af/graphics.h b/include/af/graphics.h
index 1fd9108..c6cf597 100644
--- a/include/af/graphics.h
+++ b/include/af/graphics.h
@@ -131,6 +131,15 @@ class AFAPI Window {
void image(const array& in, const char* title=NULL);
/**
+ Renders the input array as an 3d line plot to the window
+
+ \param[in] in is an \ref array
+ \param[in] title parameter is used when this function is called in grid mode
+
+ \note \p in should be 1d array of size 3n or 2d array with (3 x n) or (n x 3) channels.
+ */
+ void plot3(const array& in, const char* title=NULL);
+ /**
Renders the input arrays as a 2D plot to the window
\param[in] X is an \ref array with the x-axis data points
@@ -139,6 +148,7 @@ class AFAPI Window {
\note \p X and \p Y should be vectors.
*/
+
void plot(const array& X, const array& Y, const char* const title=NULL);
/**
@@ -293,6 +303,24 @@ AFAPI af_err af_draw_image(const af_window wind, const af_array in, const af_cel
AFAPI af_err af_draw_plot(const af_window wind, const af_array X, const af_array Y, const af_cell* const props);
/**
+ 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] 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.
+
+ \ingroup gfx_func_draw
+*/
+AFAPI af_err af_draw_plot3(const af_window wind, const af_array P, const af_cell* const props);
+
+/**
C Interface wrapper for drawing an array as a histogram
\param[in] wind is the window handle
diff --git a/src/api/c/graphics_common.cpp b/src/api/c/graphics_common.cpp
index bec27fe..eb03f25 100644
--- a/src/api/c/graphics_common.cpp
+++ b/src/api/c/graphics_common.cpp
@@ -20,13 +20,13 @@ template<typename T>
GLenum getGLType() { return GL_FLOAT; }
#define INSTANTIATE_GET_FG_TYPE(T, ForgeEnum)\
- template<> fg::FGType getGLType<T>() { return ForgeEnum; }
+ template<> fg::dtype getGLType<T>() { return ForgeEnum; }
-INSTANTIATE_GET_FG_TYPE(float, fg::FG_FLOAT);
-INSTANTIATE_GET_FG_TYPE(int , fg::FG_INT);
-INSTANTIATE_GET_FG_TYPE(unsigned, fg::FG_UNSIGNED_INT);
-INSTANTIATE_GET_FG_TYPE(char, fg::FG_BYTE);
-INSTANTIATE_GET_FG_TYPE(unsigned char, fg::FG_UNSIGNED_BYTE);
+INSTANTIATE_GET_FG_TYPE(float, fg::f32);
+INSTANTIATE_GET_FG_TYPE(int , fg::s32);
+INSTANTIATE_GET_FG_TYPE(unsigned, fg::u32);
+INSTANTIATE_GET_FG_TYPE(char, fg::s8);
+INSTANTIATE_GET_FG_TYPE(unsigned char, fg::u8);
GLenum glErrorSkip(const char *msg, const char* file, int line)
{
@@ -136,7 +136,7 @@ fg::Window* ForgeManager::getMainWindow(const bool dontCreate)
return wnd;
}
-fg::Image* ForgeManager::getImage(int w, int h, fg::ColorMode mode, fg::FGType type)
+fg::Image* ForgeManager::getImage(int w, int h, fg::ChannelFormat mode, fg::dtype type)
{
/* w, h needs to fall in the range of [0, 2^16]
* for the ForgeManager to correctly retrieve
@@ -157,7 +157,7 @@ fg::Image* ForgeManager::getImage(int w, int h, fg::ColorMode mode, fg::FGType t
return mImgMap[key];
}
-fg::Plot* ForgeManager::getPlot(int nPoints, fg::FGType type)
+fg::Plot* ForgeManager::getPlot(int nPoints, fg::dtype type)
{
/* nPoints needs to fall in the range of [0, 2^48]
* for the ForgeManager to correctly retrieve
@@ -176,7 +176,26 @@ fg::Plot* ForgeManager::getPlot(int nPoints, fg::FGType type)
return mPltMap[key];
}
-fg::Histogram* ForgeManager::getHistogram(int nBins, fg::FGType type)
+fg::Plot3* ForgeManager::getPlot3(int nPoints, fg::dtype type)
+{
+ /* nPoints 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(nPoints <= 2ll<<48);
+ long long key = ((nPoints & _48BIT) << 48) | (type & _16BIT);
+
+ Plt3MapIter iter = mPlt3Map.find(key);
+ if (iter==mPlt3Map.end()) {
+ fg::Plot3* temp = new fg::Plot3(nPoints, type);
+ mPlt3Map[key] = temp;
+ }
+
+ return mPlt3Map[key];
+}
+
+fg::Histogram* ForgeManager::getHistogram(int nBins, fg::dtype type)
{
/* nBins needs to fall in the range of [0, 2^48]
* for the ForgeManager to correctly retrieve
diff --git a/src/api/c/graphics_common.hpp b/src/api/c/graphics_common.hpp
index ac6f4c0..4f799b5 100644
--- a/src/api/c/graphics_common.hpp
+++ b/src/api/c/graphics_common.hpp
@@ -18,7 +18,7 @@
// default to f32(float) type
template<typename T>
-fg::FGType getGLType();
+fg::dtype getGLType();
// Print for OpenGL errors
// Returns 1 if an OpenGL error occurred, 0 otherwise.
@@ -45,9 +45,11 @@ static const long long _48BIT = 0x0000FFFFFFFFFFFF;
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 ImageMap_t::iterator ImgMapIter;
typedef PlotMap_t::iterator PltMapIter;
+typedef Plot3Map_t::iterator Plt3MapIter;
typedef HistogramMap_t::iterator HstMapIter;
/**
@@ -58,6 +60,7 @@ typedef HistogramMap_t::iterator HstMapIter;
* Renderables:
* fg::Image
* fg::Plot
+ * fg::Plot3
* fg::Histogram
* */
class ForgeManager
@@ -65,6 +68,7 @@ class ForgeManager
private:
ImageMap_t mImgMap;
PlotMap_t mPltMap;
+ Plot3Map_t mPlt3Map;
HistogramMap_t mHstMap;
public:
@@ -73,9 +77,10 @@ class ForgeManager
fg::Font* getFont(const bool dontCreate=false);
fg::Window* getMainWindow(const bool dontCreate=false);
- fg::Image* getImage(int w, int h, fg::ColorMode mode, fg::FGType type);
- fg::Plot* getPlot(int nPoints, fg::FGType type);
- fg::Histogram* getHistogram(int nBins, fg::FGType type);
+ fg::Image* getImage(int w, int h, fg::ChannelFormat mode, fg::dtype type);
+ fg::Plot* getPlot(int nPoints, fg::dtype type);
+ fg::Plot3* getPlot3(int nPoints, fg::dtype type);
+ fg::Histogram* getHistogram(int nBins, fg::dtype type);
protected:
ForgeManager() {}
diff --git a/src/api/c/hist.cpp b/src/api/c/hist.cpp
index 8cf1ac1..e4e3eb6 100644
--- a/src/api/c/hist.cpp
+++ b/src/api/c/hist.cpp
@@ -39,8 +39,7 @@ fg::Histogram* setup_histogram(const af_array in, const double minval, const dou
/* set x axis limits to maximum and minimum values of data
* and y axis limits to range [0, nBins]*/
hist->setAxesLimits(maxval, minval, double(freqMax), 0.0f);
- hist->setXAxisTitle("Bins");
- hist->setYAxisTitle("Frequency");
+ hist->setAxesTitles("Bins", "Frequency");
copy_histogram<T>(histogramInput, hist);
diff --git a/src/api/c/image.cpp b/src/api/c/image.cpp
index c59b31b..134e4c2 100644
--- a/src/api/c/image.cpp
+++ b/src/api/c/image.cpp
@@ -66,7 +66,7 @@ static fg::Image* convert_and_copy_image(const af_array in)
ForgeManager& fgMngr = ForgeManager::getInstance();
- fg::Image* ret_val = fgMngr.getImage(inDims[1], inDims[0], (fg::ColorMode)inDims[2], getGLType<T>());
+ fg::Image* ret_val = fgMngr.getImage(inDims[1], inDims[0], (fg::ChannelFormat)inDims[2], getGLType<T>());
copy_image<T>(normalizePerType<T>(imgData), ret_val);
@@ -233,7 +233,7 @@ af_err af_show(const af_window wind)
try {
fg::Window* wnd = reinterpret_cast<fg::Window*>(wind);
- wnd->draw();
+ wnd->swapBuffers();
}
CATCHALL;
return AF_SUCCESS;
diff --git a/src/api/c/plot.cpp b/src/api/c/plot.cpp
index 0723a30..8ac44e7 100644
--- a/src/api/c/plot.cpp
+++ b/src/api/c/plot.cpp
@@ -49,8 +49,7 @@ fg::Plot* setup_plot(const af_array X, const af_array Y)
fg::Plot* plot = fgMngr.getPlot(X_dims.elements(), getGLType<T>());
plot->setColor(1.0, 0.0, 0.0);
plot->setAxesLimits(xmax, xmin, ymax, ymin);
- plot->setXAxisTitle("X Axis");
- plot->setYAxisTitle("Y Axis");
+ plot->setAxesTitles("X Axis", "Y Axis");
copy_plot<T>(P, plot);
diff --git a/src/api/c/plot3.cpp b/src/api/c/plot3.cpp
new file mode 100644
index 0000000..91bd412
--- /dev/null
+++ b/src/api/c/plot3.cpp
@@ -0,0 +1,111 @@
+/*******************************************************
+ * 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 <plot3.hpp>
+#include <reduce.hpp>
+#include <join.hpp>
+#include <transpose.hpp>
+#include <reorder.hpp>
+#include <handle.hpp>
+#include <af/data.h>
+#include <iostream>
+
+using af::dim4;
+using namespace detail;
+
+#if defined(WITH_GRAPHICS)
+using namespace graphics;
+
+template<typename T>
+fg::Plot3* setup_plot3(const af_array P)
+{
+ Array<T> pIn = getArray<T>(P);
+ ArrayInfo Pinfo = getInfo(P);
+ af::dim4 P_dims = Pinfo.dims();
+
+ DIM_ASSERT(0, Pinfo.ndims() == 1 || Pinfo.ndims() == 2);
+ DIM_ASSERT(0, (P_dims[0] == 3 || P_dims[1] == 3) ||
+ (Pinfo.isVector() && P_dims[0]%3 == 0));
+
+ if(Pinfo.isVector()){
+ dim4 rdims(P_dims.elements()/3, 3, 1, 1);
+ pIn.modDims(rdims);
+ P_dims = pIn.dims();
+ }
+
+ T max[3], min[3];
+ if(P_dims[0] == 3) {
+ af_get_data_ptr(max, getHandle(reduce<af_max_t, T, T>(pIn, 1)));
+ af_get_data_ptr(min, getHandle(reduce<af_min_t, T, T>(pIn, 1)));
+ }
+
+ if(P_dims[1] == 3) {
+ af_get_data_ptr(max, getHandle(reduce<af_max_t, T, T>(pIn, 0)));
+ af_get_data_ptr(min, getHandle(reduce<af_min_t, T, T>(pIn, 0)));
+ }
+
+ ForgeManager& fgMngr = ForgeManager::getInstance();
+ fg::Plot3* plot3 = fgMngr.getPlot3(P_dims.elements()/3, getGLType<T>());
+ plot3->setColor(1.0, 0.0, 0.0);
+ plot3->setAxesLimits(max[0], min[0],
+ max[1], min[1],
+ max[2], min[2]);
+ plot3->setAxesTitles("X Axis", "Y Axis", "Z Axis");
+
+ if(P_dims[1] == 3){
+ pIn = transpose(pIn, false);
+ }
+ copy_plot3<T>(pIn, plot3);
+
+ return plot3;
+}
+#endif
+
+af_err af_draw_plot3(const af_window wind, const af_array P, 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 Pinfo = getInfo(P);
+ af_dtype Ptype = Pinfo.getType();
+
+ fg::Window* window = reinterpret_cast<fg::Window*>(wind);
+ window->makeCurrent();
+ fg::Plot3* plot3 = NULL;
+
+ switch(Ptype) {
+ case f32: plot3 = setup_plot3<float>(P); break;
+ case s32: plot3 = setup_plot3<int >(P); break;
+ case u32: plot3 = setup_plot3<uint >(P); break;
+ case u8 : plot3 = setup_plot3<uchar>(P); break;
+ default: TYPE_ERROR(1, Ptype);
+ }
+
+ if (props->col>-1 && props->row>-1)
+ window->draw(props->col, props->row, *plot3, props->title);
+ else
+ window->draw(*plot3);
+ }
+ 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 1272d2f..c9ff571 100644
--- a/src/api/cpp/graphics.cpp
+++ b/src/api/cpp/graphics.cpp
@@ -79,6 +79,13 @@ void Window::plot(const array& X, const array& Y, const char* const title)
AF_THROW(af_draw_plot(get(), X.get(), Y.get(), &temp));
}
+void Window::plot3(const array& P, const char* const title)
+{
+ af_cell temp{_r, _c, title, AF_COLORMAP_DEFAULT};
+ P.eval();
+ AF_THROW(af_draw_plot3(get(), P.get(), &temp));
+}
+
void Window::hist(const array& X, const double minval, const double maxval, const char* const title)
{
af_cell temp{_r, _c, title, AF_COLORMAP_DEFAULT};
diff --git a/src/api/unified/graphics.cpp b/src/api/unified/graphics.cpp
index 61ed4a9..27afb81 100644
--- a/src/api/unified/graphics.cpp
+++ b/src/api/unified/graphics.cpp
@@ -42,6 +42,11 @@ af_err af_draw_plot(const af_window wind, const af_array X, const af_array Y, co
return CALL(wind, X, Y, props);
}
+af_err af_draw_plot3(const af_window wind, const af_array P, const af_cell* const props)
+{
+ return CALL(wind, P, props);
+}
+
af_err af_draw_hist(const af_window wind, const af_array X, const double minval, const double maxval, const af_cell* const props)
{
return CALL(wind, X, minval, maxval, props);
diff --git a/src/backend/cpu/plot3.cpp b/src/backend/cpu/plot3.cpp
new file mode 100644
index 0000000..c0e26aa
--- /dev/null
+++ b/src/backend/cpu/plot3.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 <plot3.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_plot3(const Array<T> &P, fg::Plot3* plot3)
+ {
+ CheckGL("Before CopyArrayToVBO");
+
+ glBindBuffer(GL_ARRAY_BUFFER, plot3->vbo());
+ glBufferSubData(GL_ARRAY_BUFFER, 0, plot3->size(), P.get());
+ glBindBuffer(GL_ARRAY_BUFFER, 0);
+
+ CheckGL("In CopyArrayToVBO");
+ }
+
+ #define INSTANTIATE(T) \
+ template void copy_plot3<T>(const Array<T> &P, fg::Plot3* plot3);
+
+ INSTANTIATE(float)
+ INSTANTIATE(double)
+ INSTANTIATE(int)
+ INSTANTIATE(uint)
+ INSTANTIATE(uchar)
+ INSTANTIATE(short)
+ INSTANTIATE(ushort)
+}
+
+#endif // WITH_GRAPHICS
diff --git a/src/backend/cpu/plot3.hpp b/src/backend/cpu/plot3.hpp
new file mode 100644
index 0000000..1e5c97a
--- /dev/null
+++ b/src/backend/cpu/plot3.hpp
@@ -0,0 +1,22 @@
+/*******************************************************
+ * 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 <graphics_common.hpp>
+
+namespace cpu
+{
+ template<typename T>
+ void copy_plot3(const Array<T> &P, fg::Plot3* plot3);
+}
+
+#endif
+
diff --git a/src/backend/cuda/interopManager.cu b/src/backend/cuda/interopManager.cu
index fe13b46..dcee681 100644
--- a/src/backend/cuda/interopManager.cu
+++ b/src/backend/cuda/interopManager.cu
@@ -82,6 +82,23 @@ cudaGraphicsResource* InteropManager::getBufferResource(const fg::Plot* key)
return interop_maps[device][key_value];
}
+cudaGraphicsResource* InteropManager::getBufferResource(const fg::Plot3* key)
+{
+ int device = getActiveDeviceId();
+ void* key_value = (void*)key;
+
+ iter_t iter = interop_maps[device].find(key_value);
+
+ if(interop_maps[device].find(key_value) == interop_maps[device].end()) {
+ cudaGraphicsResource *cudaVBOResource;
+ // Register VBO with CUDA
+ CUDA_CHECK(cudaGraphicsGLRegisterBuffer(&cudaVBOResource, key->vbo(), cudaGraphicsMapFlagsWriteDiscard));
+ interop_maps[device][key_value] = cudaVBOResource;
+ }
+
+ return interop_maps[device][key_value];
+}
+
cudaGraphicsResource* InteropManager::getBufferResource(const fg::Histogram* key)
{
int device = getActiveDeviceId();
diff --git a/src/backend/cuda/interopManager.hpp b/src/backend/cuda/interopManager.hpp
index f6d3904..6508b08 100644
--- a/src/backend/cuda/interopManager.hpp
+++ b/src/backend/cuda/interopManager.hpp
@@ -40,6 +40,7 @@ class InteropManager
~InteropManager();
cudaGraphicsResource* getBufferResource(const fg::Image* handle);
cudaGraphicsResource* getBufferResource(const fg::Plot* handle);
+ cudaGraphicsResource* getBufferResource(const fg::Plot3* handle);
cudaGraphicsResource* getBufferResource(const fg::Histogram* handle);
protected:
diff --git a/src/backend/cuda/plot3.cu b/src/backend/cuda/plot3.cu
new file mode 100644
index 0000000..2e00ba9
--- /dev/null
+++ b/src/backend/cuda/plot3.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 <plot3.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_plot3(const Array<T> &P, fg::Plot3* plot3)
+{
+ const T *d_P = P.get();
+
+ InteropManager& intrpMngr = InteropManager::getInstance();
+
+ cudaGraphicsResource *cudaVBOResource = intrpMngr.getBufferResource(plot3);
+ // Map resource. Copy data to VBO. Unmap resource.
+ size_t num_bytes = plot3->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_plot3<T>(const Array<T> &P, fg::Plot3* plot3);
+
+INSTANTIATE(float)
+INSTANTIATE(double)
+INSTANTIATE(int)
+INSTANTIATE(uint)
+INSTANTIATE(uchar)
+
+}
+
+#endif // WITH_GRAPHICS
diff --git a/src/backend/cuda/plot3.hpp b/src/backend/cuda/plot3.hpp
new file mode 100644
index 0000000..3badb33
--- /dev/null
+++ b/src/backend/cuda/plot3.hpp
@@ -0,0 +1,22 @@
+/*******************************************************
+ * 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 <graphics_common.hpp>
+
+namespace cuda
+{
+ template<typename T>
+ void copy_plot3(const Array<T> &P, fg::Plot3* plot3);
+}
+
+#endif
+
diff --git a/src/backend/opencl/interopManager.cpp b/src/backend/opencl/interopManager.cpp
index 2b6fda0..099adb1 100644
--- a/src/backend/opencl/interopManager.cpp
+++ b/src/backend/opencl/interopManager.cpp
@@ -59,6 +59,18 @@ cl::Buffer* InteropManager::getBufferResource(const fg::Plot* plot)
return interop_maps[device][key];
}
+cl::Buffer* InteropManager::getBufferResource(const fg::Plot3* plot3)
+{
+ void * key = (void*)plot3;
+ 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, plot3->vbo(), NULL);
+
+ return interop_maps[device][key];
+}
+
cl::Buffer* InteropManager::getBufferResource(const fg::Histogram* hist)
{
void * key = (void*)hist;
diff --git a/src/backend/opencl/interopManager.hpp b/src/backend/opencl/interopManager.hpp
index 6af6d17..7bd530c 100644
--- a/src/backend/opencl/interopManager.hpp
+++ b/src/backend/opencl/interopManager.hpp
@@ -30,6 +30,7 @@ class InteropManager
~InteropManager();
cl::Buffer* getBufferResource(const fg::Image* image);
cl::Buffer* getBufferResource(const fg::Plot* plot);
+ cl::Buffer* getBufferResource(const fg::Plot3* plot3);
cl::Buffer* getBufferResource(const fg::Histogram* hist);
protected:
diff --git a/src/backend/opencl/plot3.cpp b/src/backend/opencl/plot3.cpp
new file mode 100644
index 0000000..9351498
--- /dev/null
+++ b/src/backend/opencl/plot3.cpp
@@ -0,0 +1,70 @@
+/*******************************************************
+ * 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 <plot3.hpp>
+#include <err_opencl.hpp>
+#include <debug_opencl.hpp>
+
+using af::dim4;
+
+namespace opencl
+{
+
+template<typename T>
+void copy_plot3(const Array<T> &P, fg::Plot3* plot3)
+{
+ if (isGLSharingSupported()) {
+ CheckGL("Begin OpenCL resource copy");
+ const cl::Buffer *d_P = P.get();
+ size_t bytes = plot3->size();
+
+ InteropManager& intrpMngr = InteropManager::getInstance();
+
+ cl::Buffer *clPBOResource = intrpMngr.getBufferResource(plot3);
+
+ 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, plot3->vbo());
+ GLubyte* ptr = (GLubyte*)glMapBuffer(GL_ARRAY_BUFFER, GL_WRITE_ONLY);
+ if (ptr) {
+ getQueue().enqueueReadBuffer(*P.get(), CL_TRUE, 0, plot3->size(), ptr);
+ glUnmapBuffer(GL_ARRAY_BUFFER);
+ }
+ glBindBuffer(GL_ARRAY_BUFFER, 0);
+ CheckGL("End OpenCL fallback-resource copy");
+ }
+}
+
+#define INSTANTIATE(T) \
+ template void copy_plot3<T>(const Array<T> &P, fg::Plot3* plot3);
+
+INSTANTIATE(float)
+INSTANTIATE(double)
+INSTANTIATE(int)
+INSTANTIATE(uint)
+INSTANTIATE(uchar)
+
+}
+
+#endif // WITH_GRAPHICS
diff --git a/src/backend/opencl/plot3.hpp b/src/backend/opencl/plot3.hpp
new file mode 100644
index 0000000..8609390
--- /dev/null
+++ b/src/backend/opencl/plot3.hpp
@@ -0,0 +1,22 @@
+/*******************************************************
+ * 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 <graphics_common.hpp>
+
+namespace opencl
+{
+ template<typename T>
+ void copy_plot3(const Array<T> &P, fg::Plot3* plot3);
+}
+
+#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