# CUDA - OpenGL interop

**URL:** <https://numba.discourse.group/t/cuda-opengl-interop/1865>\
**Category:** Support: How do I do ...?\
**Created:** [March 29, 2023, 4:14pm UTC](https://numba.discourse.group/t/cuda-opengl-interop/1865 "2023-03-29T16:14:30Z")\
**Posts on this page:** 1\
**Showing post:** 8

<div class="post-metadata">

**Author:** ![biosckon](https://yyz2.discourse-cdn.com/free1/user_avatar/numba.discourse.group/biosckon/32/775_2.png) [@biosckon](https://numba.discourse.group/u/biosckon)\
**Post date:** [March 30, 2023, 10:55am UTC](https://numba.discourse.group/t/cuda-opengl-interop/1865/8 "2023-03-30T10:55:11Z")

</div>

Code below does what is intended using code from [grapics api interop](https://numba.discourse.group/t/graphics-api-interop/593/8)

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`

```python
# 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)

```

---

_[View the full topic](https://numba.discourse.group/t/cuda-opengl-interop/1865)._