Skip to content

Commit

Permalink
added initialization with compute shader, fixed resizing issue and mi…
Browse files Browse the repository at this point in the history
…nor GL things
  • Loading branch information
theoden8 committed Oct 12, 2023
1 parent ca27682 commit cf739eb
Show file tree
Hide file tree
Showing 11 changed files with 266 additions and 82 deletions.
1 change: 1 addition & 0 deletions AutomatonApp.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -38,6 +38,7 @@ class AutomatonApp {
template <typename AUT>
void run(AUT &&aut, const AutOptions &opts) {
AutomatonApp &app = (*this);
w.update_size();
constexpr storage_mode storage_mode_recommended = ::use_storage_mode<AUT>::smode;
if constexpr(storage_mode_recommended == storage_mode::HOSTBUFFER) {
run_with_storage_mode<storage_mode::HOSTBUFFER>(std::forward<AUT>(aut), opts);
Expand Down
4 changes: 1 addition & 3 deletions CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -17,12 +17,10 @@ message("C++ Compiler ${CMAKE_CXX_COMPILER}")

if(UNIX AND NOT APPLE)
set(CMAKE_CXX_FLAGS "-std=c++20 -fopenmp -Wall -Wextra -Wno-unused-variable -Wno-deprecated-enum-enum-conversion -Wno-unused-parameter")
else()
if(WIN32)
elseif(WIN32)
set(CMAKE_CXX_FLAGS "/std:c++20")
else()
set(CMAKE_CXX_FLAGS "-std=c++20 -Wall -Wextra")
endif(WIN32)
endif()

set(exec automaton)
Expand Down
17 changes: 16 additions & 1 deletion File.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -17,6 +17,10 @@
#include <windows.h>
#endif

#if !defined(_POSIX_VERSION)
#include <direct.h>
#endif

namespace sys {

struct Path {
Expand Down Expand Up @@ -49,6 +53,17 @@ struct Path {
}
};

std::string get_cwd() {
#if defined(_POSIX_VERSION)
char buf[PATH_MAX + 1];
getcwd(buf, PATH_MAX);
#else
char buf[MAX_PATH + 1];
_getcwd(buf, MAX_PATH);
#endif
return buf;
}

std::string get_executable_directory(int argc, char *argv[]) {
#ifdef __linux__
std::vector<char> buf(PATH_MAX);
Expand Down Expand Up @@ -124,7 +139,7 @@ class File {
}
};
public:
File(const char *filename):
explicit File(const char *filename):
filename(filename)
{}

Expand Down
9 changes: 9 additions & 0 deletions Logger.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -66,6 +66,15 @@ class Logger {
instance->Write(fmt, argptr);
va_end(argptr);
}
static void Debug(const char *fmt, ...) {
#ifndef NDEBUG
ASSERT(instance != nullptr);
va_list argptr;
va_start(argptr, fmt);
instance->WriteFmt("DEBG: ", fmt, argptr);
va_end(argptr);
#endif
}
static void Info(const char *fmt, ...) {
ASSERT(instance != nullptr);
va_list argptr;
Expand Down
151 changes: 105 additions & 46 deletions Renderer.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -72,10 +72,12 @@ struct TexturedGridRenderer {
// init shader program
ShaderProgram::init(prog, vao, {"attrVertex"});
prog.assign_uniforms(uSampler, uNstates, uColorscheme);
init_textures(w.width(), w.height(), factor);
set_grid_size(w.width(), w.height(), factor);
init_textures();
}

virtual void init_textures(int w_, int h_, int zoom=0, const char *filename=nullptr) = 0;
virtual void set_grid_size(int w_, int h_, int zoom) = 0;
virtual void init_textures(const char *filename=nullptr) = 0;
virtual void update_state() = 0;
virtual GLuint get_current_texture_id() = 0;

Expand Down Expand Up @@ -104,6 +106,7 @@ struct TexturedGridRenderer {
ShaderAttrib::clear(attrVertex);
gl::VertexArray::clear(vao);
ShaderProgram::clear(prog);
ShaderProgram::unassign_uniforms(uSampler, uNstates, uColorscheme);
}
};

Expand Down Expand Up @@ -138,7 +141,7 @@ struct Renderer<AUT, storage_mode::HOSTBUFFER, AccessMode> : public TexturedGrid
tw(0), th(0)
{}

void init_textures(int w_, int h_, int zoom=0, const char *filename=nullptr) override {
void set_grid_size(int w_, int h_, int zoom) override {
if(zoom==0)zoom=1;
if(zoom > 0) {
w_ /= zoom, h_ /= zoom;
Expand All @@ -148,10 +151,14 @@ struct Renderer<AUT, storage_mode::HOSTBUFFER, AccessMode> : public TexturedGrid
w_ *= -zoom, h_ *= -zoom;
}
w=w_,h=h_;
Logger::Info("[%d %d] [%d %d]\n", w,h,tw,th);
if(zoom < 0) {
parent_t::colorscheme = 1;
no_states = std::min<int>(256, (w/tw) * (h/th) * (aut.no_states - 1) + 1);
}
Logger::Info("[%d %d] [%d %d]\n", w,h,tw,th);
}

void init_textures(const char *filename=nullptr) override {
if(w!=tw||h!=th||extrabuf) {
extrabuf = true;
finalbuf.init(tw, th);
Expand Down Expand Up @@ -219,16 +226,17 @@ struct Renderer<AUT, storage_mode::HOSTBUFFER, AccessMode> : public TexturedGrid
int per_x = w / tw;
int per_y = h / th;
const int area = per_x * per_y;
const float scale_states = fmax(1, float(area * (aut.no_states - 1) + 1) / no_states);
#pragma omp parallel for
for(int i = 0; i < tw*th; ++i) {
uint32_t sum = 0;
uint16_t sum = 0;
int y = i / tw, x = i % tw;
for(int iy = 0; iy < per_y; ++iy) {
for(int ix = 0; ix < per_x; ++ix) {
sum += srcbuf->buffer[(i / tw * per_y + iy) * w + (i % tw) * per_x + ix];
sum += srcbuf->buffer[(y*per_y+iy)*w + x*per_x+ix];
}
}
finalbuf.buffer[i] = std::round<uint8_t>(float(sum) / area);
finalbuf.buffer[i] = std::round<uint8_t>(float(sum) / scale_states - .01);
}
srcbuf = &finalbuf;
}
Expand Down Expand Up @@ -273,16 +281,19 @@ struct Renderer<ca::BSC, storage_mode::TEXTURES, AccessMode> : public TexturedGr
glm::ivec2 wg_per_cell = glm::ivec2(1, 1);
bool largetexture = false;

gl::Uniform<gl::UniformType::SAMPLER2D> uInitTex;
gl::Uniform<gl::UniformType::UINTEGER> uNStates, uSeed;
gl::ShaderProgram<gl::ComputeShader> computeInitSoup;
gl::Uniform<gl::UniformType::SAMPLER2D> uSrcTex, uDstTex;
gl::Uniform<gl::UniformType::UINTEGER> uBs, uSs, uC;
gl::Uniform<gl::UniformType::IVEC2> uSize, uWgPerCell;
gl::Uniform<gl::UniformType::UINTEGER> uAccessMode;
gl::ShaderProgram<gl::ComputeShader> compute;
gl::ShaderProgram<gl::ComputeShader> computeUpdate;
static constexpr int local_size = 8;
const int max_wg_invocations;
glm::ivec2 wg_size = glm::ivec2(0, 0);

using ShaderProgramCompute = decltype(compute);
using ShaderProgramCompute = decltype(computeUpdate);

storage_mode get_storage_mode() override {
return storage_mode::TEXTURES;
Expand All @@ -291,15 +302,17 @@ struct Renderer<ca::BSC, storage_mode::TEXTURES, AccessMode> : public TexturedGr
explicit Renderer(AUT &_aut, const std::string &dir):
parent_t(_aut.no_states, dir),
aut(_aut),
uSrcTex("srcbuf"s), uDstTex("dstbuf"s),
uInitTex("initTex"s), uNStates("n_states"), uSeed("seed"),
computeInitSoup({std::string(sys::Path(dir) / sys::Path("shaders"s) / sys::Path("soup.comp"s))}),
uSrcTex("srcTex"s), uDstTex("dstTex"s),
uBs("bs"s), uSs("ss"s), uC("c"s),
uSize("size"s), uWgPerCell("wg_per_cell"),
uAccessMode("access_mode"s),
compute({std::string(sys::Path(dir) / sys::Path("shaders"s) / sys::Path("bsc.comp"s))}),
max_wg_invocations(compute.get_max_wg_invocations())
computeUpdate({std::string(sys::Path(dir) / sys::Path("shaders"s) / sys::Path("bsc.comp"s))}),
max_wg_invocations(ShaderProgramCompute::get_max_wg_invocations())
{}

inline void set_grid_size(int w_, int h_, int zoom) {
void set_grid_size(int w_, int h_, int zoom) override {
w=w_,h=h_;
if(!zoom)zoom=1;
if(zoom > 0) {
Expand All @@ -309,10 +322,14 @@ struct Renderer<ca::BSC, storage_mode::TEXTURES, AccessMode> : public TexturedGr
}
Logger::Info("[w %d, h %d]\n", w, h);
largetexture = (zoom < 0);
if(zoom < 0) {
parent_t::colorscheme = 1;
}
set_work_group_sizes();
}

inline void set_work_group_sizes() {
const glm::ivec2 max_wg_size = compute.get_max_wgsize();
void set_work_group_sizes() {
const glm::ivec2 max_wg_size = ShaderProgramCompute::get_max_wgsize();
glm::ivec2 max_invocations = max_wg_size * local_size;
wg_per_cell = (glm::ivec2(w, h) + max_invocations - 1) / max_invocations + 1;
// Logger::Info("[max_invocations/w = %d/%d = %.2f wg_per_cell_x %d]\n", max_invocations, w, float(max_invocations) / float(w), wg_per_cell.x);
Expand All @@ -322,41 +339,77 @@ struct Renderer<ca::BSC, storage_mode::TEXTURES, AccessMode> : public TexturedGr
Logger::Info("[wg %dx%dx%d %dx%dx%d]\n", wg_size.x, local_size, wg_per_cell.x, wg_size.y, local_size, wg_per_cell.y);
}

void init_textures(int w_, int h_, int zoom=0, const char *filename=nullptr) override {
set_grid_size(w_, h_, zoom);
set_work_group_sizes();
if(zoom < 0) {
parent_t::colorscheme = 1;
}
std::vector<uint8_t> buf(w * h, 0);
#pragma omp parallel for
for(int i = 0; i < w * h; ++i) {
buf[i] = aut.init_state(i/w, i%w);
}
ShaderProgramCompute::compile_program(compute);
compute.assign_uniforms(
uSrcTex, uDstTex,
uBs, uSs, uC,
uSize, uWgPerCell,
uAccessMode
);
#define COMPUTE_INIT_SOUP
void init_textures(const char *filename=nullptr) override {
for(GLuint *tex_ptr : {&tex1, &tex2}) {
GLuint &tex = *tex_ptr;
gl::Texture<GL_TEXTURE_2D>::init(tex);
gl::Texture<GL_TEXTURE_2D>::bind(tex);
glTexImage2D(GL_TEXTURE_2D, 0, GL_R8UI, w, h, 0, GL_RED_INTEGER, GL_UNSIGNED_BYTE, buf.data()); GLERROR
if(tex_ptr == &tex1 && filename != nullptr) {
RenderStorage<storage_mode::HOSTBUFFER> buf;
buf.init(w, h);
// #pragma omp parallel for
// for(int i = 0; i < w * h; ++i) {
// buf.buffer[i] = aut.init_state(i/w, i%w);
// }
RLEDecoder<RenderStorage<storage_mode::HOSTBUFFER>>::read(filename, buf);
glTexImage2D(GL_TEXTURE_2D, 0, GL_R8UI, w, h, 0, GL_RED_INTEGER, GL_UNSIGNED_BYTE, buf.buffer.data()); GLERROR
} else {
glTexImage2D(GL_TEXTURE_2D, 0, GL_R8UI, w, h, 0, GL_RED_INTEGER, GL_UNSIGNED_BYTE, nullptr); GLERROR
}
gl::Texture<GL_TEXTURE_2D>::param(GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE);
gl::Texture<GL_TEXTURE_2D>::param(GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE);
gl::Texture<GL_TEXTURE_2D>::param(GL_TEXTURE_MAG_FILTER, GL_NEAREST);
gl::Texture<GL_TEXTURE_2D>::param(GL_TEXTURE_MIN_FILTER, GL_NEAREST);
gl::Texture<GL_TEXTURE_2D>::unbind();
}
compute.print_compute_capabilities();
//#ifdef COMPUTE_INIT_SOUP
if(filename == nullptr) {
ShaderProgramCompute::compile_program(computeInitSoup);
computeInitSoup.assign_uniforms(
uInitTex, uNStates,
uSize, uWgPerCell, uSeed
);
init_state_soup();
ShaderProgramCompute::clear(computeInitSoup);
ShaderProgramCompute::unassign_uniforms(
uInitTex, uNStates,
uSize, uWgPerCell, uSeed
);
}
//#endif
ShaderProgramCompute::compile_program(computeUpdate);
computeUpdate.assign_uniforms(
uSrcTex, uDstTex,
uBs, uSs, uC,
uSize, uWgPerCell,
uAccessMode
);
ShaderProgramCompute::print_compute_capabilities();
}

void set_data_compute() {
uSrcTex.set_data(current_tex ? 0 : 1);
uDstTex.set_data(current_tex ? 1 : 0);
void set_data_compute_init_soup() {
uInitTex.set_data(0);
uNStates.set_data(aut.no_states);
glm::ivec2 val_size(w, h);
uSize.set_data(val_size);
uWgPerCell.set_data(wg_per_cell);
uSeed.set_data(rand());
}

void init_state_soup() {
ShaderProgramCompute::use(computeInitSoup);
set_data_compute_init_soup();
GLuint inittex = tex1;
glBindImageTexture(0, inittex, 0, GL_FALSE, 0, GL_WRITE_ONLY, GL_R8UI); GLERROR
ShaderProgramCompute::dispatch(wg_size.x, wg_size.y, 1);
ShaderProgramCompute::barrier(GL_TEXTURE_UPDATE_BARRIER_BIT);
ShaderProgramCompute::unuse();
}

void set_data_compute_update() {
uSrcTex.set_data(0);
uDstTex.set_data(1);
int bs = int(aut.bs_bitmask.to_ulong());;
uBs.set_data(bs);
int ss = int(aut.ss_bitmask.to_ulong());;
Expand All @@ -369,15 +422,15 @@ struct Renderer<ca::BSC, storage_mode::TEXTURES, AccessMode> : public TexturedGr
}

void update_state() override {
ShaderProgramCompute::use(compute);
set_data_compute();
GLuint srctex = current_tex?tex1:tex2,
dsttex = current_tex?tex2:tex1;
ShaderProgramCompute::use(computeUpdate);
set_data_compute_update();
GLuint srctex = current_tex?tex2:tex1,
dsttex = current_tex?tex1:tex2;
glBindImageTexture(0, srctex, 0, GL_FALSE, 0, GL_READ_ONLY, GL_R8UI); GLERROR
glBindImageTexture(1, dsttex, 0, GL_FALSE, 0, GL_WRITE_ONLY, GL_R8UI); GLERROR
compute.dispatch(wg_size.x, wg_size.y, 1);
compute.barrier(GL_SHADER_IMAGE_ACCESS_BARRIER_BIT | GL_TEXTURE_UPDATE_BARRIER_BIT);
//compute.barrier(GL_ALL_BARRIER_BITS); GLERROR
ShaderProgramCompute::dispatch(wg_size.x, wg_size.y, 1);
ShaderProgramCompute::barrier(GL_SHADER_IMAGE_ACCESS_BARRIER_BIT | GL_TEXTURE_UPDATE_BARRIER_BIT);
//computeUpdate.barrier(GL_ALL_BARRIER_BITS); GLERROR
//glFinish(); GLERROR
ShaderProgramCompute::unuse();
current_tex = current_tex ? 0 : 1;
Expand All @@ -390,7 +443,13 @@ struct Renderer<ca::BSC, storage_mode::TEXTURES, AccessMode> : public TexturedGr
void clear() override {
gl::Texture<GL_TEXTURE_2D>::clear(tex1);
gl::Texture<GL_TEXTURE_2D>::clear(tex2);
ShaderProgramCompute::clear(compute);
ShaderProgramCompute::clear(computeUpdate);
ShaderProgramCompute::unassign_uniforms(
uSrcTex, uDstTex,
uBs, uSs, uC,
uSize, uWgPerCell,
uAccessMode
);
parent_t::clear();
}
};
11 changes: 11 additions & 0 deletions ShaderProgram.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -184,6 +184,17 @@ class ShaderProgram {
unroll(assign_uniform(uniforms)...);
}

template <typename ShaderUniformT>
static int unassign_uniform(ShaderUniformT &uniform) {
uniform.unset_id();
return 0;
}

template <typename... ShaderUniformTs>
static void unassign_uniforms(ShaderUniformTs& ...uniforms) {
unroll(unassign_uniform(uniforms)...);
}

static void dispatch(size_t x, size_t y, size_t z) {
glDispatchCompute(x, y, z); GLERROR
}
Expand Down
Loading

0 comments on commit cf739eb

Please sign in to comment.