Skip to content

Commit

Permalink
[opengl] [refactor] Fix memory leakages using modern C++ memory manag…
Browse files Browse the repository at this point in the history
…ement features (#696)

* [skip ci] fix memory leakage in create_glsl_root_buffer

* [skip ci] use another way to test passing

* I love C++!

* [skip ci] minor changes

* [skip ci] tmp work save

* [skip ci] apply k-ye's suggestions

* [skip ci] no gc

* [skip ci] no a

* add opengl::GLSLLauncher (like metal::KernelManager)

* [skip ci] really done GLSLLauncher

* try GLSLLaunchGuard

* fix non-GL build

* [skip ci] no dirty words

* [skip ci] use resize

* fix non-GL build again

* nongl fix again

* really fix

* [skip ci] remove unused

* [skip ci] enforce code format

* [skip ci] apply reviews

* fix non-GL build

Co-authored-by: Taichi Gardener <taichigardener@gmail.com>
  • Loading branch information
archibate and taichi-gardener authored Apr 6, 2020
1 parent 213a3c7 commit 24e76a1
Show file tree
Hide file tree
Showing 7 changed files with 158 additions and 99 deletions.
141 changes: 72 additions & 69 deletions taichi/backends/opengl/opengl_api.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -38,7 +38,7 @@ struct GLShader {
id_ = glCreateShader(type);
}

GLShader(std::string source, GLuint type = GL_COMPUTE_SHADER)
GLShader(const std::string &source, GLuint type = GL_COMPUTE_SHADER)
: GLShader(type) {
this->compile(source);
}
Expand All @@ -47,12 +47,12 @@ struct GLShader {
glDeleteShader(id_);
}

GLShader &compile(const std::string &source) {
void compile(const std::string &source) const {
const GLchar *source_cstr = source.c_str();
glShaderSource(id_, 1, &source_cstr, nullptr);

glCompileShader(id_);
GLint status = GL_TRUE;
int status = GL_TRUE;
glGetShaderiv(id_, GL_COMPILE_STATUS, &status);
if (status != GL_TRUE) {
GLsizei logLength;
Expand All @@ -63,7 +63,6 @@ struct GLShader {
TI_ERROR("[glsl] error while compiling shader:\n{}\n{}",
add_line_markers(source), log.data());
}
return *this;
}
};

Expand All @@ -77,22 +76,21 @@ struct GLProgram {
explicit GLProgram(GLuint id) : id_(id) {
}

explicit GLProgram(GLShader &shader) : GLProgram() {
explicit GLProgram(const GLShader &shader) : GLProgram() {
this->attach(shader);
}

~GLProgram() {
glDeleteProgram(id_);
}

GLProgram &attach(GLShader &shader) {
void attach(const GLShader &shader) const {
glAttachShader(id_, shader.id_);
return *this;
}

GLProgram &link() {
void link() const {
glLinkProgram(id_);
GLint status = GL_TRUE;
int status = GL_TRUE;
glGetProgramiv(id_, GL_LINK_STATUS, &status);
if (status != GL_TRUE) {
GLsizei logLength;
Expand All @@ -102,16 +100,14 @@ struct GLProgram {
log[logLength] = 0;
TI_ERROR("[glsl] error while linking program:\n{}", log.data());
}
return *this;
}

GLProgram &use() {
void use() const {
glUseProgram(id_);
return *this;
}

template <typename T>
void set_uniform(std::string name, T value) {
void set_uniform(const std::string &name, T value) const {
GLuint loc = glGetUniformLocation(id_, name.c_str());
glapi_set_uniform(loc, value);
}
Expand Down Expand Up @@ -159,31 +155,30 @@ struct GLSSBO {
used as the source for GL drawing and image specification commands.
***/

GLSSBO &bind_data(void *data, size_t size, GLuint usage = GL_STATIC_READ) {
void bind_data(void *data, size_t size, GLuint usage = GL_STATIC_READ) const {
glBindBuffer(GL_SHADER_STORAGE_BUFFER, id_);
glBufferData(GL_SHADER_STORAGE_BUFFER, size, data, usage);
return *this;
}

GLSSBO &bind_index(size_t index) {
void bind_index(size_t index) const {
// SSBO index, is `layout(std430, binding = <index>)` in shader.
// We use only one SSBO though...
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, index, id_);
return *this;
}

GLSSBO &bind_range(size_t index, size_t offset, size_t size) {
void bind_range(size_t index, size_t offset, size_t size) const {
glBindBufferRange(GL_SHADER_STORAGE_BUFFER, index, id_, offset, size);
return *this;
}

void *map(size_t offset, size_t length, GLbitfield access = GL_MAP_READ_BIT) {
void *map(size_t offset,
size_t length,
GLbitfield access = GL_MAP_READ_BIT) const {
// map GPU memory to CPU address space, offset within SSBO data
glBindBuffer(GL_SHADER_STORAGE_BUFFER, id_);
return glMapBufferRange(GL_SHADER_STORAGE_BUFFER, offset, length, access);
}

void *map(GLbitfield access = GL_MAP_READ_BIT) {
void *map(GLbitfield access = GL_MAP_READ_BIT) const {
glBindBuffer(GL_SHADER_STORAGE_BUFFER, id_);
return glMapBuffer(GL_SHADER_STORAGE_BUFFER, access);
}
Expand All @@ -208,14 +203,14 @@ void initialize_opengl() {
glfwCreateWindow(1, 1, "Make GLEW Happy", nullptr, nullptr);
if (!window) {
const char *desc = nullptr;
GLint status = glfwGetError(&desc);
int status = glfwGetError(&desc);
if (!desc)
desc = "Unknown Error";
TI_ERROR("[glsl] cannot create GLFW window: error {}: {}", status, desc);
}
glfwHideWindow(window);
glfwMakeContextCurrent(window);
GLint status = glewInit();
int status = glewInit();
if (status != GLEW_OK) {
TI_ERROR("[glsl] cannot initialize GLEW: {}", glewGetErrorString(status));
}
Expand All @@ -232,41 +227,57 @@ void initialize_opengl() {
}
}

GLProgram *compile_glsl_program(std::string source) {
GLShader shader(source);
GLProgram *program = new GLProgram(shader);
program->link();
return program;
CompiledGLSL::CompiledGLSL(const std::string &source)
: glsl(std::make_unique<GLProgram>(GLShader(source))) {
glsl->link();
}

GLSSBO *root_ssbo;
struct GLSLLauncherImpl {
std::unique_ptr<GLSSBO> root_ssbo;
std::vector<GLSSBO> ssbo;
std::vector<char> root_buffer;
};

void create_glsl_root_buffer(size_t size) {
// if (root_ssbo) return;
GLSLLauncher::GLSLLauncher(size_t size) {
initialize_opengl();
root_ssbo =
new GLSSBO; // TODO(archibate): mem leaking, use std::optional instead
impl = std::make_unique<GLSLLauncherImpl>();
impl->root_ssbo = std::make_unique<GLSSBO>();
size += 2 * sizeof(int);
void *buffer = std::calloc(size, 1);
root_ssbo->bind_data(buffer, size, GL_DYNAMIC_READ);
root_ssbo->bind_index(0);
impl->root_buffer.resize(size, 0);
impl->root_ssbo->bind_data(impl->root_buffer.data(), size, GL_DYNAMIC_READ);
impl->root_ssbo->bind_index(0);
}

std::vector<GLSSBO> ssbo;
GLSLLaunchGuard::GLSLLaunchGuard(GLSLLauncherImpl *impl,
const std::vector<IOV> &iov)
: impl(impl), iov(iov) {
impl->ssbo = std::vector<GLSSBO>(iov.size());

for (int i = 0; i < impl->ssbo.size(); i++) {
if (!iov[i].size)
continue;
impl->ssbo[i].bind_index(i + 1);
impl->ssbo[i].bind_data(iov[i].base, iov[i].size,
GL_DYNAMIC_READ); // input
}
}

void begin_glsl_kernels(const std::vector<IOV> &iov) {
ssbo = std::vector<GLSSBO>(iov.size());
GLSLLaunchGuard::~GLSLLaunchGuard() {
// glMemoryBarrier(GL_SHADER_STORAGE_BARRIER_BIT); // TODO(archibate): move to
// Program::synchroize()

for (int i = 0; i < ssbo.size(); i++) {
for (int i = 0; i < impl->ssbo.size(); i++) {
if (!iov[i].size)
continue;
ssbo[i].bind_index(i + 1);
ssbo[i].bind_data(iov[i].base, iov[i].size, GL_DYNAMIC_READ); // input
void *p = impl->ssbo[i].map(0, iov[i].size); // output
std::memcpy(iov[i].base, p, iov[i].size);
}
glUnmapBuffer(GL_SHADER_STORAGE_BUFFER);
impl->ssbo.clear();
}

void launch_glsl_kernel(GLProgram *program, int num_groups) {
program->use();
void CompiledGLSL::launch_glsl(int num_groups) const {
glsl->use();

// https://www.khronos.org/opengl/wiki/Compute_Shader
// https://community.arm.com/developer/tools-software/graphics/b/blog/posts/get-started-with-compute-shaders
Expand All @@ -278,20 +289,6 @@ void launch_glsl_kernel(GLProgram *program, int num_groups) {
glDispatchCompute(num_groups, 1, 1);
}

void end_glsl_kernels(const std::vector<IOV> &iov) {
// glMemoryBarrier(GL_SHADER_STORAGE_BARRIER_BIT); // TODO(archibate): move to
// Program::synchroize()

for (int i = 0; i < ssbo.size(); i++) {
if (!iov[i].size)
continue;
void *p = ssbo[i].map(0, iov[i].size); // output
std::memcpy(iov[i].base, p, iov[i].size);
}
glUnmapBuffer(GL_SHADER_STORAGE_BUFFER);
ssbo.clear();
}

bool is_opengl_api_available() {
return true;
}
Expand All @@ -303,19 +300,14 @@ int opengl_get_threads_per_group() {
}

#else
void create_glsl_root_buffer(size_t size) {
TI_NOT_IMPLEMENTED
}
struct GLProgram {};
struct GLSLLauncherImpl {};

void begin_glsl_kernels(const std::vector<IOV> &iov) {
GLSLLauncher::GLSLLauncher(size_t size) {
TI_NOT_IMPLEMENTED
}

void end_glsl_kernels(const std::vector<IOV> &iov) {
TI_NOT_IMPLEMENTED
}

void launch_glsl_kernel(GLProgram *program, int num_groups) {
void CompiledGLSL::launch_glsl(int num_groups) const {
TI_NOT_IMPLEMENTED
}

Expand All @@ -325,14 +317,25 @@ bool is_opengl_api_available() {

void initialize_opengl(){TI_NOT_IMPLEMENTED}

GLProgram *compile_glsl_program(std::string source) {
CompiledGLSL::CompiledGLSL(const std::string &source) {
//: glsl(std::make_unique<GLProgram>()) {
TI_NOT_IMPLEMENTED
}

int opengl_get_threads_per_group() {
int opengl_get_threads_per_group(){TI_NOT_IMPLEMENTED}

GLSLLaunchGuard::GLSLLaunchGuard(GLSLLauncherImpl *impl,
const std::vector<IOV> &iov)
: impl(impl),
iov(iov){TI_NOT_IMPLEMENTED}

GLSLLaunchGuard::~GLSLLaunchGuard() {
TI_NOT_IMPLEMENTED
}
#endif

CompiledGLSL::~CompiledGLSL() = default;
GLSLLauncher::~GLSLLauncher() = default;

} // namespace opengl
TLANG_NAMESPACE_END
16 changes: 9 additions & 7 deletions taichi/backends/opengl/opengl_api.h
Original file line number Diff line number Diff line change
Expand Up @@ -5,23 +5,25 @@
#include <string>
#include <vector>

#include "opengl_kernel_util.h"
#include "opengl_kernel_launcher.h"

TLANG_NAMESPACE_BEGIN

namespace opengl {

struct GLProgram;
void initialize_opengl();
bool is_opengl_api_available();
void create_glsl_root_buffer(size_t size);
void begin_glsl_kernels(const std::vector<IOV> &iov);
void launch_glsl_kernel(GLProgram *program, int num_groups);
void end_glsl_kernels(const std::vector<IOV> &iov);
GLProgram *compile_glsl_program(std::string source);
int opengl_get_threads_per_group();
extern bool opengl_has_GL_NV_shader_atomic_float;

struct GLProgram;
struct CompiledGLSL {
std::unique_ptr<GLProgram> glsl;
CompiledGLSL(const std::string &source);
~CompiledGLSL();
void launch_glsl(int num_groups) const;
};

} // namespace opengl

TLANG_NAMESPACE_END
30 changes: 30 additions & 0 deletions taichi/backends/opengl/opengl_kernel_launcher.h
Original file line number Diff line number Diff line change
@@ -0,0 +1,30 @@
#pragma once

#include <vector>
#include "opengl_kernel_util.h"

TLANG_NAMESPACE_BEGIN

namespace opengl {

struct GLSLLauncherImpl;

struct GLSLLaunchGuard {
GLSLLauncherImpl *impl;
const std::vector<IOV> &iov;
GLSLLaunchGuard(GLSLLauncherImpl *impl, const std::vector<IOV> &iov);
~GLSLLaunchGuard();
};

struct GLSLLauncher {
std::unique_ptr<GLSLLauncherImpl> impl;
GLSLLauncher(size_t size);
~GLSLLauncher();
GLSLLaunchGuard create_launch_guard(const std::vector<IOV> &iov) {
return GLSLLaunchGuard(impl.get(), iov);
}
};

} // namespace opengl

TLANG_NAMESPACE_END
Loading

0 comments on commit 24e76a1

Please sign in to comment.