Hello everyone. Some time ago (2010-09-11) I volunteered for fixing PyCUDA to use non-deprecated OpenGL interoperability functions. Here is THE patch.
The biggest problem is that IMO (I might be wrong - do
not know PyCUDA to such intimate details) it introduces
non-backward-compatible change.
I have split BufferObject into two:
BufferObject responsible for buffers (VBO, etc.)
ImageObject responsible for textures, etc.
Previously BufferObject was responsible for both those
cases - now anyone who wants to use CUDA on images
must use ImageObject (which is API-breaking change).
All tests and examples run without any problem.
Test undistributed/reduction-perf.py throws
out of memory on cuMemAllod, but it is not
error related to changes in attached patch.
I also attach problem sample code (Qt+OpenGL+PyCUDA)
that uses new BufferObject.
I have checked and it works after applying my patch;
I only had to change initialisation to get program
using VBO running after applying my patch.
I have not checked ImageObject.
I have also started writing documentation for new pycuda.gl
functions.
I have removed gl_init as documentation says that cuGLInit()
is deprecated and is not doing anything (Reference 4.41.4.1).
Instead I have created gl_select_device; this function however
must be called as the first (before all other CUDA functions),
is cuda*, not cu* function, does not cooperate with
CUDAPP_CALL_GUARDED - so I do not know what to do with it.
I was not able to test it - I do not have setup with two devices.
Other problems:
Function pycuda.tools.get_default_device is depreciated,
but pycuda.gl.autoinit uses it. I have changed autoinit
to use Device(0), but it is not elegant and should be changed.
Function pycuda.tools.make_default_context contains
repeated code checking for presence of CUDA device:
def make_default_context():
ndevices = cuda.Device.count()
if ndevices == 0:
errmsg = "No CUDA enabled device found. Please check your
installation."
raise RuntimeError, errmsg
ndevices = cuda.Device.count()
if ndevices == 0:
raise RuntimeError("No CUDA enabled device found. "
"Please check your installation.")
Why? Is that omission, or does this double checking serve
some purpose?
Please look at this patch and send remarks.
The biggest problem is with necessity of moving
OpenGL image operations to use ImageObject, not
BufferObject.
During this weekend we have holiday in Poland and I will
be at my parent's home, so I am not sure if I will be able
to work on example of ImageObject.
Regards.
--
Tomasz Rybak <bogom...@post.pl> GPG/PGP key ID: 2AD5 9860
Fingerprint A481 824E 7DD3 9C0E C40A 488E C654 FB33 2AD5 9860
http://member.acm.org/~tomaszrybak
diff --git a/doc/source/gl.rst b/doc/source/gl.rst
index e5d196f..b367c24 100644
--- a/doc/source/gl.rst
+++ b/doc/source/gl.rst
@@ -8,7 +8,7 @@ GL Interoperability
.. module:: pycuda.gl
-.. function :: init()
+.. function :: select_device(device)
Enable GL interoperability for the already-created (so far non-GL)
and currently active :class:`pycuda.driver.Context`.
@@ -33,7 +33,36 @@ GL Interoperability
This will fail with a rather unhelpful error message if you don't already
have a GL context created and active.
-.. class :: BufferObject(bufobj)
+.. class :: map_flags
+
+ .. attribute :: cudaGraphicsMapFlagsNone
+
+ Read and write access to mapped OpenGL object from CUDA code.
+
+ .. attribute :: cudaGraphicsMapFlagsReadOnly
+
+ Read only access to mapped OpenGL object from CUDA code.
+
+ .. attribute :: cudaGraphicsMapFlagsWriteDiscard
+
+ Write only access to mapped OpenGL object from CUDA code. Reading
+ is prohibited.
+
+.. class :: map_targets
+
+ Type of OpenGL Image object that is mapped to CUDA.
+
+ .. attribute :: GL_TEXTURE_2D
+ .. attribute :: GL_TEXTURE_RECTANGLE
+ .. attribute :: GL_TEXTURE_CUBE_MAP
+ .. attribute :: GL_TEXTURE_3D
+ .. attribute :: GL_TEXTURE_2D_ARRAY
+ .. attribute :: GL_RENDERBUFFER
+
+.. class :: BufferObject(bufobj, flags = cudaGraphicsMapFlagsNone)
+
+ Object managing mapping of OpenGL buffers to CUDA. Cannot be used to
+ map images.
.. method :: unregister()
.. method :: handle()
@@ -45,6 +74,20 @@ GL Interoperability
.. method :: device_ptr()
.. method :: size()
+.. class :: ImageObject(bufobj, target, flags = cudaGraphicsMapFlagsNone)
+
+ Object managing mapping of OpenGL textures and render buffers to CUDA.
+
+ .. method :: unregister()
+ .. method :: handle()
+ .. method :: map()
+
+.. class :: ImageObjectMapping
+
+ .. method :: unmap()
+ .. method :: device_ptr()
+ .. method :: size()
+
.. note ::
See this `post <http://forums.nvidia.com/index.php?showtopic=88152>`_ on the
diff --git a/pycuda/gl/__init__.py b/pycuda/gl/__init__.py
index a18b3d2..379b42e 100644
--- a/pycuda/gl/__init__.py
+++ b/pycuda/gl/__init__.py
@@ -3,7 +3,11 @@ import pycuda._driver as _drv
if not _drv.have_gl_ext():
raise ImportError("PyCUDA was compiled without GL extension support")
-init = _drv.gl_init
+select_device = _drv.gl_select_device
make_context = _drv.make_gl_context
+map_flags = _drv.map_flags
+target_flags = _drv.target_flags
BufferObject = _drv.BufferObject
BufferObjectMapping = _drv.BufferObjectMapping
+ImageObject = _drv.ImageObject
+ImageObjectMapping = _drv.ImageObjectMapping
diff --git a/pycuda/gl/autoinit.py b/pycuda/gl/autoinit.py
index 9fb8e88..fc0d2b0 100644
--- a/pycuda/gl/autoinit.py
+++ b/pycuda/gl/autoinit.py
@@ -5,8 +5,9 @@ import pycuda.tools
cuda.init()
assert cuda.Device.count() >= 1
-device = pycuda.tools.get_default_device()
+device = cuda.Device(0)
context = cudagl.make_context(device)
+
import atexit
atexit.register(context.pop)
diff --git a/src/cpp/cuda_gl.hpp b/src/cpp/cuda_gl.hpp
index df8d289..c6d41c3 100644
--- a/src/cpp/cuda_gl.hpp
+++ b/src/cpp/cuda_gl.hpp
@@ -11,13 +11,22 @@
#endif
#include <cudaGL.h>
+#include <cuda_gl_interop.h>
+#include <driver_types.h>
namespace cuda { namespace gl {
inline
- void gl_init()
+ void gl_select_device(int device)
{
- CUDAPP_CALL_GUARDED(cuGLInit, ());
+ CUDAPP_PRINT_CALL_TRACE("cudaGLSetGLDevice");
+ cudaError_t cu_status_code = cudaSuccess;
+// cu_status_code = cudaGLSetGLDevice(device);
+ CUDAPP_PRINT_ERROR_TRACE("cudaGLSetGLDevice", cu_status_code);
+ if (cu_status_code != cudaSuccess) {
+ throw cuda::error("cudaGLSetGLDevice", CUDA_SUCCESS);
+ }
+// CUDAPP_CALL_GUARDED(cudaGLSetGLDevice, (device));
}
@@ -40,12 +49,13 @@ namespace cuda { namespace gl {
{
private:
GLuint m_handle;
+ CUgraphicsResource m_resource;
bool m_valid;
public:
- buffer_object(GLuint handle)
+ buffer_object(GLuint handle, cudaGraphicsMapFlags flags = cudaGraphicsMapFlagsNone)
: m_handle(handle), m_valid(true)
- { CUDAPP_CALL_GUARDED(cuGLRegisterBufferObject, (handle)); }
+ { CUDAPP_CALL_GUARDED(cuGraphicsGLRegisterBuffer, (&m_resource, handle, flags)); }
~buffer_object()
{
@@ -56,6 +66,9 @@ namespace cuda { namespace gl {
GLuint handle()
{ return m_handle; }
+ CUgraphicsResource * resource()
+ { return &m_resource; }
+
void unregister()
{
if (m_valid)
@@ -63,7 +76,7 @@ namespace cuda { namespace gl {
try
{
scoped_context_activation ca(get_context());
- CUDAPP_CALL_GUARDED_CLEANUP(cuGLUnregisterBufferObject, (m_handle));
+ CUDAPP_CALL_GUARDED_CLEANUP(cuGraphicsUnregisterResource, (m_resource));
m_valid = false;
}
CUDAPP_CATCH_CLEANUP_ON_DEAD_CONTEXT(buffer_object);
@@ -104,7 +117,9 @@ namespace cuda { namespace gl {
try
{
scoped_context_activation ca(get_context());
- CUDAPP_CALL_GUARDED_CLEANUP(cuGLUnmapBufferObject, (m_buffer_object->handle()));
+// Stream as third parameter
+// Is that a problem, or all main tasks are done in stream 0 in PyCUDA?
+ CUDAPP_CALL_GUARDED_CLEANUP(cuGraphicsUnmapResources, (1, m_buffer_object->resource(), NULL));
m_valid = false;
}
CUDAPP_CATCH_CLEANUP_ON_DEAD_CONTEXT(buffer_object_mapping)
@@ -128,10 +143,123 @@ namespace cuda { namespace gl {
{
CUdeviceptr devptr;
pycuda_size_t size;
- CUDAPP_CALL_GUARDED(cuGLMapBufferObject, (&devptr, &size, bobj->handle()));
+// Stream as third parameter
+// Is that a problem, or all main tasks are done in stream 0 in PyCUDA?
+ CUDAPP_CALL_GUARDED_CLEANUP(cuGraphicsMapResources, (1, bobj->resource(), NULL));
+ CUDAPP_CALL_GUARDED(cuGraphicsResourceGetMappedPointer, (&devptr, &size, *(bobj->resource())));
return new buffer_object_mapping(bobj, devptr, size);
}
+
+
+
+
+ class image_object : public context_dependent
+ {
+ private:
+ GLuint m_handle;
+ GLenum m_target;
+ CUgraphicsResource m_resource;
+ bool m_valid;
+
+ public:
+ image_object(GLuint handle, GLenum target, cudaGraphicsMapFlags flags = cudaGraphicsMapFlagsNone)
+ : m_handle(handle), m_target(target), m_resource(NULL), m_valid(true)
+ { CUDAPP_CALL_GUARDED(cuGraphicsGLRegisterBuffer, (&m_resource, handle, flags)); }
+
+ ~image_object()
+ {
+ if (m_valid)
+ unregister();
+ }
+
+ GLuint handle()
+ { return m_handle; }
+
+ CUgraphicsResource * resource()
+ { return &m_resource; }
+
+ void unregister()
+ {
+ if (m_valid)
+ {
+ try
+ {
+ scoped_context_activation ca(get_context());
+ CUDAPP_CALL_GUARDED_CLEANUP(cuGraphicsUnregisterResource, (m_resource));
+ m_valid = false;
+ }
+ CUDAPP_CATCH_CLEANUP_ON_DEAD_CONTEXT(image_object);
+ }
+ else
+ throw cuda::error("image_object::unregister", CUDA_ERROR_INVALID_HANDLE);
+ }
+ };
+
+
+
+ class image_object_mapping : public context_dependent
+ {
+ private:
+ boost::shared_ptr<image_object> m_image_object;
+ CUdeviceptr m_devptr;
+ unsigned int m_size;
+ bool m_valid;
+
+ public:
+ image_object_mapping(
+ boost::shared_ptr<image_object> iobj,
+ CUdeviceptr devptr,
+ unsigned int size)
+ : m_image_object(iobj), m_devptr(devptr), m_size(size), m_valid(true)
+ { }
+
+ ~image_object_mapping()
+ {
+ if (m_valid)
+ unmap();
+ }
+
+ void unmap()
+ {
+ if (m_valid)
+ {
+ try
+ {
+ scoped_context_activation ca(get_context());
+// Stream as third parameter
+// Is that a problem, or all main tasks are done in stream 0 in PyCUDA?
+ CUDAPP_CALL_GUARDED_CLEANUP(cuGraphicsUnmapResources, (1, m_image_object->resource(), NULL));
+ m_valid = false;
+ }
+ CUDAPP_CATCH_CLEANUP_ON_DEAD_CONTEXT(image_object_mapping)
+ }
+ else
+ throw cuda::error("image_object_mapping::unmap", CUDA_ERROR_INVALID_HANDLE);
+ }
+
+ CUdeviceptr device_ptr() const
+ { return m_devptr; }
+
+ unsigned int size() const
+ { return m_size; }
+ };
+
+
+
+
+ inline image_object_mapping *map_image_object(
+ boost::shared_ptr<image_object> iobj)
+ {
+ CUdeviceptr devptr;
+ pycuda_size_t size;
+// Stream as third parameter
+// Is that a problem, or all main tasks are done in stream 0 in PyCUDA?
+ CUDAPP_CALL_GUARDED_CLEANUP(cuGraphicsMapResources, (1, iobj->resource(), NULL));
+ CUDAPP_CALL_GUARDED(cuGraphicsResourceGetMappedPointer, (&devptr, &size, *(iobj->resource())));
+
+ return new image_object_mapping(iobj, devptr, size);
+ }
} }
diff --git a/src/wrapper/wrap_cudagl.cpp b/src/wrapper/wrap_cudagl.cpp
index 8e9ba02..2532262 100644
--- a/src/wrapper/wrap_cudagl.cpp
+++ b/src/wrapper/wrap_cudagl.cpp
@@ -1,5 +1,6 @@
#include <cuda.hpp>
#include <cuda_gl.hpp>
+#include <driver_types.h>
#include "tools.hpp"
#include "wrap_helpers.hpp"
@@ -19,13 +20,29 @@ void pycuda_expose_gl()
using py::arg;
using py::args;
- DEF_SIMPLE_FUNCTION(gl_init);
+ py::def("gl_select_device", gl_select_device, (arg("device")));
py::def("make_gl_context", make_gl_context, (arg("dev"), arg("flags")=0));
+// Reference Manual 4.9 and 4.23.3.6
+ py::enum_<cudaGraphicsMapFlags>("map_flags")
+ .value("cudaGraphicsMapFlagsNone", cudaGraphicsMapFlagsNone)
+ .value("cudaGraphicsMapFlagsReadOnly", cudaGraphicsMapFlagsReadOnly)
+ .value("cudaGraphicsMapFlagsWriteDiscard", cudaGraphicsMapFlagsWriteDiscard)
+ ;
+
+ py::enum_<GLenum>("target_flags")
+ .value("GL_TEXTURE_2D", GL_TEXTURE_2D)
+ .value("GL_TEXTURE_RECTANGLE", GL_TEXTURE_RECTANGLE)
+ .value("GL_TEXTURE_CUBE_MAP", GL_TEXTURE_CUBE_MAP)
+ .value("GL_TEXTURE_3D", GL_TEXTURE_3D)
+ .value("GL_TEXTURE_2D_ARRAY", GL_TEXTURE_2D_ARRAY)
+ .value("GL_RENDERBUFFER", GL_RENDERBUFFER)
+ ;
+
{
typedef buffer_object cl;
- py::class_<cl, shared_ptr<cl> >("BufferObject", py::init<GLuint>())
+ py::class_<cl, shared_ptr<cl> >("BufferObject", py::init<GLuint, py::optional<cudaGraphicsMapFlags> >())
.DEF_SIMPLE_METHOD(handle)
.DEF_SIMPLE_METHOD(unregister)
.def("map", map_buffer_object,
@@ -41,4 +58,24 @@ void pycuda_expose_gl()
.DEF_SIMPLE_METHOD(size)
;
}
+
+
+ {
+ typedef image_object cl;
+ py::class_<cl, shared_ptr<cl> >("ImageObject", py::init<GLuint, GLenum, py::optional<cudaGraphicsMapFlags> >())
+ .DEF_SIMPLE_METHOD(handle)
+ .DEF_SIMPLE_METHOD(unregister)
+ .def("map", map_image_object,
+ py::return_value_policy<py::manage_new_object>())
+ ;
+ }
+
+ {
+ typedef image_object_mapping cl;
+ py::class_<cl>("ImageObjectMapping", py::no_init)
+ .DEF_SIMPLE_METHOD(unmap)
+ .DEF_SIMPLE_METHOD(device_ptr)
+ .DEF_SIMPLE_METHOD(size)
+ ;
+ }
}
#! /usr/bin/python
import sys
import array
import math
import numpy
from OpenGL.GL import *
import pycuda
import pycuda.compiler
import pycuda.driver
import pycuda.gl
import pycuda.gpuarray
#pycuda.init()
#import pycuda.autoinit
from PyQt4 import QtCore, QtGui, QtOpenGL
vertexSource = """#version 150 core
uniform mat4 projectionMatrix;
in vec3 vertexCoordinate;
void main() {
vec4 a = vec4(vertexCoordinate, 1.0);
a.z = a.z - 2.0;
gl_Position = projectionMatrix * a;
}
"""
fragmentSource = """#version 150 core
void main() {
gl_FragColor = vec4(1.0, 1.0, 1.0, 1.0);
}
"""
kernel = """
__global__ void compute(float3 *pos) {
unsigned int x = threadIdx.x;
// if (x < N) {
if (x % 2 == 0) {
pos[x].z = pos[x].z+1;
} else {
pos[x].z = pos[x].z-1;
}
// }
}
"""
class Widget(QtOpenGL.QGLWidget):
def __init__(self, parent = None):
super(Widget, self).__init__(parent)
def minimumSizeHint(self):
return QtCore.QSize(400, 300)
def sizeHint(self):
return QtCore.QSize(800, 600)
def initializeGL(self):
self.qglClearColor(QtCore.Qt.black)
glEnable(GL_DEPTH_TEST)
program = glCreateProgram()
vertexShader = glCreateShader(GL_VERTEX_SHADER)
glShaderSource(vertexShader, [vertexSource])
glCompileShader(vertexShader)
glAttachShader(program, vertexShader)
fragmentShader = glCreateShader(GL_FRAGMENT_SHADER)
glShaderSource(fragmentShader, [fragmentSource])
glCompileShader(fragmentShader)
glAttachShader(program, fragmentShader)
glBindAttribLocation(program, 0, 'vertexCoordinate')
glLinkProgram(program)
glValidateProgram(program)
glUseProgram(program)
self.projectionMatrixLocation = glGetUniformLocation(program, 'projectionMatrix')
self.data = numpy.zeros((99, 3), numpy.float32)
for i in range(33):
self.data[i*3+0, 0] = 0
self.data[i*3+0, 1] = 0
self.data[i*3+0, 2] = 0
self.data[i*3+1, 0] = 1*math.cos(math.pi*(i+0)/16)
self.data[i*3+1, 1] = 1*math.sin(math.pi*(i+0)/16)
self.data[i*3+1, 2] = 0
self.data[i*3+2, 0] = 1*math.cos(math.pi*(i+1)/16)
self.data[i*3+2, 1] = 1*math.sin(math.pi*(i+1)/16)
self.data[i*3+2, 2] = 0
self.gl_buffer = glGenBuffers(1)
glBindBuffer(GL_ARRAY_BUFFER, self.gl_buffer)
glBufferData(GL_ARRAY_BUFFER, self.data, GL_DYNAMIC_DRAW)
glVertexAttribPointer(0, 3, GL_FLOAT, False, 0, None)
glEnableVertexAttribArray(0)
import pycuda.gl.autoinit
# pycuda.gl.init(0)
k = pycuda.compiler.SourceModule(kernel)
self.cuda_function = k.get_function("compute")
self.cuda_function.prepare("P", (99, 1, 1))
self.cuda_buffer = pycuda.gl.BufferObject(long(self.gl_buffer))
def paintGL(self):
cuda_object = self.cuda_buffer.map()
self.cuda_function.prepared_call((1, 1), cuda_object.device_ptr())
pycuda.driver.Context.synchronize()
cuda_object.unmap()
glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT)
glDrawArrays(GL_TRIANGLES, 0, 99)
def resizeGL(self, width, height):
side = min(width, height)
glViewport((width-side)/2, (height-side)/2, side, side)
projectionMatrix = numpy.zeros((4, 4), 'f')
top = +1.5
bottom = -1.5
left = -1.5
right = +1.5
near = 1.0
far = 3.0
projectionMatrix[0, 0] = (2.0)/(right-left)
projectionMatrix[0, 3] = -(right+left)/(right-left)
projectionMatrix[1, 1] = (2.0)/(top-bottom)
projectionMatrix[1, 3] = -(top+bottom)/(top-bottom)
projectionMatrix[2, 2] = -(2.0)/(far-near)
projectionMatrix[2, 3] = -(far+near)/(far-near)
projectionMatrix[3, 3] = 1.0
glUniformMatrix4fv(self.projectionMatrixLocation, 1, True, projectionMatrix)
def closeEvent(self, event):
glDeleteBuffers(1, long(self.gl_buffer))
self.gl_buffer = None
self.cuda_buffer = None
super(Widget, self).closeEvent(event)
class Window(QtGui.QWidget):
def __init__(self):
super(Window, self).__init__()
self.widget = Widget()
self.setWindowTitle("CUDA OpenGL")
layout = QtGui.QHBoxLayout()
layout.addWidget(self.widget)
self.setLayout(layout)
def closeEvent(self, event):
self.widget.close()
super(Window, self).closeEvent(event)
app = QtGui.QApplication(sys.argv)
window = Window()
window.show()
app.exec_()
signature.asc
Description: This is a digitally signed message part
_______________________________________________ PyCUDA mailing list PyCUDA@tiker.net http://lists.tiker.net/listinfo/pycuda
