Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Basic WebGL Backend #672

Merged
merged 48 commits into from
Jan 20, 2018
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
48 commits
Select commit Hold shift + click to select a range
c900170
OpenGL backend skeleton
phisiart Nov 5, 2017
12b1fd7
Extend OpenGL backend to run "empty" functions.
phisiart Nov 19, 2017
527a89d
Forgot build_opengl.cc; Modify glfw3 cmake for Linux.
phisiart Nov 21, 2017
fcc696e
Very naive OpenGL codegen just to get one demo correct.
phisiart Nov 22, 2017
0ee9db1
More log.
phisiart Nov 22, 2017
db01e8c
Add opengl schedule. This schedule just fuses all dimensions into one…
phisiart Nov 25, 2017
edd2dd5
OpenGL codegen checks itervar binding.
phisiart Nov 25, 2017
2016c68
change opencl text to opengl
Nov 23, 2017
e9f90cf
import glfw and glad in cmakelist.txt
Nov 24, 2017
917f8a8
copy opengl initialization code into opengl_device_api.cc
Nov 24, 2017
fa7aff5
add memory allocation and memory copy code
Nov 24, 2017
cd918c1
use tvm logs
Nov 24, 2017
f599fd9
partially add create program functions
Nov 24, 2017
102e9b4
short circuit threadIdx in codegen for now
Nov 24, 2017
745ab5e
add render function
Nov 25, 2017
be22151
introduct Texture class
Nov 25, 2017
a15201f
introduce Program class
Nov 25, 2017
10a28e0
final touch of the runtime
Nov 25, 2017
a94e897
change glad to be submodule
phisiart Nov 25, 2017
9a77ad6
Format OpenGL code.
phisiart Nov 25, 2017
d1b1eb3
Use dmlc/dlpack again. We should send a PR to dmlc/dlpack later.
phisiart Nov 26, 2017
bec8cfc
Remove OpenGLWorkspace::Init(). Create glfw invisible window.
phisiart Nov 26, 2017
e6bf910
Several cleanups.
phisiart Nov 26, 2017
f59aa64
Support putting partial texture data.
phisiart Nov 26, 2017
c2948e7
GLSL output now has type float.
phisiart Nov 26, 2017
cfdf727
Slightly improve comments.
phisiart Nov 26, 2017
871861d
Some fixes according to review.
phisiart Dec 15, 2017
5b78914
remove const of PrintType
phisiart Dec 15, 2017
55147cb
Fix code style according to lint.
phisiart Dec 15, 2017
4d55284
Trying to add OpenGL runtime to emscripten'ed web runtime.
phisiart Dec 27, 2017
ad6c8dc
remove glad, add temporary rpc opengl test
phisiart Jan 2, 2018
f17ac76
Add 'type' parameter to AllocDataSpace.
phisiart Jan 5, 2018
4ee83ce
Improve OpenGL texture. Now test_runtime_ndarray.py passes.
phisiart Jan 6, 2018
9df0c34
- address review comments;
phisiart Jan 12, 2018
a0662a3
add remote.download("myadd.tvm_meta.json")
phisiart Jan 12, 2018
c7c30f0
Now tests/webgl/test_remote_save_load.py succeeds.
phisiart Jan 13, 2018
c0b04a8
Savepoint that works.
phisiart Jan 13, 2018
da21d5c
Remove temporary RPC test.
phisiart Jan 13, 2018
3d669fd
Correctly handle OpenGL argument.
phisiart Jan 16, 2018
3c90dde
Change emcc optmization flag back to -Oz.
phisiart Jan 17, 2018
d390676
Support reduction. Now this program runs correctly:
phisiart Jan 17, 2018
96145f9
Generate tvm_get_texel function in GLSL for cleaner code.
phisiart Jan 18, 2018
d289979
Add tests/webgl/test_local_gemm.py
phisiart Jan 18, 2018
20087e3
Directly return if threadIdx.x < thread_extent.
phisiart Jan 18, 2018
8dbd632
Cleanup OpenGLArgKind and OpenGLShader.
phisiart Jan 19, 2018
2acb00e
Enable local OpenGL tests in cpu & gpu Jenkins.
phisiart Jan 19, 2018
d30d6a5
Cleanup.
phisiart Jan 19, 2018
8b98b59
Address review comments.
phisiart Jan 20, 2018
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
18 changes: 16 additions & 2 deletions CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -22,6 +22,7 @@ endif()

