Basic WebGL Backend #672

phisiart · 2017-11-26T20:33:15Z

TLDR

Currently the following demo program runs and gets the correct result.

from __future__ import absolute_import, print_function

import tvm
import numpy as np

n = tvm.var("n")
A = tvm.placeholder((n,), name='A')
B = tvm.placeholder((n,), name='B')
C = tvm.compute(A.shape, lambda i: A[i] + B[i], name="C")

s = tvm.create_schedule(C.op)
s[C].opengl()

fadd_gl = tvm.build(s, [A, B, C], "opengl", name="myadd")
print("------opengl code------")
print(fadd_gl.imported_modules[0].get_source(fmt="gl"))

ctx = tvm.opengl(0)
n = 10
a = tvm.nd.array(np.random.uniform(size=(n)).astype(A.dtype), ctx)
b = tvm.nd.array(np.random.uniform(size=(n)).astype(B.dtype), ctx)
c = tvm.nd.array(np.zeros((n), dtype=C.dtype), ctx)
fadd_gl(a, b, c)

np.testing.assert_allclose(c.asnumpy(), a.asnumpy() + b.asnumpy())

The corresponding fragment shader is

#version 330 core
uniform sampler2D A;
uniform sampler2D B;
out float C;
void main() {
  ivec2 threadIdx = ivec2(gl_FragCoord.xy);
  C = (texelFetch(A, ivec2(threadIdx.x, 0), 0).r + texelFetch(B, ivec2(threadIdx.x, 0), 0).r);
}

Current Status

OpenGL Tensor Storage

We store tensors in OpenGL textures.
No matter what dimensions a tensor has, we always store it in a 2D texture with height=1.
The reason we are not using 1D textures is that texelFetch in GLSL only supports 2D textures.
We support uint{8,16,32}, int{8,16,32} and float32 textures.

OpenGL Schedule

We added an opengl schedule, which basically fuses all dimensions into one and binds that single dimension to threadIdx.x.

Codegen

We haven't changed lowering at all. When the IR says Store(buffer, index) and the buffer happens to be the output texture, we check that index must be threadIdx.x, and emit code to output a pixel.
The codegen part has only been started. We will go through all the AST nodes as our next step.

Task Items

~~Remove the dependency on glfw and glad.~~ Emscripten supports glfw. Glad removed.
~~Investigate emscripten.~~ Runtime runnable.
~~Support other types than float.~~ Done.
~~Use full RGBA channels.~~ Not critical now.
Don't let height always be 1. The reason is that OpenGL textures have stupid size limitations. For example, you can have a 2500x2500 2D texture but not a 6250000x1 2D texture. One possible way is just let height be the maximum supported value. Not critical now.
Go through the AST to properly do codegen.
~~Don't fuse reduction.~~ Gemm runnable.
~~Fix argument name generation.~~ Done.

tqchen · 2017-11-27T19:40:41Z

.gitmodules

@@ -7,3 +7,7 @@
 [submodule "dlpack"]
 	path = dlpack
 	url = https://github.com/dmlc/dlpack
+[submodule "glad"]


What is glad used for, is it possible to only use OpenGL standard and not rely on new libraries?

We actually rely on 2 libraries.

glad: This is the library to help you use the correct version of OpenGL.
For example, suppose your OS supports OpenGL 3.3, then the OS is able to give you any OpenGL version below 3.3 on your request. However, if you just #include <GL/gl.h> or whatever header your OS provides, the function prototypes are for OpenGL 1.3 or something, so you can't use OpenGL 3.3 even your OS supports it. The OS provides a way of querying the function pointers of OpenGL 3.3 APIs. That's what glad does. It is possible that we don't rely on glad, but with a huge cost - we must write our own code to get all the function pointers.

glfw: This is the library to help you create an OpenGL context in a cross-platform way.
You can't use any OpenGL API before creating a context. GLFW wraps the context creation code of different OS'es to provide a unified API. It is possible that we don't rely on glfw, but with a huge cost - we must write context initialization code for different platforms.

