diff --git a/.wordlist.txt b/.wordlist.txt index 55bf87e5e5..b3b8686678 100644 --- a/.wordlist.txt +++ b/.wordlist.txt @@ -73,7 +73,7 @@ iGPU inlined inplace interop -Interoperation +interoperation interoperate interoperation Interprocess @@ -162,6 +162,8 @@ unintuitive UMM unmap unmapped +unmapping +unregister upscaled variadic vulkan diff --git a/docs/how-to/hip_runtime_api.rst b/docs/how-to/hip_runtime_api.rst index 913c379e8c..223f1b296e 100644 --- a/docs/how-to/hip_runtime_api.rst +++ b/docs/how-to/hip_runtime_api.rst @@ -44,4 +44,5 @@ Here are the various HIP Runtime API high level functions: * :doc:`./hip_runtime_api/hipgraph` * :doc:`./hip_runtime_api/call_stack` * :doc:`./hip_runtime_api/multi_device` +* :doc:`./hip_runtime_api/opengl_interop` * :doc:`./hip_runtime_api/external_interop` \ No newline at end of file diff --git a/docs/how-to/hip_runtime_api/opengl_interop.rst b/docs/how-to/hip_runtime_api/opengl_interop.rst new file mode 100644 index 0000000000..45f34de257 --- /dev/null +++ b/docs/how-to/hip_runtime_api/opengl_interop.rst @@ -0,0 +1,94 @@ +.. meta:: + :description: HIP provides an OpenGL interoperability API that allows + efficient data sharing between HIP's computing power and + OpenGL's graphics rendering. + :keywords: AMD, ROCm, HIP, OpenGL, interop, interoperability + +******************************************************************************* +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 +:cpp:func:`hipGraphicsGLRegisterBuffer` or :cpp:func:`hipGraphicsGLRegisterImage` +then they can be mapped to HIP with :cpp:func:`hipGraphicsMapResources` +function. + +After mapping, the :cpp:func:`hipGraphicsResourceGetMappedPointer` or +:cpp:func:`hipGraphicsSubResourceGetMappedArray` functions used to retrieve a +device pointer to the mapped resource, which can then be used in HIP kernels. + +Unmapping resources with :cpp:func:`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 + + .. + + .. literalinclude:: ../../tools/example_codes/opengl_interop.hip + :start-after: // [Sphinx opengl functions load start] + :end-before: // [Sphinx opengl functions load end] + :language: cpp + + .. + +The OpenGL buffer is imported to HIP in the following way: + +.. + +.. literalinclude:: ../../tools/example_codes/opengl_interop.hip + :start-after: // [Sphinx buffer register and get start] + :end-before: // [Sphinx buffer register and get end] + :language: cpp + +.. + +The imported pointer is manipulated in the sinewave kernel as shown in the +following example: + +.. + +.. literalinclude:: ../../tools/example_codes/opengl_interop.hip + :start-after: /// [Sphinx sinewave kernel start] + :end-before: /// [Sphinx sinewave kernel end] + :language: cpp + +.. literalinclude:: ../../tools/example_codes/opengl_interop.hip + :start-after: // [Sphinx buffer use in kernel start] + :end-before: // [Sphinx buffer use in kernel end] + :language: cpp + +.. + +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: + +.. + +.. literalinclude:: ../../tools/example_codes/opengl_interop.hip + :start-after: // [Sphinx unregister start] + :end-before: // [Sphinx unregister end] + :language: cpp + +.. diff --git a/docs/sphinx/_toc.yml.in b/docs/sphinx/_toc.yml.in index f0ce57f9ec..703b65969c 100644 --- a/docs/sphinx/_toc.yml.in +++ b/docs/sphinx/_toc.yml.in @@ -53,6 +53,7 @@ subtrees: - file: how-to/hip_runtime_api/hipgraph - file: how-to/hip_runtime_api/call_stack - file: how-to/hip_runtime_api/multi_device + - file: how-to/hip_runtime_api/opengl_interop - file: how-to/hip_runtime_api/external_interop - file: how-to/hip_porting_guide - file: how-to/hip_porting_driver_api diff --git a/docs/tools/example_codes/opengl_interop.hip b/docs/tools/example_codes/opengl_interop.hip new file mode 100644 index 0000000000..64ece9ddf2 --- /dev/null +++ b/docs/tools/example_codes/opengl_interop.hip @@ -0,0 +1,628 @@ +// MIT License +// +// Copyright (c) 2022-2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#include "nvidia_hip_fix.hpp" + +#include "example_utils.hpp" + +#include "glad/glad.h" + +#include +#include +#include + +#include +#include +#include +#include +#include + +/// \brief The number of triangles that the example's grid is in width. +constexpr uint32_t grid_width = 256; +/// \brief The number of triangles that the example's grid is in height. +constexpr uint32_t grid_height = 256; + +/// \brief The OpenGL vertex shader that is used to render the triangles in this example. +/// The grid x- and y-positions are used to set the triangle coordinates in clip space. +/// The height value is passed on to the fragment shader. +constexpr const char* vertex_shader = R"( +#version 330 core + +in float in_height; +in vec2 in_xy; + +out float frag_height; + +void main() +{ + gl_Position = vec4(in_xy, 0, 1); + frag_height = in_height; +} +)"; + +/// \brief The OpenGL fragment shader that is used to render the triangles in this example. +/// The "height" value is used to shade the vertex. Its values are interpolated linearly +/// between the vertex and fragment shaders. +constexpr const char* fragment_shader = R"( +#version 330 core + +in float frag_height; + +void main() +{ + gl_FragColor = vec4(vec3(frag_height * 0.5 + 0.5), 1.0); +} +)"; + +/// \brief Initialize a GLFW window with initial dimensions. +GLFWwindow* create_window(const int initial_width, const int initial_height) +{ + /// [Sphinx-create-window] + glfwWindowHint(GLFW_CONTEXT_VERSION_MAJOR, 3); + glfwWindowHint(GLFW_CONTEXT_VERSION_MINOR, 3); + glfwWindowHint(GLFW_OPENGL_PROFILE, GLFW_OPENGL_CORE_PROFILE); + glfwWindowHint(GLFW_OPENGL_DEBUG_CONTEXT, GLFW_TRUE); + + GLFWwindow* window = glfwCreateWindow(initial_width, + initial_height, + "OpenGL-HIP interop example", + nullptr, + nullptr); + if(window == nullptr) + { + std::cerr << "Failed to create GLFW window\n"; + std::exit(error_exit_code); + } + /// [Sphinx-create-window] + return window; +} + +/// \brief Select a HIP device that is compatible with the current OpenGL context. +/// \returns A HIP device-id that is capable of rendering the example. If no +/// suitable device is found, an error is printed and the program is exited. +int pick_hip_device() +{ + /// [Sphinx-pick device] + unsigned int gl_device_count; + int hip_device; + HIP_CHECK( + hipGLGetDevices(&gl_device_count, &hip_device, 1, hipGLDeviceList::hipGLDeviceListAll)); + + if(gl_device_count == 0) + { + std::cerr << "System has no OpenGL-capable HIP devices" << std::endl; + std::exit(error_exit_code); + } + /// [Sphinx-pick device] + + return hip_device; +} + +/// \brief Utility function to compile shader source into an OpenGL shader. +/// If the shader could not be compiled, this function prints the compile log +/// and exits the program. +/// \param type - The OpenGL shader type for this shader, for example +/// \p GL_VERTEX_SHADER or \p GL_FRAGMENT_SHADER. +/// \param source - The GLSL source code for the shader. +GLuint compile_shader(const GLenum type, const char* const source) +{ + const GLuint shader = glCreateShader(type); + + const GLint length = static_cast(std::strlen(source)); + glShaderSource(shader, 1, &source, &length); + glCompileShader(shader); + + GLint compile_status; + glGetShaderiv(shader, GL_COMPILE_STATUS, &compile_status); + + if(compile_status != GL_TRUE) + { + // Compiling failed, get the shader log and print it to the user. + GLint log_length; + glGetShaderiv(shader, GL_INFO_LOG_LENGTH, &log_length); + std::vector log(log_length); + glGetShaderInfoLog(shader, length, nullptr, log.data()); + std::cerr << "Failed to compile shader:\n"; + std::cerr.write(log.data(), log.size()) << std::endl; + std::exit(error_exit_code); + } + + return shader; +} + +/// \brief Utility function to compile and link a vertex and fragment shader into an OpenGL +/// shader program. +/// If the shaders could not be compiled, a log is printed and the program is exited. +/// \param vert_src - The GLSL source code for the shader program's vertex shader. +/// \param frag_src - The GLSL source code for the shader program's fragment shader. +GLuint compile_shader_program(const char* const vert_src, const char* const frag_src) +{ + const GLuint program = glCreateProgram(); + + const GLuint vert = compile_shader(GL_VERTEX_SHADER, vert_src); + const GLuint frag = compile_shader(GL_FRAGMENT_SHADER, frag_src); + + glAttachShader(program, frag); + glAttachShader(program, vert); + + glLinkProgram(program); + + GLint link_status; + glGetProgramiv(program, GL_LINK_STATUS, &link_status); + if(link_status != GL_TRUE) + { + // Linking failed, get the program link log and print it to the user. + GLint log_length; + glGetProgramiv(program, GL_INFO_LOG_LENGTH, &log_length); + std::vector log(log_length); + glGetProgramInfoLog(program, log_length, nullptr, log.data()); + std::cerr << "Failed to link program:\n"; + std::cerr.write(log.data(), log.size()) << std::endl; + std::exit(error_exit_code); + } + + glDetachShader(program, frag); + glDetachShader(program, vert); + + glDeleteShader(frag); + glDeleteShader(vert); + + return program; +} + +/// \brief This structure contains the OpenGL handles that this example uses to render the +/// triangle grid to the screen. +/// +/// Three buffers are used to render the triangle grid, the color of which is determined by +/// a HIP compulation in \p simulator: +/// - One buffer contains the height of each triangle (rendered as color). +/// - One buffer holds the x- and y-coordinates for each of the corners of the triangle. Note: these +/// coordinates are unique, as the triangles that are made up from these points are defined by the +/// - Index buffer, that holds indices into the former two buffers to make up a list of triangles. +struct renderer +{ + /// The total number of vertices for the triangles. + constexpr static size_t num_verts = grid_width * grid_height; + /// The number of bytes in the x- and y-coordinates buffer. Each x/y coordinate is encoded as + /// a pair of floats, which are stored in a packed array-of-structures format: | x | y | x | y | ... |. + constexpr static size_t grid_buffer_size = num_verts * sizeof(float) * 2; + /// The number of bytes in the height buffer. Each height is encoded as a floating point value. + /// This buffer will be shared with HIP, which is why these coordinates are + /// stored in a separate buffer. + constexpr static size_t height_buffer_size = num_verts * sizeof(float); + + /// The number of indices in the index buffer. Each triangle has 3 points, each square in the grid + /// is made up of 2 triangles. There are (width - 1) by (height - 1) squares in the grid. + constexpr static size_t num_indices = (grid_width - 1) * (grid_height - 1) * 3 * 2; + /// The number of bytes in the index buffer. Each index is encoded as a 32-bit int. + constexpr static size_t index_buffer_size = num_indices * sizeof(uint32_t); + + /// An OpenGL handle to a Vertex Array Object, which has the grid and height buffers + /// bound to the corresponding attribute in the shader program (program) used for rendering. + GLuint vao; + + /// Handle to the buffer that holds the indices for the triangles to render. + GLuint index_buffer; + + /// Handle to the buffer that holds the x- and y-coordinates for each grid point. + GLuint grid_buffer; + + /// Handle to the buffer that holds the heights each grid point. This buffer is shared with HIP. + GLuint height_buffer; + + /// Handle to the OpenGL shader program that this example uses to render the triangles to the screen. + GLuint program; + + /// Counters used to keep track of the rendering performance. + uint32_t fps_frame = 0; + std::chrono::high_resolution_clock::time_point fps_start_time; + + /// \brief Initialize OpenGL rendering resources. + renderer() + { + // Create a vertex array used to bind the attribute buffers. + glGenVertexArrays(1, &this->vao); + + // Also generate the buffers in question. + GLuint buffers[3]; + glGenBuffers(std::size(buffers), buffers); + this->index_buffer = buffers[0]; + this->grid_buffer = buffers[1]; + this->height_buffer = buffers[2]; + + // Compile the shader program used to render the triangles. + this->program = compile_shader_program(vertex_shader, fragment_shader); + + // Upload the initial data to the buffers. + this->initialize_buffer_data(); + + // Set up the VAO by binding the height and grid buffers to the attribute locations + // in the shader program. + glBindVertexArray(this->vao); + + // Note - keep variable "in_height" in sync with shader. + glBindBuffer(GL_ARRAY_BUFFER, this->height_buffer); + const GLuint height_attrib = glGetAttribLocation(this->program, "in_height"); + glVertexAttribPointer(height_attrib, 1, GL_FLOAT, GL_FALSE, 0, 0); + glEnableVertexAttribArray(height_attrib); + + // Note - keep variable "in_xy" in sync with shader. + const GLuint grid_attrib = glGetAttribLocation(this->program, "in_xy"); + glBindBuffer(GL_ARRAY_BUFFER, this->grid_buffer); + glVertexAttribPointer(grid_attrib, 2, GL_FLOAT, GL_FALSE, 0, 0); + glEnableVertexAttribArray(grid_attrib); + + this->fps_start_time = std::chrono::high_resolution_clock::now(); + } + + renderer(const renderer&) = delete; + renderer& operator=(const renderer&) = delete; + + renderer(renderer&&) = delete; + renderer& operator=(renderer&&) = delete; + + ~renderer() + { + glDeleteProgram(this->program); + GLuint buffers[] = {this->index_buffer, this->grid_buffer, this->height_buffer}; + glDeleteBuffers(std::size(buffers), buffers); + glDeleteVertexArrays(1, &this->vao); + } + + /// \brief Upload the initial values for each buffer to Vulkan. + void initialize_buffer_data() const + { + // Initialize the height buffer. + glBindBuffer(GL_ARRAY_BUFFER, this->height_buffer); + // We do not need to fill it, as that is going to be done from HIP, but we + // do need to allocate it from OpenGL. This is done simply by passing `nullptr` as + // initial data pointer. + // GL_DYNAMIC_DRAW is passed because this buffer is going to be updated every frame, + // and is going to be used to hold vertex data for drawing - this may help the driver + // to render more efficiently. + glBufferData(GL_ARRAY_BUFFER, height_buffer_size, nullptr, GL_DYNAMIC_DRAW); + + // Initialize the grid buffer. + { + glBindBuffer(GL_ARRAY_BUFFER, this->grid_buffer); + // Avoid having to allocate on host by allocating the buffer in OpenGL and then mapping it + // into host-memory to initialize it. + // This buffer is going to be initialized once and is going to be used for drawing, + // so pass GL_STATIC_DRAW as usage hint. + glBufferData(GL_ARRAY_BUFFER, grid_buffer_size, nullptr, GL_STATIC_DRAW); + + float* grid = reinterpret_cast(glMapBuffer(GL_ARRAY_BUFFER, GL_WRITE_ONLY)); + for(uint32_t y = 0; y < grid_height; ++y) + { + for(uint32_t x = 0; x < grid_width; ++x) + { + *grid++ = (2.0f * x) / (grid_width - 1) - 1; + *grid++ = (2.0f * y) / (grid_height - 1) - 1; + } + } + + // Let OpenGL know that we are done with this buffer. + glUnmapBuffer(GL_ARRAY_BUFFER); + } + + // Initialize the index buffer + { + glBindBuffer(GL_ELEMENT_ARRAY_BUFFER, this->index_buffer); + // Similar as the grid buffer, this buffer is going to be initialized once and is then used + // for drawing. + glBufferData(GL_ELEMENT_ARRAY_BUFFER, index_buffer_size, nullptr, GL_STATIC_DRAW); + + uint32_t* indices + = reinterpret_cast(glMapBuffer(GL_ELEMENT_ARRAY_BUFFER, GL_WRITE_ONLY)); + for(uint32_t y = 0; y < grid_height - 1; ++y) + { + for(uint32_t x = 0; x < grid_width - 1; ++x) + { + *indices++ = (y + 0) * grid_width + (x + 0); + *indices++ = (y + 1) * grid_width + (x + 0); + *indices++ = (y + 0) * grid_width + (x + 1); + *indices++ = (y + 1) * grid_width + (x + 0); + *indices++ = (y + 1) * grid_width + (x + 1); + *indices++ = (y + 0) * grid_width + (x + 1); + } + } + + glUnmapBuffer(GL_ELEMENT_ARRAY_BUFFER); + } + } + + /// \brief Bind the OpenGL pipeline state for this renderer. + void bind() const + { + glBindVertexArray(this->vao); + glBindBuffer(GL_ELEMENT_ARRAY_BUFFER, this->index_buffer); + glUseProgram(this->program); + } + + /// \brief Draw the next frame to the window. This requires the render state be bound using + /// bind. + void draw() + { + glDrawElements(GL_TRIANGLES, num_indices, GL_UNSIGNED_INT, nullptr); + + // Output a native performance measurement. + ++this->fps_frame; + const auto frame_time = std::chrono::high_resolution_clock::now(); + const auto time_diff = frame_time - this->fps_start_time; + if(time_diff > std::chrono::seconds{5}) + { + const auto time_diff_sec + = std::chrono::duration_cast>(time_diff).count(); + std::cout << "Average FPS (over " << double_precision(time_diff_sec, 2, true) + << " seconds): " << double_precision(this->fps_frame / time_diff_sec, 2, true) + << " (" << double_precision((time_diff_sec * 1000) / this->fps_frame, 2, true) + << " ms per frame, " << this->fps_frame << " frames)" << std::endl; + this->fps_frame = 0; + this->fps_start_time = frame_time; + } + } +}; + +/// [Sphinx sinewave kernel start] +/// \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); + } +} +/// [Sphinx sinewave kernel end] + +/// \brief This structure contains the HIP state and functionality used to advance the simulation. +/// Initializing a \p simulator fetches the OpenGL height buffer from the corresponding renderer, +/// and imports it as a HIP device pointer. This pointer is then passed to the simulation kernel +/// (sinewave_kernel), which updates the values in it. When renderer::draw is called, +/// the updated values are read from the buffer in OpenGL and used to render the triangle grid. +struct simulator +{ + /// The HIP stream used to advance the simulation. This must be created from an OpenGL-interop + /// capable device, see pick_hip_device. + hipStream_t hip_stream; + /// A HIP graphics resource that is imported from the OpenGL height buffer to simulate. + hipGraphicsResource_t hip_height_buffer; + /// A device pointer to the height buffer, imported from the OPenGL height buffer. + float* hip_height_ptr; + + /// The start time of the program, used for the simulation. + std::chrono::high_resolution_clock::time_point start_time; + + /// \brief Initialize a simulator, that uses a particular HIP device. + /// \param renderer - The renderer that will be used to render the example. Its height buffer + /// is imported to HIP for use with this simulator. + explicit simulator(const int hip_device, const renderer& renderer) + { + // Create a HIP stream for the target device. + HIP_CHECK(hipSetDevice(hip_device)); + HIP_CHECK(hipStreamCreate(&this->hip_stream)); + + // [Sphinx buffer register and get start] + // 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(&this->hip_height_ptr), + &size, + this->hip_height_buffer)); + // [Sphinx buffer register and get end] + + this->start_time = std::chrono::high_resolution_clock::now(); + } + + simulator(const simulator&) = delete; + simulator& operator=(const simulator&) = delete; + + simulator(simulator&&) = delete; + simulator& operator=(simulator&&) = delete; + + ~simulator() + { + // [Sphinx unregister start] + 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)); + // [Sphinx unregister end] + } + + /// \brief Advance the simulation one step. + void step() + { + const auto now = std::chrono::high_resolution_clock::now(); + const float time + = std::chrono::duration(now - this->start_time) + .count(); + + // [Sphinx buffer use in kernel start] + // 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<<hip_stream>>>(this->hip_height_ptr, time); + + // Check that no errors occured while launching the kernel. + HIP_CHECK(hipGetLastError()); + // [Sphinx buffer use in kernel end] + } +}; + +/// \brief GLFW window resize callback: If the window is resized then we need to re-size +/// the OpenGL viewport. +void resize_callback(GLFWwindow* const window, const int width, const int height) +{ + (void)window; + glViewport(0, 0, width, height); +} + +/// \brief Program entry point. +int main() +{ + // The initial width of the GLFW window when the example is first started. + constexpr int initial_window_width = 1280; + // The initial height of the GLFW window. + constexpr int initial_window_height = 800; + + // Initialize GLFW. + glfwSetErrorCallback( + [](int code, const char* const message) + { std::cerr << "A glfw error encountered: " << message << "(" << code << ")\n"; }); + + if(glfwInit() != GLFW_TRUE) + { + std::cerr << "failed to initialize GLFW\n"; + return error_exit_code; + } + + // Initialize the GLFW window used to render the example. + GLFWwindow* const window = create_window(initial_window_width, initial_window_height); + + // Ensure that we are using the OpenGL context associated to the Window. + glfwMakeContextCurrent(window); + + // [Sphinx opengl functions load start] + // 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(glfwGetProcAddress))) + { + std::cerr << "Failed to load OpenGL function pointers" << std::endl; + return error_exit_code; + } + // [Sphinx opengl functions load end] + + // Disable vsync. + glfwSwapInterval(0); + + // If the OpenGL GL_ARB_debug_output extension is present, set a callback that is called + // whenever an OpenGL error occurs. This saves us calling glGetError after every OpenGL function. + if(GLAD_GL_ARB_debug_output) + { + glDebugMessageCallbackARB( + [](GLenum, + GLenum, + GLuint, + GLenum severity, + GLsizei length, + const GLchar* message, + const void*) + { + std::cerr << "[OpenGL] "; + std::cerr.write(message, length) << std::endl; + if(severity == GL_DEBUG_SEVERITY_HIGH_ARB) + { + std::exit(error_exit_code); + } + }, + nullptr); + // We just want the errors: First disable all messaging, and then enable just the + // most severe ones. + glDebugMessageControlARB(GL_DONT_CARE, GL_DONT_CARE, GL_DONT_CARE, 0, NULL, GL_FALSE); + glDebugMessageControlARB(GL_DONT_CARE, + GL_DONT_CARE, + GL_DEBUG_SEVERITY_HIGH_ARB, + 0, + NULL, + GL_TRUE); + // Report errors synchronously instead of asynchronously. + glEnable(GL_DEBUG_OUTPUT_SYNCHRONOUS_ARB); + } + + // Figure out which HIP device we need to use. + // This device needs to be interop-capable (see pick_hip_device). + const int hip_device = pick_hip_device(); + + // Let the user know which device we are using, on both the OpenGL and HIP sides. + hipDeviceProp_t hip_props; + HIP_CHECK(hipGetDeviceProperties(&hip_props, hip_device)); + const GLubyte* const device_name = glGetString(GL_RENDERER); + std::cout << "Using device " << device_name << " (hip device " << hip_device + << ", compute capability " << hip_props.major << "." << hip_props.minor << ")\n"; + + // Sub-scope to call destructors before terminating GLFW. + { + renderer renderer; + simulator simulator(hip_device, renderer); + + // There are no other renderers, so we can bind the OpenGL state once. + renderer.bind(); + + glfwSetFramebufferSizeCallback(window, resize_callback); + glClearColor(0, 0, 0, 1); + + // The main rendering loop. + // Repeat for as long as the window is not closed. + while(glfwWindowShouldClose(window) == GLFW_FALSE) + { + glClear(GL_COLOR_BUFFER_BIT); + + // First step the simulation so that the height buffer is ready + // for the next frame. + simulator.step(); + + // Draw the example to the window's framebuffer. + renderer.draw(); + + // Present the framebuffer on screen. + glfwSwapBuffers(window); + glfwPollEvents(); + } + } + + // Clean up GLFW. + glfwDestroyWindow(window); + glfwTerminate(); +} diff --git a/docs/tools/update_example_codes.py b/docs/tools/update_example_codes.py index ae74bc4e8c..32a4320750 100644 --- a/docs/tools/update_example_codes.py +++ b/docs/tools/update_example_codes.py @@ -1,3 +1,4 @@ import urllib.request +urllib.request.urlretrieve("https://raw.githubusercontent.com/ROCm/rocm-examples/refs/heads/develop/HIP-Basic/opengl_interop/main.hip", "docs/tools/example_codes/opengl_interop.hip") urllib.request.urlretrieve("https://raw.githubusercontent.com/ROCm/rocm-examples/refs/heads/develop/HIP-Basic/vulkan_interop/main.hip", "docs/tools/example_codes/external_interop.hip")