Skip to content

Commit bd5871a

Browse files
austinvhuangaustinvhuang
committed
refine state handling in fasthtml experiment, begin work on shadertui+fasthtml (WIP)
1 parent d1bb95e commit bd5871a

5 files changed

Lines changed: 186 additions & 60 deletions

File tree

‎experimental/fasthtml/Makefile‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -11,10 +11,6 @@ NO_MODULARIZE_FLAGS=-s EXPORTED_FUNCTIONS=['_executeKernel'] -s EXPORTED_RUNTIME
1111

1212
default: server
1313

14-
# build/run.html: check-emsdk run.cpp
15-
# em++ run.cpp -o build/run.html \
16-
# $(COMMON_FLAGS) $(JS_FLAGS) --shell-file ./custom_shell.html
17-
1814
build/run.js: check-emsdk run.cpp
1915
em++ run.cpp -o build/run.js \
2016
$(COMMON_FLAGS) $(JS_FLAGS) $(MODULARIZE_FLAGS)
@@ -23,6 +19,10 @@ build/run.wasm: check-emsdk run.cpp
2319
em++ run.cpp -o build/run.wasm \
2420
$(COMMON_FLAGS) $(WASM_FLAGS)
2521

22+
watch:
23+
@echo "Watching for changes..."
24+
ls run.cpp | entr -c make build/run.js
25+
2626
server: build/run.js
2727
python3 run.py
2828

‎experimental/fasthtml/components/code_editor.py‎

Lines changed: 8 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -2,46 +2,33 @@
22

33
from fasthtml.common import *
44
from .toolbar import Toolbar
5+
import json
56

