Skip to content

Commit 7fe8d8b

Browse files
austinvhuangaustinvhuang
committed
add emscripten hello world implementation (working), update matrix show() to use snprintf instead of sprintf
1 parent dbaeef0 commit 7fe8d8b

File tree

17 files changed

+220
-21
lines changed

17 files changed

+220
-21
lines changed

‎experimental/wasm/emscripten/CMakeLists.txt

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
cmake_minimum_required(VERSION 3.13)
2+
project(run)
3+
set(CMAKE_CXX_STANDARD 20)
4+
5+
add_executable(run "run.cpp")
6+
7+
set_target_properties(run PROPERTIES SUFFIX ".html")
8+
set(SHELL_FILE_PATH "../custom_shell.html")
9+
10+
# target_link_options(run PRIVATE "-sUSE_WEBGPU=1" "-sUSE_GLFW=3")
11+
12+
# asyncify for emscripten_sleep()
13+
# see https://emscripten.org/docs/api_reference/emscripten.h.html#pseudo-synchronous-functions
14+
target_link_options(run PRIVATE "-sUSE_WEBGPU=1" "-sASYNCIFY=1" "--shell-file=${SHELL_FILE_PATH}")
15+
16+
# em++ compiler arguments
17+
target_compile_options(run PRIVATE "-sUSE_WEBGPU=1" "-sASYNCIFY=1" "--shell-file=${SHELL_FILE_PATH}")
18+
19+
20+
target_include_directories(run PRIVATE "../../../")
21+
target_include_directories(run PRIVATE "../../../numeric_types")
22+
target_include_directories(run PRIVATE "../../../utils")

