Skip to content

Commit

Permalink
[opengl] Remove listgen support (taichi-dev#3257)
Browse files Browse the repository at this point in the history 
* [opengl] Remove listgen support

* fix
  • Loading branch information
@k-ye
k-ye authored Oct 26, 2021
1 parent 3a9c1f3 commit f1fe2c6
Show file tree
Hide file tree
Showing 5 changed files with 17 additions and 147 deletions.
82 changes: 0 additions & 82 deletions taichi/backends/opengl/codegen_opengl.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,6 @@ namespace shaders {
#define TI_INSIDE_OPENGL_CODEGEN
#include "taichi/backends/opengl/shaders/atomics_macro_f32.glsl.h"
#include "taichi/backends/opengl/shaders/runtime.h"
#include "taichi/backends/opengl/shaders/listman.h"
#include "taichi/backends/opengl/shaders/random.glsl.h"
#include "taichi/backends/opengl/shaders/fast_pow.glsl.h"
#include "taichi/backends/opengl/shaders/print.glsl.h"
Expand Down Expand Up @@ -148,8 +147,6 @@ class KernelGen : public IRVisitor {
// clang-format off
if (used.print) // the runtime buffer is only used for print now..
line_appender_header_.append_raw(shaders::kOpenGlRuntimeSourceCode);
if (used.listman)
line_appender_header_.append_raw(shaders::kOpenGLListmanSourceCode);

std::string kernel_header;
#define DEFINE_LAYOUT(layout, restype, name, id, dt, dtype) \
Expand Down Expand Up @@ -949,24 +946,6 @@ class KernelGen : public IRVisitor {
emit("}}\n");
}

void generate_struct_for_kernel(OffloadedStmt *stmt) {
TI_ASSERT(stmt->task_type == OffloadedStmt::TaskType::struct_for);
used.listman = true;
const std::string glsl_kernel_name = make_kernel_name();
emit("void {}()", glsl_kernel_name);
this->glsl_kernel_name_ = glsl_kernel_name;
emit("{{ // struct for {}", stmt->snode->node_type_name);
{
ScopedIndent _s(line_appender_);
workgroup_size_ = stmt->block_dim;
num_workgroups_ = stmt->grid_dim;
ScopedGridStrideLoop _gsl(this, "_list_len_");
emit("int _itv = _list_[_sid];");
stmt->body->accept(this);
}
emit("}}\n");
}

size_t get_snode_base_address(const SNode *snode) {
if (snode->type == SNodeType::root)
return 0;
Expand All @@ -984,57 +963,6 @@ class KernelGen : public IRVisitor {
return addr;
}

void generate_listgen_for_dynamic(const SNode *snode) {
TI_ASSERT(snode->type == SNodeType::dynamic);
// the `length` field of a dynamic SNode is at it's end:
// | x[0] | x[1] | x[2] | x[3] | ... | len |
TI_ASSERT_INFO(snode->parent->type == SNodeType::root,
"Non-top-level dynamic not supported yet on OpenGL");
size_t addr = get_snode_meta_address(snode);
used.int32 = true;
emit("_list_len_ = _data_i32_[{} >> 2];", addr);
emit("for (int i = 0; i < _list_len_; i++) {{");
{
ScopedIndent _s(line_appender_);
emit("_list_[i] = i;");
}
emit("}}");
}

void generate_listgen_for_dense(const SNode *snode) {
TI_ASSERT(snode->type == SNodeType::dense);
// the `length` field of a dynamic SNode is at it's end:
// | x[0] | x[1] | x[2] | x[3] | ... | len |
emit("_list_len_ = {};",
struct_compiled_->snode_map.at(snode->node_type_name).length);
emit("for (int i = 0; i < _list_len_; i++) {{");
{
ScopedIndent _s(line_appender_);
emit("_list_[i] = i;");
}
emit("}}");
}

void generate_listgen_kernel(OffloadedStmt *stmt) {
TI_ASSERT(stmt->task_type == OffloadedStmt::TaskType::listgen);
const std::string glsl_kernel_name = make_kernel_name();
emit("void {}()", glsl_kernel_name);
this->glsl_kernel_name_ = glsl_kernel_name;
used.listman = true;
emit("{{ // listgen {}", stmt->snode->node_type_name);
{
ScopedIndent _s(line_appender_);
if (stmt->snode->type == SNodeType::dense) {
generate_listgen_for_dense(stmt->snode);
} else if (stmt->snode->type == SNodeType::dynamic) {
generate_listgen_for_dynamic(stmt->snode);
} else {
TI_NOT_IMPLEMENTED
}
}
emit("}}\n");
}

void visit(GlobalTemporaryStmt *stmt) override {
TI_ASSERT(stmt->width() == 1);
used.buf_gtmp = true;
Expand Down Expand Up @@ -1113,10 +1041,6 @@ class KernelGen : public IRVisitor {
generate_serial_kernel(stmt);
} else if (stmt->task_type == Type::range_for) {
generate_range_for_kernel(stmt);
} else if (stmt->task_type == Type::struct_for) {
generate_struct_for_kernel(stmt);
} else if (stmt->task_type == Type::listgen) {
generate_listgen_kernel(stmt);
} else {
// struct_for is automatically lowered to ranged_for for dense snodes
// (#378). So we only need to support serial and range_for tasks.
Expand All @@ -1131,12 +1055,6 @@ class KernelGen : public IRVisitor {
TI_ERROR("[glsl] Struct for cannot be nested under OpenGL for now");
}

void visit(ClearListStmt *stmt) override {
used.listman = true;
emit("// clear list {}", stmt->snode->node_type_name);
emit("_list_len_ = 0;");
}

void visit(IfStmt *if_stmt) override {
emit("if ({} != 0) {{", if_stmt->cond->short_name());
if (if_stmt->true_statements) {
Expand Down
12 changes: 1 addition & 11 deletions taichi/backends/opengl/opengl_api.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -7,7 +7,6 @@
#include "taichi/program/py_print_buffer.h"
#include "taichi/util/environ_config.h"
#include "taichi/backends/opengl/shaders/runtime.h"
#include "taichi/backends/opengl/shaders/listman.h"
#include "taichi/ir/transforms.h"

#ifdef TI_WITH_OPENGL
Expand Down Expand Up @@ -50,17 +49,14 @@ static std::string add_line_markers(std::string x) {
struct OpenGlRuntimeImpl {
struct {
DeviceAllocation runtime = kDeviceNullAllocation;
DeviceAllocation listman = kDeviceNullAllocation;
DeviceAllocation root = kDeviceNullAllocation;
DeviceAllocation gtmp = kDeviceNullAllocation;
} core_bufs;

OpenGlRuntimeImpl() {
}

std::unique_ptr<GLSLRuntime> runtime;
std::unique_ptr<GLSLListman> listman;

std::unique_ptr<GLSLRuntime> runtime{nullptr};
std::vector<std::unique_ptr<DeviceCompiledProgram>> programs;
};

Expand Down Expand Up @@ -319,7 +315,6 @@ void DeviceCompiledProgram::launch(Context &ctx, OpenGlRuntime *runtime) const {
auto binder = compiled_pipeline_[i]->resource_binder();
auto &core_bufs = runtime->impl->core_bufs;
binder->buffer(0, int(GLBufId::Runtime), core_bufs.runtime);
binder->buffer(0, int(GLBufId::Listman), core_bufs.listman);
binder->buffer(0, int(GLBufId::Root), core_bufs.root);
binder->buffer(0, int(GLBufId::Gtmp), core_bufs.gtmp);
if (program_.args_buf_size)
Expand Down Expand Up @@ -408,17 +403,12 @@ OpenGlRuntime::OpenGlRuntime() {
impl->core_bufs.runtime = device->allocate_memory(
{sizeof(GLSLRuntime), /*host_write=*/false, /*host_read=*/true});

impl->listman = std::make_unique<GLSLListman>();
impl->core_bufs.listman = device->allocate_memory({sizeof(GLSLListman)});

impl->core_bufs.gtmp =
device->allocate_memory({taichi_global_tmp_buffer_size});

auto cmdlist = device->get_compute_stream()->new_command_list();
cmdlist->buffer_fill(impl->core_bufs.runtime.get_ptr(0), sizeof(GLSLRuntime),
0);
cmdlist->buffer_fill(impl->core_bufs.listman.get_ptr(0), sizeof(GLSLListman),
0);
cmdlist->buffer_fill(impl->core_bufs.gtmp.get_ptr(0),
taichi_global_tmp_buffer_size, 0);
device->get_compute_stream()->submit_synced(cmdlist.get());
Expand Down
15 changes: 7 additions & 8 deletions taichi/backends/opengl/opengl_kernel_util.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -6,7 +6,8 @@

#include "taichi/ir/snode.h"

TLANG_NAMESPACE_BEGIN
namespace taichi {
namespace lang {

class SNode;

Expand All @@ -33,7 +34,6 @@ struct UsedFeature {

// utilties:
bool fast_pow{false};
bool listman{false};
bool random{false};
bool print{false};
bool reduction{false};
Expand All @@ -46,19 +46,18 @@ struct UsedFeature {

enum class GLBufId {
Root = 0,
Runtime = 6,
Listman = 7,
Gtmp = 1,
Args = 2,
Retr = 3,
Extr = 4,
Runtime = 5,
};

struct IOV {
void *base;
size_t size;
void *base{nullptr};
size_t size{0};
};

} // namespace opengl

TLANG_NAMESPACE_END
} // namespace lang
} // namespace taichi
29 changes: 0 additions & 29 deletions taichi/backends/opengl/shaders/listman.h
View file

This file was deleted.

26 changes: 9 additions & 17 deletions tests/python/test_dynamic.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -3,15 +3,7 @@
import taichi as ti


def ti_support_dynamic(test):
return ti.archs_excluding(ti.cc, ti.vulkan)(test)


def ti_support_non_top_dynamic(test):
return ti.archs_excluding(ti.opengl, ti.cc, ti.vulkan)(test)


@ti.test(exclude=[ti.cc, ti.vulkan])
@ti.test(require=ti.extension.sparse)
def test_dynamic():
x = ti.field(ti.f32)
n = 128
Expand All @@ -29,7 +21,7 @@ def func():
assert x[i] == i


@ti.test(exclude=[ti.cc, ti.vulkan])
@ti.test(require=ti.extension.sparse)
def test_dynamic2():
x = ti.field(ti.f32)
n = 128
Expand All @@ -47,7 +39,7 @@ def func():
assert x[i] == i


@ti.test(exclude=[ti.cc, ti.vulkan])
@ti.test(require=ti.extension.sparse)
def test_dynamic_matrix():
x = ti.Matrix.field(2, 1, dtype=ti.i32)
n = 8192
Expand All @@ -70,7 +62,7 @@ def func():
assert b == 0


@ti.test(exclude=[ti.cc, ti.vulkan])
@ti.test(require=ti.extension.sparse)
def test_append():
x = ti.field(ti.i32)
n = 128
Expand All @@ -92,7 +84,7 @@ def func():
assert elements[i] == i


@ti.test(exclude=[ti.cc, ti.vulkan])
@ti.test(require=ti.extension.sparse)
def test_length():
x = ti.field(ti.i32)
y = ti.field(ti.f32, shape=())
Expand All @@ -116,7 +108,7 @@ def get_len():
assert y[None] == n


@ti.test(exclude=[ti.cc, ti.vulkan])
@ti.test(require=ti.extension.sparse)
def test_append_ret_value():
x = ti.field(ti.i32)
y = ti.field(ti.i32)
Expand All @@ -141,7 +133,7 @@ def func():
assert x[i] + 3 == z[i]


@ti.test(exclude=[ti.opengl, ti.cc, ti.vulkan])
@ti.test(require=ti.extension.sparse)
def test_dense_dynamic():
# The spin lock implementation has triggered a bug in CUDA, the end result
# being that appending to Taichi's dynamic node messes up its length. See
Expand Down Expand Up @@ -172,7 +164,7 @@ def func():
assert l[i] == n


@ti.test(exclude=[ti.opengl, ti.cc, ti.vulkan])
@ti.test(require=ti.extension.sparse)
def test_dense_dynamic_len():
n = 128
x = ti.field(ti.i32)
Expand All @@ -191,7 +183,7 @@ def func():
assert l[i] == 0


@ti.test(exclude=[ti.cc, ti.vulkan])
@ti.test(require=ti.extension.sparse)
def test_dynamic_activate():
ti.init(arch=ti.metal)
# record the lengths
Expand Down

0 comments on commit f1fe2c6

Please sign in to comment.