6-
editor_script = Script("""
7+
def editor_script(initial_content: str) -> Script:
8+
return Script("""
79
let editor;
810
let completionTippy;
911
let currentCompletion = '';
1012
11-
1213
function initEditor() {
1314
editor = ace.edit("editor");
1415
editor.setTheme("ace/theme/monokai");
1516
editor.session.setMode("ace/mode/javascript");
1617
editor.setOptions({
1718
fontSize: "14px",
1819
showPrintMargin: false,
19-
// disable showing errors in gutter Ace WGSL parser is out of date
20+
// disable showing errors in gutter, Ace's WGSL parser is out of date
2021
showGutter: false,
2122
highlightActiveLine: true,
2223
wrap: true,
2324
});
2425
editor.setKeyboardHandler("ace/keyboard/vim");
2526
26-
editor.setValue(
27-
"// Start editing here to see the results.\\n\// Warning: You are in vim mode.\\n\
28-
\\n\
29-
@group(0) @binding(0) var<storage, read_write> input: array<f32>;\\n\
30-
@group(0) @binding(1) var<storage, read_write> output : array<f32>;\\n\
31-
@compute @workgroup_size(256)\\n\
32-
fn main(\\n\
33-
@builtin(local_invocation_id) LocalInvocationID: vec3<u32>) {\\n\
34-
let local_idx = LocalInvocationID.x;\\n\
35-
if (local_idx < arrayLength(&input)) {\\n\
36-
output[local_idx] = input[local_idx] + 1;\\n\
37-
}\\n\
38-
}\\n\
39-
");
40-
27+
editor.setValue(""" + json.dumps(initial_content) + ");" +
28+
"""
4129
window.addEventListener('resize', function() {
4230
editor.resize();
4331
});
44-
4532
document.getElementById('language').addEventListener('change', function(e) {
4633
let mode = "ace/mode/" + e.target.value;
4734
editor.session.setMode(mode);
@@ -149,7 +136,7 @@
149136
document.addEventListener('DOMContentLoaded', initEditor);
150137
""")
151138

152-
def CodeEditor():
139+
def CodeEditor(initial_content: str):
153140
return (
154141
Div(
155142
Toolbar(),
@@ -167,5 +154,5 @@ def CodeEditor():
167154
# cls="flex flex-col h-screen w-full", style="height: 100vh; overflow: hidden;"
168155
style="height: 100vh; overflow: hidden;"
169156
),
170-
editor_script
157+
editor_script(initial_content)
171158
)

‎experimental/fasthtml/run.cpp‎

Lines changed: 11 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -49,17 +49,23 @@ void executeKernel(const char *kernelCode) {
4949
toCPU(ctx, output, outputArr.data(), sizeof(outputArr));
5050

5151
char buffer[1024];
52-
for (int i = 0; i < 12; ++i) {
53-
snprintf(buffer, sizeof(buffer), " kernel(%.1f) = %.4f", inputArr[i],
54-
outputArr[i]);
52+
for (int i = 0; i < 10; ++i) {
53+
snprintf(buffer, sizeof(buffer), " [%d] kernel(%.1f) = %.4f", i,
54+
inputArr[i], outputArr[i]);
5555
js_print(buffer);
5656
}
5757
snprintf(buffer, sizeof(buffer), " ...");
5858
js_print(buffer);
59+
for (int i = N - 10; i < N; ++i) {
60+
snprintf(buffer, sizeof(buffer), " [%d] kernel(%.1f) = %.4f", i,
61+
inputArr[i], outputArr[i]);
62+
js_print(buffer);
63+
}
5964
snprintf(buffer, sizeof(buffer), "Computed %zu values", N);
6065
js_print(buffer);
61-
}
62-
}
66+
} // executeKernel
67+
68+
} // extern "C"
6369

6470
#ifndef STANDALONE_WASM
6571
#include "emscripten/bind.h"

‎experimental/fasthtml/run.py‎

Lines changed: 64 additions & 30 deletions
Original file line numberDiff line numberDiff line change
@@ -7,34 +7,42 @@
77
TARGET = os.getenv("TARGET", "debug")
88

99
ace_editor = Script(src="https://cdnjs.cloudflare.com/ajax/libs/ace/1.4.12/ace.js")
10-
gpucpp_runtime = Script(src="/build/run.js")
11-
gpucpp_wasm = Script(src="/build/run.wasm")
10+
gpucpp_runtime = Script(src="/build/run.js")
11+
gpucpp_wasm = Script(src="/build/run.wasm")
1212
tippy_css = Link(rel="stylesheet", href="https://unpkg.com/tippy.js@6/dist/tippy.css")
1313
tippy_js = Script(src="https://unpkg.com/@popperjs/core@2")
1414
tippy_js2 = Script(src="https://unpkg.com/tippy.js@6")
15-
xterm_css = Link(rel="stylesheet", href="https://cdn.jsdelivr.net/npm/xterm/css/xterm.css")
15+
xterm_css = Link(
16+
rel="stylesheet", href="https://cdn.jsdelivr.net/npm/xterm/css/xterm.css"
17+
)
1618
xterm_js = Script(src="https://cdn.jsdelivr.net/npm/xterm/lib/xterm.js")
17-
xterm_fit_js = Script(src="https://cdn.jsdelivr.net/npm/xterm-addon-fit/lib/xterm-addon-fit.js")
19+
xterm_fit_js = Script(
20+
src="https://cdn.jsdelivr.net/npm/xterm-addon-fit/lib/xterm-addon-fit.js"
21+
)
1822

19-
global_style = Style("""
23+
global_style = Style(
24+
"""
2025
#editor {
2126
height: 50vh;
2227
width: 50vw;
2328
}
24-
""")
29+
"""
30+
)
2531

26-
terminal_init = Script("""
27-
console.log("Terminal initialized");
32+
terminal_init = \
33+
"""
2834
const terminal = new Terminal();
2935
const fitAddon = new FitAddon.FitAddon();
3036
terminal.loadAddon(fitAddon);
3137
// terminal.open(document.getElementById('output'));
3238
window.terminal = terminal;
33-
fitAddon.fit();
39+
// fitAddon.fit();
3440
console.log("Terminal initialized");
35-
""");
41+
"""
3642

37-
print_script = Script("""
43+
print_script = (
44+
Script(
45+
"""
3846
window.customPrint = function(text) {
3947
console.log(text);
4048
if (window.terminal) {
@@ -48,24 +56,43 @@
4856
Module.printErr = window.customPrint;
4957
window.Module = Module;
5058
console.log("Module ready");
59+
// window.Module.executeKernel(editor.getValue());
5160
});
52-
"""),
53-
54-
bind_terminal = Script("window.terminal.open(document.getElementById('output'));")
61+
"""
62+
),
63+
)
5564

56-
# TODO(avh): Global state handling of terminal binding, module creation, etc.
65+
bind_terminal = """
66+
window.terminal.open(document.getElementById('output'));
67+
fitAddon.fit();
68+
"""
69+
70+
gelu_kernel = """// Start editing here to see the results.
71+
// Warning: You are in vim mode.
72+
@group(0) @binding(0) var<storage, read_write> input : array<f32>;
73+
@group(0) @binding(1) var<storage, read_write> output : array<f32>;
74+
@compute @workgroup_size(256)
75+
fn main(
76+
@builtin(global_invocation_id) GlobalInvocationID: vec3<u32>) {
77+
let i: u32 = GlobalInvocationID.x;
78+
if (i < arrayLength(&input)) {
79+
output[i] = input[i] + 1;
80+
}
81+
}
82+
"""
83+
84+
# TODO(avh) : Global state handling of terminal binding, module creation, etc.
5785
# could be improved
5886

5987
HDRS = (
60-
picolink,
88+
picolink,
6189
ace_editor,
6290
xterm_css,
6391
xterm_js,
6492
xterm_fit_js,
65-
terminal_init,
93+
Script(terminal_init),
6694
gpucpp_runtime,
6795
print_script,
68-
bind_terminal,
6996
global_style,
7097
tippy_css,
7198
tippy_js,
@@ -78,7 +105,7 @@
78105
image="",
79106
url="https://gpucpp.answer.ai",
80107
),
81-
)
108+
)
82109

83110
if TARGET == "release":
84111
app = FastHTML(hdrs=HDRS)
@@ -87,29 +114,36 @@
87114

88115
rt = app.route
89116

117+
90118
@app.get("/build/run.js")
91119
async def serve_wasm(fname: str, ext: str):
92120
return FileResponse(f"build/run.js")
93121

122+
94123
@app.get("/build/run.wasm")
95124
async def serve_wasm(fname: str, ext: str):
96125
return FileResponse(f"build/run.wasm")
97126

127+
98128
def page():
99-
return Title("GPU Puzzles"), Body(
100-
Div(
129+
return (
130+
Title("GPU Puzzles"),
131+
Body(
101132
Div(
102-
CodeEditor(),
103-
style="width: 66vw; height:100vh; background-color: #333; float: left;",
104-
),
105-
Div(
106-
"Output",
107-
id="output",
108-
style="width: 34vw; height:100vh; background-color: #444; float: right;",
133+
Div(
134+
CodeEditor(initial_content=gelu_kernel),
135+
style="width: 66vw; height:100vh; background-color: #333; float: left;",
136+
),
137+
Div(
138+
"Output",
139+
id="output",
140+
style="width: 34vw; height:100vh; background-color: #444; float: right;",
141+
),
109142
),
143+
Script(bind_terminal),
144+
style="height: 100vh; overflow: hidden;",
110145
),
111-
bind_terminal,
112-
style="height: 100vh; overflow: hidden;"),
146+
)
113147

114148

115149
@rt("/")

‎experimental/fasthtml/shadertui_web.cpp‎

Lines changed: 99 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,99 @@
1+
#include "gpu.h"
2+
#include <array>
3+
#include <cstdio>
4+
#include <emscripten/emscripten.h>
5+
#include <future>
6+
#include <chrono>
7+
#include <string>
8+
#include <algorithm>
9+
10+
using namespace gpu;
11+
12+
static constexpr size_t kRows = 64;
13+
static constexpr size_t kCols = 96;
14+
15+
template <size_t rows, size_t cols>
16+
void rasterize(const std::array<float, rows * cols> &values,
17+
std::array<char, rows *(cols + 1)> &raster) {
18+
static const char intensity[] = " .`'^-+=*x17X$8#%@";
19+
for (size_t i = 0; i < rows; ++i) {
20+
for (size_t j = 0; j < cols; ++j) {
21+
size_t index =
22+
std::min(sizeof(intensity) - 2,
23+
std::max(0ul, static_cast<size_t>(values[i * cols + j] *
24+
(sizeof(intensity) - 2))));
25+
raster[i * (cols + 1) + j] = intensity[index];
26+
}
27+
raster[i * (cols + 1) + cols] = '\n';
28+
}
29+
}
30+
31+
float getCurrentTimeInMilliseconds(
32+
std::chrono::time_point<std::chrono::high_resolution_clock> &zeroTime) {
33+
std::chrono::duration<float> duration =
34+
std::chrono::high_resolution_clock::now() - zeroTime;
35+
return duration.count() * 1000.0f; // Convert to milliseconds
36+
}
37+
38+
EM_JS(void, js_print, (const char *str), {
39+
if (typeof window != 'undefined' && window.customPrint) {
40+
window.customPrint(UTF8ToString(str));
41+
} else {
42+
console.log(UTF8ToString(str));
43+
}
44+
});
45+
46+
extern "C" {
47+
EMSCRIPTEN_KEEPALIVE
48+
void executeShader(const char *shaderCode) {
49+
Context ctx = createContext();
50+
51+
std::array<float, kRows * kCols> screenArr;
52+
Tensor screen = createTensor(ctx, {kRows, kCols}, kf32, screenArr.data());
53+
54+
struct Params {
55+
float time;
56+
uint32_t screenWidth;
57+
uint32_t screenHeight;
58+
} params = {0.0f, kCols, kRows};
59+
60+
KernelCode shader{shaderCode, Shape{16, 16, 1}};
61+
Kernel renderKernel =
62+
createKernel(ctx, shader, Bindings{screen},
63+
cdiv({kCols, kRows, 1}, shader.workgroupSize), params);
64+
65+
std::array<char, kRows *(kCols + 1)> raster;
66+
67+
auto start = std::chrono::high_resolution_clock::now();
68+
69+
// Render a few frames
70+
for (int frame = 0; frame < 10; ++frame) {
71+
params.time = getCurrentTimeInMilliseconds(start);
72+
toGPU(ctx, params, renderKernel);
73+
74+
std::promise<void> promise;
75+
std::future<void> future = promise.get_future();
76+
dispatchKernel(ctx, renderKernel, promise);
77+
wait(ctx, future);
78+
79+
toCPU(ctx, screen, screenArr.data());
80+
rasterize<kRows, kCols>(screenArr, raster);
81+
82+
// Print the rasterized frame
83+
js_print(raster.data());
84+
85+
// Add a small delay between frames
86+
emscripten_sleep(100);
87+
}
88+
}
89+
}
90+
91+
#ifndef STANDALONE_WASM
92+
#include <emscripten/bind.h>
93+
EMSCRIPTEN_BINDINGS(module) {
94+
emscripten::function("executeShader", emscripten::optional_override(
95+
[](const std::string &shaderCode) {
96+
executeShader(shaderCode.c_str());
97+
}));
98+
}
99+
#endif

0 commit comments

Comments
 (0)