tvm_option(USE_CUDA "Build with CUDA" OFF)
tvm_option(USE_OPENCL "Build with OpenCL" OFF)
tvm_option(USE_OPENGL "Build with OpenGL" OFF)
tvm_option(USE_METAL "Build with Metal" OFF)
tvm_option(USE_RPC "Build with RPC" ON)
tvm_option(USE_GRAPH_RUNTIME "Build with tiny graph runtime" ON)
Expand Down Expand Up @@ -61,8 +62,8 @@ if(MSVC)
else(MSVC)
include(CheckCXXCompilerFlag)
check_cxx_compiler_flag("-std=c++11" SUPPORT_CXX11)
set(CMAKE_C_FLAGS "-O3 -Wall -std=c++11 -fPIC")
set(CMAKE_CXX_FLAGS ${CMAKE_C_FLAGS})
set(CMAKE_C_FLAGS "-O3 -Wall -fPIC")
set(CMAKE_CXX_FLAGS "${CMAKE_C_FLAGS} -std=c++11")
endif(MSVC)

# add source group
Expand All @@ -87,6 +88,7 @@ file(GLOB RUNTIME_SRCS src/runtime/*.cc)
file(GLOB COMPILER_LLVM_SRCS src/codegen/llvm/*.cc)
file(GLOB RUNTIME_CUDA_SRCS src/runtime/cuda/*.cc)
file(GLOB RUNTIME_OPENCL_SRCS src/runtime/opencl/*.cc)
file(GLOB RUNTIME_OPENGL_SRCS src/runtime/opengl/*.cc)
file(GLOB RUNTIME_METAL_SRCS src/runtime/metal/*.mm)
file(GLOB RUNTIME_RPC_SRCS src/runtime/rpc/*.cc)
file(GLOB RUNTIME_GRAPH_SRCS src/runtime/graph/*.cc)
Expand Down Expand Up @@ -135,6 +137,18 @@ else(USE_OPENCL)
add_definitions(-DTVM_OPENCL_RUNTIME=0)
endif(USE_OPENCL)

if(USE_OPENGL)
find_package(OpenGL QUIET REQUIRED)
find_package(glfw3 QUIET REQUIRED)
message(STATUS "Build with OpenGL support")
include_directories(${OPENGL_INCLUDE_DIRS})
list(APPEND TVM_RUNTIME_LINKER_LIBS ${OpenGL_LIBRARIES} glfw)
list(APPEND RUNTIME_SRCS ${RUNTIME_OPENGL_SRCS})
add_definitions(-DTVM_OPENGL_RUNTIME=1)
else(USE_OPENGL)
add_definitions(-DTVM_OPENGL_RUNTIME=0)
endif(USE_OPENGL)

if(USE_METAL)
find_package(OpenCL QUIET REQUIRED)
message(STATUS "Build with Metal support")
Expand Down
2 changes: 2 additions & 0 deletions Jenkinsfile
Original file line number Diff line number Diff line change
Expand Up @@ -88,6 +88,7 @@ stage('Build') {
echo USE_CUDNN=1 >> config.mk
echo USE_CUDA=1 >> config.mk
echo USE_OPENCL=1 >> config.mk
echo USE_OPENGL=1 >> config.mk
echo LLVM_CONFIG=llvm-config-4.0 >> config.mk
echo USE_RPC=1 >> config.mk
echo USE_GRAPH_RUNTIME=1 >> config.mk
Expand Down Expand Up @@ -120,6 +121,7 @@ stage('Build') {
echo USE_CUDA=0 >> config.mk
echo USE_OPENCL=0 >> config.mk
echo USE_RPC=0 >> config.mk
echo USE_OPENGL=1 >> config.mk
echo LLVM_CONFIG=llvm-config-4.0 >> config.mk
"""
make('cpu', '-j2')
Expand Down
17 changes: 16 additions & 1 deletion Makefile
Original file line number Diff line number Diff line change
Expand Up @@ -32,8 +32,8 @@ OBJCFLAGS = -fno-objc-arc
EMCC_FLAGS= -std=c++11 -DDMLC_LOG_STACK_TRACE=0\
-Oz -s RESERVED_FUNCTION_POINTERS=2 -s MAIN_MODULE=1 -s NO_EXIT_RUNTIME=1\
-s EXTRA_EXPORTED_RUNTIME_METHODS="['cwrap','getValue','setValue','addFunction']"\
-s USE_GLFW=3 -s USE_WEBGL2=1 -lglfw\
$(INCLUDE_FLAGS)

# llvm configuration
ifdef LLVM_CONFIG
LLVM_VERSION=$(shell $(LLVM_CONFIG) --version| cut -b 1,3)
Expand All @@ -54,6 +54,7 @@ METAL_SRC = $(wildcard src/runtime/metal/*.mm)
CUDA_SRC = $(wildcard src/runtime/cuda/*.cc)
ROCM_SRC = $(wildcard src/runtime/rocm/*.cc)
OPENCL_SRC = $(wildcard src/runtime/opencl/*.cc)
OPENGL_SRC = $(wildcard src/runtime/opengl/*.cc)
RPC_SRC = $(wildcard src/runtime/rpc/*.cc)
GRAPH_SRC = $(wildcard src/runtime/graph/*.cc)
RUNTIME_SRC = $(wildcard src/runtime/*.cc)
Expand All @@ -65,6 +66,7 @@ METAL_OBJ = $(patsubst src/%.mm, build/%.o, $(METAL_SRC))
CUDA_OBJ = $(patsubst src/%.cc, build/%.o, $(CUDA_SRC))
ROCM_OBJ = $(patsubst src/%.cc, build/%.o, $(ROCM_SRC))
OPENCL_OBJ = $(patsubst src/%.cc, build/%.o, $(OPENCL_SRC))
OPENGL_OBJ = $(patsubst src/%.cc, build/%.o, $(OPENGL_SRC))
RPC_OBJ = $(patsubst src/%.cc, build/%.o, $(RPC_SRC))
GRAPH_OBJ = $(patsubst src/%.cc, build/%.o, $(GRAPH_SRC))
CC_OBJ = $(patsubst src/%.cc, build/%.o, $(CC_SRC)) $(LLVM_OBJ)
Expand Down Expand Up @@ -119,6 +121,19 @@ else
CFLAGS += -DTVM_OPENCL_RUNTIME=0
endif

ifeq ($(USE_OPENGL), 1)
CFLAGS += -DTVM_OPENGL_RUNTIME=1
EMCC_FLAGS += -DTVM_OPENGL_RUNTIME=1
ifeq ($(UNAME_S), Darwin)
FRAMEWORKS += -framework OpenGL
else
LDFLAGS += -lGL -lglfw
endif
RUNTIME_DEP += $(OPENGL_OBJ)
else
CFLAGS += -DTVM_OPENGL_RUNTIME=0
endif

ifeq ($(USE_METAL), 1)
CFLAGS += -DTVM_METAL_RUNTIME=1
LDFLAGS += -lobjc
Expand Down
4 changes: 3 additions & 1 deletion include/tvm/runtime/c_runtime_api.h
Original file line number Diff line number Diff line change
Expand Up @@ -55,9 +55,11 @@ typedef int64_t tvm_index_t;

/*! \brief Extension device types in TVM */
typedef enum {
kOpenGL = 11,

// Extension DRAM type, used for quickly test extension device
// The device api can differ depending on the xpu driver registered.
kExtDev = 12
kExtDev = 12,
// AddExtraTVMType which is not in DLPack here
} TVMDeviceExtType;

Expand Down
11 changes: 8 additions & 3 deletions include/tvm/runtime/device_api.h
Original file line number Diff line number Diff line change
Expand Up @@ -55,11 +55,16 @@ class DeviceAPI {
/*!
* \brief Allocate a data space on device.
* \param ctx The device context to perform operation.
* \param size The size of the memory
* \param nbytes The number of bytes in memory.
* \param alignment The alignment of the memory.
* \return The allocated device pointer
* \param type_hint The type of elements. Only needed by certain backends such
* as OpenGL, as nbytes & alignment are sufficient for most backends.
* \return The allocated device pointer.
*/
virtual void* AllocDataSpace(TVMContext ctx, size_t size, size_t alignment) = 0;
virtual void* AllocDataSpace(TVMContext ctx,
size_t nbytes,
size_t alignment,
TVMType type_hint) = 0;
/*!
* \brief Free a data space on device.
* \param ctx The device context to perform operation.
Expand Down
5 changes: 5 additions & 0 deletions include/tvm/schedule.h
Original file line number Diff line number Diff line change
Expand Up @@ -213,6 +213,11 @@ class Stage : public NodeRef {
* \return reference to self.
*/
Stage& double_buffer(); // NOLINT(*)
/*!
* \brief Schedule for OpenGL fragment shader.
* \return reference to self.
*/
Stage& opengl(); // NOLINT(*)
/*!
* \brief whether the stage has been scheduled.
* \return whether the stage has been scheduled.
Expand Down
2 changes: 1 addition & 1 deletion python/tvm/__init__.py
Original file line number Diff line number Diff line change
Expand Up @@ -17,7 +17,7 @@
from . import target

from . import ndarray as nd
from .ndarray import context, cpu, gpu, opencl, cl, metal, mtl, vpi, rocm, ext_dev
from .ndarray import context, cpu, gpu, opencl, cl, metal, mtl, vpi, rocm, opengl, ext_dev

from ._ffi.runtime_ctypes import TypeCode
from ._ffi.function import Function
Expand Down
2 changes: 2 additions & 0 deletions python/tvm/_ffi/runtime_ctypes.py
Original file line number Diff line number Diff line change
Expand Up @@ -97,6 +97,7 @@ class TVMContext(ctypes.Structure):
8 : 'metal',
9 : 'vpi',
10: 'rocm',
11: 'opengl',
12: 'ext_dev',
}
STR2MASK = {
Expand All @@ -111,6 +112,7 @@ class TVMContext(ctypes.Structure):
'metal': 8,
'vpi': 9,
'rocm': 10,
'opengl': 11,
'ext_dev': 12,
}
def __init__(self, device_type, device_id):
Expand Down
4 changes: 4 additions & 0 deletions python/tvm/contrib/rpc.py
Original file line number Diff line number Diff line change
Expand Up @@ -285,6 +285,10 @@ def metal(self, dev_id=0):
"""Construct remote Metal device."""
return self.context(8, dev_id)

def opengl(self, dev_id=0):
"""Construct remote OpenGL device."""
return self.context(11, dev_id)

def ext_dev(self, dev_id=0):
"""Construct remote extension device."""
return self.context(12, dev_id)
Expand Down
15 changes: 15 additions & 0 deletions python/tvm/ndarray.py
Original file line number Diff line number Diff line change
Expand Up @@ -120,6 +120,21 @@ def vpi(dev_id=0):
"""
return TVMContext(9, dev_id)

def opengl(dev_id=0):
"""Construct a OpenGL device

Parameters
----------
dev_id : int, optional
The integer device id

Returns
-------
ctx : TVMContext
The created context
"""
return TVMContext(11, dev_id)

def ext_dev(dev_id=0):
"""Construct a extension device

Expand Down
7 changes: 7 additions & 0 deletions python/tvm/schedule.py
Original file line number Diff line number Diff line change
Expand Up @@ -611,4 +611,11 @@ def double_buffer(self):
"""
_api_internal._StageDoubleBuffer(self)

def opengl(self):
"""The special OpenGL schedule

Maps each output element to a pixel.
"""
_api_internal._StageOpenGL(self)

_init_api("tvm.schedule")
4 changes: 3 additions & 1 deletion python/tvm/target.py
Original file line number Diff line number Diff line change
Expand Up @@ -67,7 +67,7 @@ class Target(object):

Parameters
----------
target_name : {"llvm", "cuda", "opencl", "metal", "rocm", "stackvm", "ext_dev"}
target_name : {"llvm", "cuda", "opencl", "metal", "rocm", "stackvm", "opengl", "ext_dev"}
The major target name.

options : list of str, optional
Expand Down Expand Up @@ -119,6 +119,8 @@ def __init__(self,
elif target_name in ("metal",):
self.keys += ("gpu",)
self.max_num_threads = 256
elif target_name in ("opengl",):
self.keys += ("opengl",)
elif target_name in ("stackvm", "ext_dev"):
# Do not now class for stacvm or ext_dev
pass
Expand Down
5 changes: 5 additions & 0 deletions src/api/api_lang.cc
Original file line number Diff line number Diff line change
Expand Up @@ -399,6 +399,11 @@ TVM_REGISTER_API("_StageDoubleBuffer")
args[0].operator Stage().double_buffer();
});

TVM_REGISTER_API("_StageOpenGL")
.set_body([](TVMArgs args, TVMRetValue *ret) {
args[0].operator Stage().opengl();
});

TVM_REGISTER_API("_ScheduleNormalize")
.set_body([](TVMArgs args, TVMRetValue* ret) {
*ret = args[0].operator Schedule()
Expand Down
35 changes: 35 additions & 0 deletions src/codegen/build_opengl.cc
Original file line number Diff line number Diff line change
@@ -0,0 +1,35 @@
/*!
* Copyright (c) 2017 by Contributors
* Build opengl modules from source.
* \file build_opengl.cc
*/
#include <tvm/base.h>
#include "./codegen_opengl.h"
#include "./build_common.h"

namespace tvm {
namespace codegen {

runtime::Module BuildOpenGL(Array<LoweredFunc> funcs) {
bool output_ssa = false;
CodeGenOpenGL cg;
cg.Init(output_ssa);
for (LoweredFunc f : funcs) {
cg.AddFunction(f);
}
auto shaders = cg.Finish();
#if TVM_OPENGL_RUNTIME
return OpenGLModuleCreate(shaders, "gl", ExtractFuncInfo(funcs));
#else
LOG(WARNING) << "OpenGL runtime not enabled, return a source module...";
auto data = ToJSON(shaders);
return DeviceSourceModuleCreate(data, "gl", ExtractFuncInfo(funcs), "opengl");
#endif // TVM_OPENGL_RUNTIME
}

TVM_REGISTER_API("codegen.build_opengl")
.set_body([](TVMArgs args, TVMRetValue* rv) {
*rv = BuildOpenGL(args[0]);
});
} // namespace codegen
} // namespace tvm
2 changes: 1 addition & 1 deletion src/codegen/codegen_c.h
Original file line number Diff line number Diff line change
Expand Up @@ -150,7 +150,7 @@ class CodeGenC :
std::string GetStructRef(
Type t, const Expr& buffer, const Expr& index, int kind);
// print reference to a buffer as type t in index.
std::string GetBufferRef(
virtual std::string GetBufferRef(
Type t, const Variable* buffer, Expr index);
/*!
* \brief If buffer is allocated as type t.
Expand Down
Loading