CUDA - OpenGL interop

Numba Support: How do I do ...?
8

Code below does what is intended using code from grapics api interop

I’ll proceed with extension improvement…

I have a small problem running the code in vscode jupyter notebook and as a file from vscode… I’ll describe it in the follow up post, after I reduce it to the minimum reproducible.

Runs with no errors if launched as:
python example.py

# example.py
import ctypes
import numpy as np
from PIL import Image
from cuda import cudart
from cuda.cudart import cudaGraphicsRegisterFlags as GLFlags
import numba.cuda as nbcuda
import moderngl as mgl

#####################################################
# add 0.2 to the value
@nbcuda.jit
def shift_02(x):
    idx = nbcuda.grid(1)
    x[idx] = x[idx] + 0.2

###################################################
# wrap the raw pointer into DeviceNDArray
def gl_NDArray(ptr, shape, dtype=np.float32, strides=None, order="C"):
    shape, strides, dtype = nbcuda.api.prepare_shape_strides_dtype(
        shape, strides, dtype, order
    )
    datasize = nbcuda.driver.memory_size_from_info(shape, strides, dtype.itemsize)

    ctx = nbcuda.current_context()
    c_ptr = ctypes.c_uint64(ptr)
    mem = nbcuda.driver.MemoryPointer(ctx, c_ptr, datasize)
    return nbcuda.cudadrv.devicearray.DeviceNDArray(shape, strides, dtype, gpu_data=mem)

#####################################
# save fbo image into file
def save_image(fbo, name):
    data = fbo.read(components=3)
    image = Image.frombytes("RGB", fbo.size, data)
    image.save(name)


##################################################################
# these shaders do nothing just passing the data along the pipeline

vert_code = """
#version 330

in vec2 in_vert;
in vec3 in_color;

out vec3 color;

void main() {
    gl_Position = vec4(in_vert, 0.0, 1.0);
    color = in_color;
}
"""

frag_code = """
#version 330

in vec3 color;

out vec4 fragColor;

void main(){
    fragColor = vec4(color, 1.0);
}
"""

ctx = mgl.create_context(standalone=True)
prog = ctx.program(vertex_shader=vert_code, fragment_shader=frag_code)


# vertices for triangle with Red Green and Blue colors
vertices = np.array(
    # X     Y                   x     y                   x    y
    [-0.6, -0.6, 1.0, 0.0, 0.0, 0.6, -0.6, 0.0, 1.0, 0.0, 0.0, 0.6, 0.0, 0.0, 1.0],
    dtype="f4",
)

# init VBO, VAO and FBO
vbo = ctx.buffer(vertices)
vao = ctx.simple_vertex_array(prog, vbo, "in_vert", "in_color")
fbo = ctx.framebuffer(color_attachments=[ctx.texture((512, 512), 4)])
fbo.use()

# Get CUDA pointer to VBO
err, gl_resource = cudart.cudaGraphicsGLRegisterBuffer(
    vbo.glo, GLFlags.cudaGraphicsRegisterFlagsNone 
)
(err,) = cudart.cudaGraphicsMapResources(1, gl_resource, 0)

# get raw pointer to device array
(err, dev_ptr, dev_ptr_size) = cudart.cudaGraphicsResourceGetMappedPointer(gl_resource)
# wrap into DeviceNDArray
vbo_arr = gl_NDArray(dev_ptr, vertices.size)

# rendering

# 1st render as is
ctx.clear()
vao.render(mgl.TRIANGLES)
save_image(fbo, "image1.jpg")

# 2nd render after numba kernel call
ctx.clear()
shift_02[1,2](vbo_arr) # shift coordinates (x,y) of first point
nbcuda.synchronize()
vao.render(mgl.TRIANGLES)
save_image(fbo, "image2.jpg")

# CUDA cleanup
(err,) = cudart.cudaGraphicsUnmapResources(1, gl_resource, 0)
(err,) = cudart.cudaGraphicsUnregisterResource(gl_resource)