I can remove the dependencies on those libraries, but that can only happen after Dec 12 or something. We are focusing on producing a final-report-able version currently.

Sure, I think these are fine as long as we fix this at the final merge.

AFAIK GLFW itself wraps up all the functions. Maybe GLAD is not needed since GLFW is already there. If my memory went wrong please inform me.

tqchen · 2017-11-27T19:41:31Z

include/tvm/runtime/c_runtime_api.h

  // AddExtraTVMType which is not in DLPack here
+  kOpenGL = 11,


put kOpenGL before kExtDev

tqchen · 2017-11-28T17:55:44Z

src/schedule/schedule_lang.cc

+  CHECK(!is_scheduled()) << "Must be a fresh schedule";
+  StageNode *self = operator->();
+
+  auto all_iter_vars = self->all_iter_vars; // curr version of all_iter_vars


be careful that this can include reduction variables, which should not be part of the fusing stage

Will go back to this later. Added a TODO for now.

tqchen · 2017-11-28T17:58:09Z

src/runtime/opengl/opengl_device_api.cc

+    assert(false);
+  }
+
+  LOG_INFO.stream() << "GLFW says OpenGL version: "


Can directly use LOG(INFO)

tqchen · 2017-11-28T17:58:58Z

src/codegen/codegen_opengl.cc

