OpenGL interoperability

Contents

OpenGL interoperability#

The HIP–OpenGL interoperation involves mapping OpenGL resources, such as buffers and textures, for HIP to interact with OpenGL. This mapping process enables HIP to utilize these resources directly, bypassing the need for costly data transfers between the CPU and GPU. This capability is useful in applications that require both intensive GPU computation and real-time visualization.

The graphics resources must be registered using functions like hipGraphicsGLRegisterBuffer() or hipGraphicsGLRegisterImage() then they can be mapped to HIP with hipGraphicsMapResources() function.

After mapping, the hipGraphicsResourceGetMappedPointer() or hipGraphicsSubResourceGetMappedArray() functions used to retrieve a device pointer to the mapped resource, which can then be used in HIP kernels.

Unmapping resources with hipGraphicsUnmapResources() after computations ensure proper resource management.

Example#

ROCm examples have a HIP–OpenGL interoperation example, where a simple HIP kernel is used to simulate a sine wave and rendered to a window as a grid of triangles using OpenGL. For a working example, there are multiple initialization steps needed like creating and opening a window, initializing OpenGL or selecting the OpenGL-capable device. After the initialization in the example, the kernel simulates the sinewave and updates the window’s framebuffer in a cycle until the window is closed.

Note

The more recent OpenGL functions are loaded with OpenGL loader, as these are not loaded by default on all platforms. The use of a custom loader is shown in the following example

    // Make GLFW use a custom loader - we need this for the more recent OpenGL functions,
    // as these are not loaded by default on all platforms.
    if(!gladLoadGLLoader(reinterpret_cast<GLADloadproc>(glfwGetProcAddress)))
    {
        std::cerr << "Failed to load OpenGL function pointers" << std::endl;
        return error_exit_code;
    }

The OpenGL buffer is imported to HIP in the following way:

        // Import the OpenGL height buffer into a HIP graphics resource.
        HIP_CHECK(hipGraphicsGLRegisterBuffer(
            &this->hip_height_buffer,
            renderer.height_buffer,
            // We are going to write to this buffer from HIP,
            // but we do not need to read from it.
            // As an optimization we can pass hipGraphicsRegisterFlagsWriteDiscard,
            // so that the driver knows that we do not need the old values of
            // the buffer.
            hipGraphicsRegisterFlagsWriteDiscard));

        // After importing the OpenGL height buffer into HIP, map it into HIP memory so that we can use it.
        HIP_CHECK(hipGraphicsMapResources(1, &this->hip_height_buffer, this->hip_stream));

        // Fetch the device pointer that points to the OpenGL buffer's memory.
        // This function also fetches the size of the buffer. We already know it, but we still need to pass
        // a valid pointer to hipGraphicsResourceGetMappedPointer.
        size_t size;
        HIP_CHECK(
            hipGraphicsResourceGetMappedPointer(reinterpret_cast<void**>(&this->hip_height_ptr),
                                                &size,
                                                this->hip_height_buffer));

The imported pointer is manipulated in the sinewave kernel as shown in the following example:

/// \brief The main HIP kernel for this example - computes a simple sine wave over a
/// 2-dimensional grid of points.
/// \param height_map - the grid of points to compute a sine wave for. It is expected to be
///   a \p grid_width by \p grid_height array packed into memory.(y on the inner axis).
/// \param time - The current time relative to the start of the program.
__global__ void sinewave_kernel(float* height_map, const float time)
{
    const float        freq = 10.f;
    const unsigned int x    = blockIdx.x * blockDim.x + threadIdx.x;
    const unsigned int y    = blockIdx.y * blockDim.y + threadIdx.y;
    const float        u    = (2.f * x) / grid_width - 1.f;
    const float        v    = (2.f * y) / grid_height - 1.f;

    if(x < grid_width && y < grid_height)
    {
        height_map[x * grid_width + y] = sinf(u * freq + time) * cosf(v * freq + time);
    }
}

        // The tile size to be used for each block of the computation. A tile is
        // tile_size by tile_size threads in this case, since we are invoking the
        // computation over a 2D-grid.
        constexpr size_t tile_size = 8;

        // Launch the HIP kernel to advance the simulation.
        sinewave_kernel<<<dim3(ceiling_div(grid_width, tile_size),
                               ceiling_div(grid_height, tile_size)),
                          dim3(tile_size, tile_size),
                          0,
                          this->hip_stream>>>(this->hip_height_ptr, time);

        // Check that no errors occured while launching the kernel.
        HIP_CHECK(hipGetLastError());

The HIP graphics resource that is imported from the OpenGL buffer and is not needed anymore should be unmapped and unregistered as shown in the following way:

        HIP_CHECK(hipStreamSynchronize(this->hip_stream));
        HIP_CHECK(hipGraphicsUnmapResources(1, &this->hip_height_buffer, this->hip_stream));
        HIP_CHECK(hipGraphicsUnregisterResource(this->hip_height_buffer));
        HIP_CHECK(hipStreamDestroy(this->hip_stream));
Contents