‎experimental/wasm/emscripten/Makefile

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,19 @@
1+
2+
FLAGS=--target=wasm32 -nostdlib -Wl,--no-entry -Wl,--export-all -Wl,--import-memory -Wl,--allow-undefined -fexceptions -std=c++17 -O3
3+
4+
5+
# Make sure you have:
6+
# 1) emsdk installed
7+
# 2) `source emsdk_env.sh` in the emsdk directory
8+
build/run.wasm:
9+
em++ -sUSE_WEBGPU=1 -sUSE_GLFW=3 run.cpp -o build/run.wasm
10+
11+
cmake:
12+
emcmake cmake -B build && cmake --build build -j4
13+
npx http-server
14+
15+
browser:
16+
open http://127.0.0.1:8000/build/run.html
17+
18+
clean:
19+
rm -rf build/*

‎experimental/wasm/build/.gitkeep renamed to ‎experimental/wasm/emscripten/build/.gitkeep

File renamed without changes.

‎experimental/wasm/emscripten/custom_shell.html

Lines changed: 53 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,53 @@
1+
<!DOCTYPE html>
2+
<html>
3+
<head>
4+
<meta charset="UTF-8">
5+
<title>gpu.cpp</title>
6+
<!-- Include xterm.js CSS -->
7+
<link rel="stylesheet" href="https://cdn.jsdelivr.net/npm/xterm/css/xterm.css" />
8+
<style>
9+
body, html {
10+
margin: 0;
11+
padding: 0;
12+
height: 100%;
13+
width: 100%;
14+
}
15+
#terminal {
16+
height: 100%;
17+
width: 100%;
18+
}
19+
</style>
20+
</head>
21+
<body>
22+
<!-- Terminal container -->
23+
<div id="terminal"></div>
24+
<!-- Include xterm.js and xterm-addon-fit libraries -->
25+
<script src="https://cdn.jsdelivr.net/npm/xterm/lib/xterm.js"></script>
26+
<script src="https://cdn.jsdelivr.net/npm/xterm-addon-fit/lib/xterm-addon-fit.js"></script>
27+
<script type="text/javascript">
28+
// Initialize xterm.js
29+
const terminal = new Terminal();
30+
const fitAddon = new FitAddon.FitAddon();
31+
terminal.loadAddon(fitAddon);
32+
terminal.open(document.getElementById('terminal'));
33+
fitAddon.fit();
34+
35+
// This function captures stdout and displays it in the xterm.js terminal
36+
var Module = {
37+
preRun: [],
38+
postRun: [],
39+
print: function(text) {
40+
if (arguments.length > 1) text = Array.prototype.slice.call(arguments).join(' ');
41+
console.log(text); // Log to the console
42+
terminal.writeln(text); // Write to the terminal
43+
},
44+
printErr: function(text) {
45+
if (arguments.length > 1) text = Array.prototype.slice.call(arguments).join(' ');
46+
console.error(text); // Log to the console as an error
47+
terminal.writeln(text); // Write to the terminal as well
48+
}
49+
};
50+
</script>
51+
{{{ SCRIPT }}}
52+
</body>
53+
</html>

‎experimental/wasm/emscripten/run.cpp

Lines changed: 62 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,62 @@
1+
#include <array>
2+
#include <cstdio>
3+
#include <future>
4+
#include <memory>
5+
#include "gpu.h"
6+
7+
// #include <webgpu/webgpu.h>
8+
#include "emscripten/emscripten.h"
9+
10+
using namespace gpu; // createContext, createTensor, createKernel,
11+
// createShader, dispatchKernel, wait, toCPU
12+
// Tensor, Kernel, Context, Shape, kf32
13+
14+
static const char *kGelu = R"(
15+
const GELU_SCALING_FACTOR: f32 = 0.7978845608028654; // sqrt(2.0 / PI)
16+
@group(0) @binding(0) var<storage, read_write> inp: array<{{precision}}>;
17+
@group(0) @binding(1) var<storage, read_write> out: array<{{precision}}>;
18+
@group(0) @binding(1) var<storage, read_write> dummy: array<{{precision}}>;
19+
@compute @workgroup_size({{workgroupSize}})
20+
fn main(
21+
@builtin(global_invocation_id) GlobalInvocationID: vec3<u32>) {
22+
let i: u32 = GlobalInvocationID.x;
23+
if (i < arrayLength(&inp)) {
24+
let x: f32 = inp[i];
25+
out[i] = select(0.5 * x * (1.0 + tanh(GELU_SCALING_FACTOR
26+
* (x + .044715 * x * x * x))), x, x > 10.0);
27+
}
28+
}
29+
)";
30+
31+
int main(int argc, char **argv) {
32+
printf("\033[2J\033[1;1H");
33+
printf("\nHello gpu.cpp!\n");
34+
printf("--------------\n\n");
35+
36+
// const WGPUInstanceDescriptor descriptor = { };
37+
// std::unique_ptr<WGPUInstanceDescriptor> descriptor = std::make_unique<WGPUInstanceDescriptor>();
38+
39+
// WGPUInstance instance = wgpuCreateInstance({});
40+
Context ctx = createContext({});
41+
static constexpr size_t N = 5000;
42+
std::array<float, N> inputArr, outputArr;
43+
for (int i = 0; i < N; ++i) {
44+
inputArr[i] = static_cast<float>(i) / 10.0; // dummy input data
45+
}
46+
Tensor input = createTensor(ctx, Shape{N}, kf32, inputArr.data());
47+
Tensor output = createTensor(ctx, Shape{N}, kf32);
48+
std::promise<void> promise;
49+
std::future<void> future = promise.get_future();
50+
Kernel op = createKernel(ctx, {kGelu, 256, kf32},
51+
Bindings{input, output},
52+
{cdiv(N, 256), 1, 1});
53+
dispatchKernel(ctx, op, promise);
54+
wait(ctx, future);
55+
toCPU(ctx, output, outputArr.data(), sizeof(outputArr));
56+
for (int i = 0; i < 12; ++i) {
57+
printf(" gelu(%.2f) = %.2f\n", inputArr[i], outputArr[i]);
58+
}
59+
printf(" ...\n\n");
60+
printf("Computed %zu values of GELU(x)\n\n", N);
61+
return 0;
62+
}

‎experimental/wasm/Makefile renamed to ‎experimental/wasm/pure_wasm/Makefile

File renamed without changes.

‎experimental/wasm/pure_wasm/README.md

Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,17 @@
1+
Test of pure wasm control of webgpu w/o emscripten runtime.
2+
3+
First, build the wasm binary:
4+
5+
```
6+
make build/run.wasm
7+
```
8+
9+
Run the static server:
10+
```
11+
make server
12+
```
13+
14+
Open browser to localhost:8000.
15+
16+
You should see in the developer console `WebAssembly module loaded ... Hello
17+
World!` printed from the wasm binary.

‎experimental/wasm/pure_wasm/build/.gitkeep

Whitespace-only changes.

‎experimental/wasm/favicon.ico renamed to ‎experimental/wasm/pure_wasm/favicon.ico

File renamed without changes.

‎experimental/wasm/gpu.js renamed to ‎experimental/wasm/pure_wasm/gpu.js

File renamed without changes.

0 commit comments

Comments
 (0)