@@ -0,0 +1,138 @@
+/*!
+ *  Copyright (c) 2017 by Contributors
+ * \file codegen_opengl.cc


Add a specific comment that we are targeting subset of OpenGL that is working for WebGL2,(no compute shaders)

tqchen · 2017-11-28T17:59:39Z

python/tvm/target.py

@@ -114,6 +114,9 @@ def __init__(self,
        elif target_name in ("metal",):
            self.keys += ("gpu",)
            self.max_num_threads = 256
+        elif target_name in ("opengl"):


use opengl as key, as opengl schedule need to be different from gpu schedules

PENGUINLIONG · 2017-12-02T02:42:25Z

python/tvm/target.py

@@ -114,6 +114,9 @@ def __init__(self,
        elif target_name in ("metal",):
            self.keys += ("gpu",)
            self.max_num_threads = 256
+        elif target_name in ("opengl"):


Shouldn't it be ("opengl",)? It ought to be searching for target name in a tuple.

phisiart · 2017-12-03T03:23:06Z

I'm currently stuck on other final projects for courses, so I won't be able to work on this PR until mid December. Then I should have roughly a month after the end of this semester.

phisiart · 2017-12-15T23:16:40Z

Next step: Remove dependency on glfw and glad. I will focus on Ubuntu, since it should be the tricky one. Windows and Mac should have official APIs to create OpenGL contexts.

I will see what Ubuntu Desktop has by default and only depend on that.
On Ubuntu Server, OpenGL might not exist by default (because it doesn't have GUI). If that's the case I will see what minimum packages are needed to have OpenGL, and only depend on them.

PENGUINLIONG · 2017-12-16T01:54:40Z

Use wgl* functions you can create a context on Windows.

On Linux, if MESA is installed, OpenGL soft-rendering should be available as well as accelerated graphics, if there is a GPU. Compute Shader should be supported as it's a part of the OpenGL 4.3 spec.

FYI, without GLFW it will be your responsibility to query if a GL function is available. Using GL_ABR_compute_shader, the minimum GL version requirement should be (Core) 4.3. On initialization, GLFW stores function pointers to GL APIs so that the query results can be reused.

tqchen · 2017-12-16T02:19:19Z

One of the the main goal here is to use WebGL, this means

so the standard version should be less than OpenGL2.
we cannot use compute shader

Please try to emscripten the GL runtime to see if we can successfully build it

PENGUINLIONG · 2017-12-16T04:44:56Z

I see. But it would be nice if we can utilize the functionalities provided by OpenGL of higher versions - Compute Shader. We have Compute Shader support in OpenGL (Core) 4.3 and in OpenGL ES 3.1 but not yet in WebGL. In the current latest commit, you can see there is a dummy vertex shader that doesn't need to but has to be executed.

Or a strategy can be devised so that both shaders can do some work.

tqchen · 2017-12-16T05:17:59Z

For most other devices, opencl is also available, which is more desirable than OpenGL itself

phisiart · 2017-12-19T00:53:44Z

After some experiments, I found that the following is the minimum requirements for find_package(OpenGL REQUIRED) in cmake to succeed.

Dockerfile:

FROM ubuntu

RUN apt-get update --fix-missing

RUN apt-get install -y --no-install-recommends \
  make cmake g++ libgl1-mesa-dev

CMake is looking for libGL.so, and on Ubuntu there are 3 packages that provide this (see here):

mesa: This is what comes by default when you install Ubuntu desktop;
nvidia: The NVIDIA driver;
fglrx: The AMD driver.

Therefore, we should depend on libgl1-mesa-dev.

phisiart · 2017-12-19T01:10:14Z

Okay, I just tested and found out that on a fresh installation of Ubuntu Desktop you also need the package libgl1-mesa-dev, for find_package(OpenGL REQUIRED) in cmake to succeed (because Ubuntu doesn't come with the 'dev' version of the package).

PENGUINLIONG · 2017-12-20T02:31:38Z

Since the goal for the OpenGL port has shifted to NN deployment on web apps.. Does it still necessarily need mesa?

Emscripten is compiling LLVM intermediate to JavaScript. It should require no linking to native libraries. Maybe we can just extern "C"-declare those APIs we need, so that the feature can be implemented with no dependency added.

See this for GLES2 APIs.

phisiart · 2017-12-20T03:14:45Z

@PENGUINLIONG I think we want to support both running natively and in the browser.

A funny thing is that emscripten has direct support for glfw. If we remove the dependency on glfw then we would be calling native glx functions. I'm not sure whether emscripten has good support for those. I will look into it and report my findings.

PENGUINLIONG · 2017-12-20T04:06:11Z

@phisiart It seems Emscripten simply map the calls. See this. Then it's possible that we simply declare them and they will work.

However, it seems not possible we can create context without dependencies.

tqchen · 2017-12-23T02:53:43Z

I think the first milestone is to confirm emscripten with simple gl runtime works. You can likely utilize the RPC module here https://github.com/dmlc/tvm/tree/master/web for testing

phisiart · 2018-01-01T03:13:37Z

I tried RPC. I have successfully crashed in the middle of OpenGL initialization, after glfwInit(). It seems like incompatibility with glad. Since I'm going to remove the dependency on glad anyway, let me do it and then continue with RPC.

phisiart · 2018-01-01T23:59:17Z

Update: Now OpenGL initialization succeeds.

phisiart · 2018-01-02T00:02:58Z

Update: Now rendering succeeds.

tqchen · 2018-01-02T01:09:30Z

here are some trackable milestones changes that I think could be useful. It would be very helpful to create a series of test functions under tests/web/webgl that relies on web proxy rpc and can be run manually.

Runtime array copy pass, something like https://github.com/dmlc/tvm/blob/master/tests/python/unittest/test_runtime_ndarray.py
Simple ewise pass via RPC, something like https://github.com/dmlc/tvm/blob/master/tests/python/unittest/test_runtime_rpc.py#L134
Pass the basic gemm
Enable gl in topi, so we can pass end to end resnet in nnvm compiler

tqchen · 2018-01-06T04:40:38Z

Makefile

@@ -30,10 +30,10 @@ CFLAGS = -std=c++11 -Wall -O2 $(INCLUDE_FLAGS) -fPIC
 FRAMEWORKS =
 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\
+	-O0 -s RESERVED_FUNCTION_POINTERS=2 -s MAIN_MODULE=1 -s NO_EXIT_RUNTIME=1\


is -O0 necessary? or can we use -Oz?

Oh this is just for debugging purposes. Will be changed back when I cleanup.

Changed back.

tqchen · 2018-01-06T04:41:28Z

include/tvm/runtime/device_api.h

   * \param alignment The alignment of the memory.
   * \return The allocated device pointer
   */
-  virtual void* AllocDataSpace(TVMContext ctx, size_t size, size_t alignment) = 0;
+  virtual void* AllocDataSpace(TVMContext ctx, TVMType type, size_t nbytes,


let us put type->type_hint and put it as last parameter. Use comment to say type_hint is only needed by a few backend such as GL

tqchen · 2018-01-06T04:43:05Z

src/runtime/cpu_device_api.cc

@@ -20,13 +20,14 @@ class CPUDeviceAPI final : public DeviceAPI {
      *rv = 1;
    }
  }
-  void* AllocDataSpace(TVMContext ctx, size_t size, size_t alignment) final {
+  void* AllocDataSpace(TVMContext ctx, TVMType type, size_t nbytes,


if we have multiple line breaks, it might be more natural to break each argument into one line

tqchen · 2018-01-06T04:43:54Z

src/runtime/rpc/rpc_session.h

@@ -48,6 +48,7 @@ enum class RPCCode : int {
  kModuleFree,
  kModuleGetFunc,
  kModuleGetSource,
+  kTestRemoteOpenGL,


this can be safely removed when we upstream?

Yes, this is just for my debugging purposes. Will remove when I cleanup.

tqchen · 2018-01-06T04:46:37Z

include/tvm/schedule.h

@@ -213,6 +213,11 @@ class Stage : public NodeRef {
   * \return reference to self.
   */
  Stage& double_buffer();   // NOLINT(*)
+  /*!
+   * \brief Schedule for GpenGL fragment shader.


tqchen · 2018-01-06T04:49:30Z

for the case when we need 2D texture, we might be able to use the following workaround.

Always allocate memory as i = 2^k * y + x, where x and y are two dimensional axis. k is the folding ideal packing length over the x dimension, say 10

so we can easily use bit operation to get y and x from global index i, and reconstruct vice versa

The Opengl runtime will need this information for textures.

- fix opengl func param retrieval; - can save opengl module locally; - working on loading opengl module in browser.

Known issue: cannot retrieve integer texture data in webgl.

from __future__ import absolute_import, print_function import tvm import numpy as np n = tvm.var("n") m = tvm.var("m") A = tvm.placeholder((n, m), name='A') k = tvm.reduce_axis((0, m), "k") B = tvm.compute((n,), lambda i: tvm.sum(A[i, k], axis=k), name="B") s = tvm.create_schedule(B.op) s[B].opengl() fadd_gl = tvm.build(s, [A, B], "opengl", name="myadd") print("------opengl code------") print(fadd_gl.imported_modules[0].get_source(fmt="gl")) ctx = tvm.opengl(0) n = 10 m = 10 a = tvm.nd.array(np.random.uniform(size=(n, m)).astype(A.dtype), ctx) b = tvm.nd.array(np.random.uniform(size=(n,)).astype(B.dtype), ctx) fadd_gl(a, b) np.testing.assert_allclose(b.asnumpy(), np.sum(a.asnumpy(), axis=1))

phisiart · 2018-01-20T20:13:01Z

Review comments are addressed.

tqchen · 2018-01-20T21:51:19Z

Thanks, this concludes the basic support of GL, and this PR is merged. Let us start new PRs to make further improvements

Basic WebGL Backend

tqchen reviewed Nov 27, 2017

View reviewed changes

tqchen reviewed Nov 28, 2017

View reviewed changes

tqchen requested changes Nov 28, 2017

View reviewed changes

PENGUINLIONG reviewed Dec 2, 2017

View reviewed changes

phisiart force-pushed the opengl branch from e5c159d to 107bdbd Compare December 15, 2017 22:36

phisiart force-pushed the opengl branch from e588170 to 57bda69 Compare December 27, 2017 03:36

phisiart force-pushed the opengl branch 3 times, most recently from 95ccc12 to f3d1724 Compare January 2, 2018 04:04

tqchen requested changes Jan 6, 2018

View reviewed changes

phisiart added 19 commits January 20, 2018 14:50

Trying to add OpenGL runtime to emscripten'ed web runtime.

4d55284

remove glad, add temporary rpc opengl test

ad6c8dc

Add 'type' parameter to AllocDataSpace.

f17ac76

The Opengl runtime will need this information for textures.

Improve OpenGL texture. Now test_runtime_ndarray.py passes.

4ee83ce

- address review comments;

9df0c34

- fix opengl func param retrieval; - can save opengl module locally; - working on loading opengl module in browser.

add remote.download("myadd.tvm_meta.json")

a0662a3

Now tests/webgl/test_remote_save_load.py succeeds.

c7c30f0

Known issue: cannot retrieve integer texture data in webgl.

Savepoint that works.

c0b04a8

Remove temporary RPC test.

da21d5c

Correctly handle OpenGL argument.

3d669fd

Change emcc optmization flag back to -Oz.

3c90dde

Generate tvm_get_texel function in GLSL for cleaner code.

96145f9

Add tests/webgl/test_local_gemm.py

d289979

Directly return if threadIdx.x < thread_extent.

20087e3

Cleanup OpenGLArgKind and OpenGLShader.

8dbd632

Enable local OpenGL tests in cpu & gpu Jenkins.

2acb00e

Cleanup.

d30d6a5

Address review comments.

8b98b59

phisiart force-pushed the opengl branch from 686ad2a to 8b98b59 Compare January 20, 2018 19:50

tqchen approved these changes Jan 20, 2018

View reviewed changes

tqchen changed the title ~~[WIP] WebGL Backend~~ Basic WebGL Backend Jan 20, 2018

tqchen merged commit 7009496 into apache:master Jan 20, 2018

tqchen mentioned this pull request Jan 20, 2018

WebGL Followup #799

Closed

5 tasks

tqchen pushed a commit to tqchen/tvm that referenced this pull request Jul 6, 2018

[WIP] WebGL Backend (apache#672)

589831d

Basic WebGL Backend

sergei-mironov pushed a commit to sergei-mironov/tvm that referenced this pull request Aug 8, 2018

[WIP] WebGL Backend (apache#672)

9533d3d

Basic WebGL Backend

archibate mentioned this pull request Feb 17, 2020

[Backend] New OpenGL Compute Shader Backend taichi-dev/taichi#492

Closed

Basic WebGL Backend #672

Basic WebGL Backend #672

Conversation

phisiart commented Nov 26, 2017 • edited Loading

TLDR

Current Status

OpenGL Tensor Storage

OpenGL Schedule

Codegen

Task Items

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

phisiart commented Dec 3, 2017

phisiart commented Dec 15, 2017

PENGUINLIONG commented Dec 16, 2017

tqchen commented Dec 16, 2017 • edited Loading

PENGUINLIONG commented Dec 16, 2017

tqchen commented Dec 16, 2017

phisiart commented Dec 19, 2017

phisiart commented Dec 19, 2017

PENGUINLIONG commented Dec 20, 2017

phisiart commented Dec 20, 2017

PENGUINLIONG commented Dec 20, 2017

tqchen commented Dec 23, 2017

phisiart commented Jan 1, 2018

phisiart commented Jan 1, 2018

phisiart commented Jan 2, 2018

tqchen commented Jan 2, 2018 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

tqchen commented Jan 6, 2018

phisiart commented Jan 20, 2018

tqchen commented Jan 20, 2018

phisiart commented Nov 26, 2017 •

edited

Loading

tqchen commented Dec 16, 2017 •

edited

Loading

tqchen commented Jan 2, 2018 •

edited

Loading