summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--CMakeLists.txt8
-rw-r--r--assets/final/demo_assets.txt1
-rw-r--r--assets/final/shaders/compute/gen_noise.wgsl26
-rw-r--r--src/gpu/effects/shaders.cc4
-rw-r--r--src/gpu/effects/shaders.h1
-rw-r--r--src/gpu/texture_manager.cc196
-rw-r--r--src/gpu/texture_manager.h20
-rw-r--r--src/tests/test_gpu_procedural.cc56
-rw-r--r--src/util/asset_manager.h1
-rw-r--r--tools/asset_packer.cc61
10 files changed, 371 insertions, 3 deletions
diff --git a/CMakeLists.txt b/CMakeLists.txt
index e90cb4a..26818c3 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -593,6 +593,14 @@ if(DEMO_BUILD_TESTS)
target_link_libraries(test_texture_manager PRIVATE 3d gpu audio procedural util ${DEMO_LIBS})
add_dependencies(test_texture_manager generate_demo_assets)
+ # GPU Procedural Texture Test
+ add_demo_test(test_gpu_procedural GpuProceduralTest
+ src/tests/test_gpu_procedural.cc
+ ${PLATFORM_SOURCES}
+ ${GEN_DEMO_CC})
+ target_link_libraries(test_gpu_procedural PRIVATE 3d gpu audio procedural util ${DEMO_LIBS})
+ add_dependencies(test_gpu_procedural generate_demo_assets)
+
# Gantt chart output test (bash script)
add_test(
NAME GanttOutputTest
diff --git a/assets/final/demo_assets.txt b/assets/final/demo_assets.txt
index 05eee17..1fdf2b4 100644
--- a/assets/final/demo_assets.txt
+++ b/assets/final/demo_assets.txt
@@ -52,6 +52,7 @@ SHADER_MESH, NONE, shaders/mesh_render.wgsl, "Mesh Rasterization Shader"
MESH_CUBE, NONE, test_mesh.obj, "A simple cube mesh"
DODECAHEDRON, NONE, dodecahedron.obj, "A dodecahedron mesh"
SHADER_VIGNETTE, NONE, shaders/vignette.wgsl, "Vignette Shader"
+SHADER_COMPUTE_GEN_NOISE, NONE, shaders/compute/gen_noise.wgsl, "GPU Noise Compute Shader"
CIRCLE_MASK_COMPUTE_SHADER, NONE, shaders/circle_mask_compute.wgsl, "Circle mask compute shader"
CIRCLE_MASK_RENDER_SHADER, NONE, shaders/circle_mask_render.wgsl, "Circle mask render shader"
MASKED_CUBE_SHADER, NONE, shaders/masked_cube.wgsl, "Masked cube shader"
diff --git a/assets/final/shaders/compute/gen_noise.wgsl b/assets/final/shaders/compute/gen_noise.wgsl
new file mode 100644
index 0000000..5c0babd
--- /dev/null
+++ b/assets/final/shaders/compute/gen_noise.wgsl
@@ -0,0 +1,26 @@
+// GPU procedural noise texture generator.
+// Uses compute shader for parallel texture generation.
+
+#include "math/noise"
+
+struct NoiseParams {
+ width: u32,
+ height: u32,
+ seed: f32,
+ frequency: f32,
+}
+
+@group(0) @binding(0) var output_tex: texture_storage_2d<rgba8unorm, write>;
+@group(0) @binding(1) var<uniform> params: NoiseParams;
+
+@compute @workgroup_size(8, 8, 1)
+fn main(@builtin(global_invocation_id) id: vec3<u32>) {
+ if (id.x >= params.width || id.y >= params.height) { return; }
+
+ let uv = vec2<f32>(f32(id.x) / f32(params.width),
+ f32(id.y) / f32(params.height));
+ let p = uv * params.frequency + params.seed;
+ let noise = noise_2d(p);
+
+ textureStore(output_tex, id.xy, vec4<f32>(noise, noise, noise, 1.0));
+}
diff --git a/src/gpu/effects/shaders.cc b/src/gpu/effects/shaders.cc
index 2e1cfe5..5c6dd37 100644
--- a/src/gpu/effects/shaders.cc
+++ b/src/gpu/effects/shaders.cc
@@ -99,6 +99,10 @@ const char* chroma_aberration_shader_wgsl =
SafeGetAsset(AssetId::ASSET_SHADER_CHROMA_ABERRATION);
+const char* gen_noise_compute_wgsl =
+
+ SafeGetAsset(AssetId::ASSET_SHADER_COMPUTE_GEN_NOISE);
+
const char* vignette_shader_wgsl =
SafeGetAsset(AssetId::ASSET_SHADER_VIGNETTE);
diff --git a/src/gpu/effects/shaders.h b/src/gpu/effects/shaders.h
index 50b4f32..b629e30 100644
--- a/src/gpu/effects/shaders.h
+++ b/src/gpu/effects/shaders.h
@@ -18,3 +18,4 @@ extern const char* solarize_shader_wgsl;
extern const char* distort_shader_wgsl;
extern const char* chroma_aberration_shader_wgsl;
extern const char* vignette_shader_wgsl;
+extern const char* gen_noise_compute_wgsl;
diff --git a/src/gpu/texture_manager.cc b/src/gpu/texture_manager.cc
index 0c30c94..aff106a 100644
--- a/src/gpu/texture_manager.cc
+++ b/src/gpu/texture_manager.cc
@@ -2,7 +2,10 @@
// It implements the TextureManager.
#include "gpu/texture_manager.h"
+#include "gpu/effects/shader_composer.h"
+#include "platform/platform.h"
#include <cstdio>
+#include <cstring>
#include <vector>
#if defined(DEMO_CROSS_COMPILE_WIN32)
@@ -18,6 +21,7 @@
void TextureManager::init(WGPUDevice device, WGPUQueue queue) {
device_ = device;
queue_ = queue;
+ noise_compute_pipeline_ = nullptr;
}
void TextureManager::shutdown() {
@@ -26,6 +30,10 @@ void TextureManager::shutdown() {
wgpuTextureRelease(pair.second.texture);
}
textures_.clear();
+ if (noise_compute_pipeline_) {
+ wgpuComputePipelineRelease(noise_compute_pipeline_);
+ noise_compute_pipeline_ = nullptr;
+ }
}
void TextureManager::create_procedural_texture(
@@ -112,3 +120,191 @@ WGPUTextureView TextureManager::get_texture_view(const std::string& name) {
}
return nullptr;
}
+
+void TextureManager::dispatch_noise_compute(WGPUTexture target,
+ const GpuProceduralParams& params) {
+ // Lazy-init compute pipeline
+ if (!noise_compute_pipeline_) {
+ extern const char* gen_noise_compute_wgsl;
+
+ // Resolve #include directives using ShaderComposer
+ ShaderComposer& composer = ShaderComposer::Get();
+ std::string resolved_shader = composer.Compose({}, gen_noise_compute_wgsl);
+
+ WGPUShaderSourceWGSL wgsl_src = {};
+ wgsl_src.chain.sType = WGPUSType_ShaderSourceWGSL;
+ wgsl_src.code = str_view(resolved_shader.c_str());
+ WGPUShaderModuleDescriptor shader_desc = {};
+ shader_desc.nextInChain = &wgsl_src.chain;
+ WGPUShaderModule shader_module =
+ wgpuDeviceCreateShaderModule(device_, &shader_desc);
+
+ // Bind group layout (storage texture + uniform)
+ WGPUBindGroupLayoutEntry bgl_entries[2] = {};
+ bgl_entries[0].binding = 0;
+ bgl_entries[0].visibility = WGPUShaderStage_Compute;
+ bgl_entries[0].storageTexture.access = WGPUStorageTextureAccess_WriteOnly;
+ bgl_entries[0].storageTexture.format = WGPUTextureFormat_RGBA8Unorm;
+ bgl_entries[0].storageTexture.viewDimension = WGPUTextureViewDimension_2D;
+
+ bgl_entries[1].binding = 1;
+ bgl_entries[1].visibility = WGPUShaderStage_Compute;
+ bgl_entries[1].buffer.type = WGPUBufferBindingType_Uniform;
+ bgl_entries[1].buffer.minBindingSize = 16; // sizeof(NoiseParams)
+
+ WGPUBindGroupLayoutDescriptor bgl_desc = {};
+ bgl_desc.entryCount = 2;
+ bgl_desc.entries = bgl_entries;
+ WGPUBindGroupLayout bind_group_layout =
+ wgpuDeviceCreateBindGroupLayout(device_, &bgl_desc);
+
+ WGPUPipelineLayoutDescriptor pl_desc = {};
+ pl_desc.bindGroupLayoutCount = 1;
+ pl_desc.bindGroupLayouts = &bind_group_layout;
+ WGPUPipelineLayout pipeline_layout =
+ wgpuDeviceCreatePipelineLayout(device_, &pl_desc);
+
+ WGPUComputePipelineDescriptor pipeline_desc = {};
+ pipeline_desc.layout = pipeline_layout;
+ pipeline_desc.compute.module = shader_module;
+ pipeline_desc.compute.entryPoint = str_view("main");
+
+ noise_compute_pipeline_ =
+ wgpuDeviceCreateComputePipeline(device_, &pipeline_desc);
+
+ wgpuPipelineLayoutRelease(pipeline_layout);
+ wgpuBindGroupLayoutRelease(bind_group_layout);
+ wgpuShaderModuleRelease(shader_module);
+ }
+
+ // Create uniform buffer (width, height, seed, frequency)
+ struct NoiseParams {
+ uint32_t width;
+ uint32_t height;
+ float seed;
+ float frequency;
+ };
+ NoiseParams uniform_data = {(uint32_t)params.width, (uint32_t)params.height,
+ params.params[0], params.params[1]};
+ WGPUBufferDescriptor buf_desc = {};
+ buf_desc.size = sizeof(NoiseParams);
+ buf_desc.usage = WGPUBufferUsage_Uniform | WGPUBufferUsage_CopyDst;
+ buf_desc.mappedAtCreation = WGPUOptionalBool_True;
+ WGPUBuffer uniform_buf = wgpuDeviceCreateBuffer(device_, &buf_desc);
+ void* mapped = wgpuBufferGetMappedRange(uniform_buf, 0, sizeof(NoiseParams));
+ memcpy(mapped, &uniform_data, sizeof(NoiseParams));
+ wgpuBufferUnmap(uniform_buf);
+
+ // Create storage texture view
+ WGPUTextureViewDescriptor view_desc = {};
+ view_desc.format = WGPUTextureFormat_RGBA8Unorm;
+ view_desc.dimension = WGPUTextureViewDimension_2D;
+ view_desc.mipLevelCount = 1;
+ view_desc.arrayLayerCount = 1;
+ WGPUTextureView target_view = wgpuTextureCreateView(target, &view_desc);
+
+ // Create bind group layout entries (must match pipeline)
+ WGPUBindGroupLayoutEntry bgl_entries[2] = {};
+ bgl_entries[0].binding = 0;
+ bgl_entries[0].visibility = WGPUShaderStage_Compute;
+ bgl_entries[0].storageTexture.access = WGPUStorageTextureAccess_WriteOnly;
+ bgl_entries[0].storageTexture.format = WGPUTextureFormat_RGBA8Unorm;
+ bgl_entries[0].storageTexture.viewDimension = WGPUTextureViewDimension_2D;
+ bgl_entries[1].binding = 1;
+ bgl_entries[1].visibility = WGPUShaderStage_Compute;
+ bgl_entries[1].buffer.type = WGPUBufferBindingType_Uniform;
+ bgl_entries[1].buffer.minBindingSize = 16;
+
+ WGPUBindGroupLayoutDescriptor bgl_desc = {};
+ bgl_desc.entryCount = 2;
+ bgl_desc.entries = bgl_entries;
+ WGPUBindGroupLayout bind_group_layout =
+ wgpuDeviceCreateBindGroupLayout(device_, &bgl_desc);
+
+ // Create bind group
+ WGPUBindGroupEntry bg_entries[2] = {};
+ bg_entries[0].binding = 0;
+ bg_entries[0].textureView = target_view;
+ bg_entries[1].binding = 1;
+ bg_entries[1].buffer = uniform_buf;
+ bg_entries[1].size = sizeof(NoiseParams);
+
+ WGPUBindGroupDescriptor bg_desc = {};
+ bg_desc.layout = bind_group_layout;
+ bg_desc.entryCount = 2;
+ bg_desc.entries = bg_entries;
+ WGPUBindGroup bind_group = wgpuDeviceCreateBindGroup(device_, &bg_desc);
+
+ // Dispatch compute
+ WGPUCommandEncoderDescriptor enc_desc = {};
+ WGPUCommandEncoder encoder = wgpuDeviceCreateCommandEncoder(device_, &enc_desc);
+ WGPUComputePassEncoder pass = wgpuCommandEncoderBeginComputePass(encoder, nullptr);
+ wgpuComputePassEncoderSetPipeline(pass, noise_compute_pipeline_);
+ wgpuComputePassEncoderSetBindGroup(pass, 0, bind_group, 0, nullptr);
+ wgpuComputePassEncoderDispatchWorkgroups(pass, (params.width + 7) / 8,
+ (params.height + 7) / 8, 1);
+ wgpuComputePassEncoderEnd(pass);
+
+ WGPUCommandBufferDescriptor cmd_desc = {};
+ WGPUCommandBuffer cmd = wgpuCommandEncoderFinish(encoder, &cmd_desc);
+ wgpuQueueSubmit(queue_, 1, &cmd);
+
+ // Cleanup
+ wgpuCommandBufferRelease(cmd);
+ wgpuCommandEncoderRelease(encoder);
+ wgpuComputePassEncoderRelease(pass);
+ wgpuBindGroupRelease(bind_group);
+ wgpuBindGroupLayoutRelease(bind_group_layout);
+ wgpuBufferRelease(uniform_buf);
+ wgpuTextureViewRelease(target_view);
+}
+
+void TextureManager::create_gpu_noise_texture(
+ const std::string& name, const GpuProceduralParams& params) {
+ // Create storage texture
+ WGPUTextureDescriptor tex_desc = {};
+ tex_desc.usage =
+ WGPUTextureUsage_StorageBinding | WGPUTextureUsage_TextureBinding;
+ tex_desc.dimension = WGPUTextureDimension_2D;
+ tex_desc.size = {(uint32_t)params.width, (uint32_t)params.height, 1};
+ tex_desc.format = WGPUTextureFormat_RGBA8Unorm;
+ tex_desc.mipLevelCount = 1;
+ tex_desc.sampleCount = 1;
+ WGPUTexture texture = wgpuDeviceCreateTexture(device_, &tex_desc);
+
+ // Generate via compute
+ dispatch_noise_compute(texture, params);
+
+ // Create view for sampling
+ WGPUTextureViewDescriptor view_desc = {};
+ view_desc.format = WGPUTextureFormat_RGBA8Unorm;
+ view_desc.dimension = WGPUTextureViewDimension_2D;
+ view_desc.mipLevelCount = 1;
+ view_desc.arrayLayerCount = 1;
+ WGPUTextureView view = wgpuTextureCreateView(texture, &view_desc);
+
+ // Store texture
+ GpuTexture gpu_tex;
+ gpu_tex.texture = texture;
+ gpu_tex.view = view;
+ gpu_tex.width = params.width;
+ gpu_tex.height = params.height;
+ textures_[name] = gpu_tex;
+
+#if !defined(STRIP_ALL)
+ printf("Generated GPU noise texture: %s (%dx%d)\n", name.c_str(),
+ params.width, params.height);
+#endif
+}
+
+#if !defined(STRIP_ALL)
+WGPUTextureView TextureManager::get_or_generate_gpu_texture(
+ const std::string& name, const GpuProceduralParams& params) {
+ auto it = textures_.find(name);
+ if (it != textures_.end()) {
+ return it->second.view;
+ }
+ create_gpu_noise_texture(name, params);
+ return textures_[name].view;
+}
+#endif
diff --git a/src/gpu/texture_manager.h b/src/gpu/texture_manager.h
index 23fdbe8..0cffe0c 100644
--- a/src/gpu/texture_manager.h
+++ b/src/gpu/texture_manager.h
@@ -23,6 +23,13 @@ struct GpuTexture {
int height;
};
+struct GpuProceduralParams {
+ int width;
+ int height;
+ const float* params;
+ int num_params;
+};
+
class TextureManager {
public:
void init(WGPUDevice device, WGPUQueue queue);
@@ -36,11 +43,24 @@ class TextureManager {
void create_texture(const std::string& name, int width, int height,
const uint8_t* data);
+ // GPU procedural generation
+ void create_gpu_noise_texture(const std::string& name,
+ const GpuProceduralParams& params);
+
+#if !defined(STRIP_ALL)
+ // On-demand lazy generation (stripped in final builds)
+ WGPUTextureView get_or_generate_gpu_texture(const std::string& name,
+ const GpuProceduralParams& params);
+#endif
+
// Retrieves a texture view by name (returns nullptr if not found)
WGPUTextureView get_texture_view(const std::string& name);
private:
+ void dispatch_noise_compute(WGPUTexture target,
+ const GpuProceduralParams& params);
WGPUDevice device_;
WGPUQueue queue_;
std::map<std::string, GpuTexture> textures_;
+ WGPUComputePipeline noise_compute_pipeline_ = nullptr;
};
diff --git a/src/tests/test_gpu_procedural.cc b/src/tests/test_gpu_procedural.cc
new file mode 100644
index 0000000..f68eb4f
--- /dev/null
+++ b/src/tests/test_gpu_procedural.cc
@@ -0,0 +1,56 @@
+// This file is part of the 64k demo project.
+// Tests GPU procedural texture generation.
+
+#include "gpu/gpu.h"
+#include "gpu/texture_manager.h"
+#include "platform/platform.h"
+#include <cstdio>
+
+int main() {
+ printf("GPU Procedural Test: Starting...\n");
+
+ // Minimal GPU initialization for testing
+ PlatformState platform = platform_init(false, 256, 256);
+ if (!platform.window) {
+ fprintf(stderr, "Error: Failed to create window\n");
+ return 1;
+ }
+
+ gpu_init(&platform);
+ const GpuContext* ctx = gpu_get_context();
+
+ // Initialize shader composer (needed for #include resolution)
+ extern void InitShaderComposer();
+ InitShaderComposer();
+
+ // Create TextureManager
+ TextureManager tex_mgr;
+ tex_mgr.init(ctx->device, ctx->queue);
+
+ // Test GPU noise generation
+ GpuProceduralParams params = {};
+ params.width = 256;
+ params.height = 256;
+ float proc_params[2] = {0.0f, 4.0f}; // seed, frequency
+ params.params = proc_params;
+ params.num_params = 2;
+
+ tex_mgr.create_gpu_noise_texture("test_noise", params);
+
+ // Verify texture exists
+ WGPUTextureView view = tex_mgr.get_texture_view("test_noise");
+ if (!view) {
+ fprintf(stderr, "Error: GPU noise texture not created\n");
+ tex_mgr.shutdown();
+ gpu_shutdown();
+ return 1;
+ }
+
+ printf("SUCCESS: GPU noise texture created (256x256)\n");
+
+ // Cleanup
+ tex_mgr.shutdown();
+ gpu_shutdown();
+ platform_shutdown(&platform);
+ return 0;
+}
diff --git a/src/util/asset_manager.h b/src/util/asset_manager.h
index 1e0638c..168bfca 100644
--- a/src/util/asset_manager.h
+++ b/src/util/asset_manager.h
@@ -10,6 +10,7 @@ struct AssetRecord {
size_t size; // Size of the asset data
bool is_procedural; // True if data was dynamically allocated by a procedural
// generator
+ bool is_gpu_procedural; // True if GPU compute shader generates texture
const char* proc_func_name_str; // Name of procedural generation function
// (string literal)
const float* proc_params; // Parameters for procedural generation (static,
diff --git a/tools/asset_packer.cc b/tools/asset_packer.cc
index 0d26cf6..72592ae 100644
--- a/tools/asset_packer.cc
+++ b/tools/asset_packer.cc
@@ -52,6 +52,7 @@ struct AssetBuildInfo {
std::string name;
std::string filename; // Original filename for static assets
bool is_procedural;
+ bool is_gpu_procedural;
std::string proc_func_name; // Function name string
std::vector<float> proc_params; // Parameters for procedural function
@@ -182,9 +183,62 @@ int main(int argc, char* argv[]) {
info.params_array_name = "ASSET_PROC_PARAMS_" + info.name;
info.func_name_str_name = "ASSET_PROC_FUNC_STR_" + info.name;
info.is_procedural = false;
+ info.is_gpu_procedural = false;
- if (compression_type_str.rfind("PROC(", 0) == 0) {
+ if (compression_type_str.rfind("PROC_GPU(", 0) == 0) {
info.is_procedural = true;
+ info.is_gpu_procedural = true;
+ size_t open_paren = compression_type_str.find('(');
+ size_t close_paren = compression_type_str.rfind(')');
+ if (open_paren == std::string::npos ||
+ close_paren == std::string::npos) {
+ fprintf(stderr,
+ "Error: Invalid PROC_GPU() syntax for asset: %s, string: [%s]\n",
+ info.name.c_str(), compression_type_str.c_str());
+ return 1;
+ }
+ std::string func_and_params_str = compression_type_str.substr(
+ open_paren + 1, close_paren - open_paren - 1);
+
+ size_t params_start = func_and_params_str.find(',');
+ if (params_start != std::string::npos) {
+ std::string params_str = func_and_params_str.substr(params_start + 1);
+ info.proc_func_name = func_and_params_str.substr(0, params_start);
+
+ size_t current_pos = 0;
+ while (current_pos < params_str.length()) {
+ size_t comma_pos = params_str.find(',', current_pos);
+ std::string param_val_str =
+ (comma_pos == std::string::npos)
+ ? params_str.substr(current_pos)
+ : params_str.substr(current_pos, comma_pos - current_pos);
+ param_val_str.erase(0, param_val_str.find_first_not_of(" \t\r\n"));
+ param_val_str.erase(param_val_str.find_last_not_of(" \t\r\n") + 1);
+ try {
+ info.proc_params.push_back(std::stof(param_val_str));
+ } catch (...) {
+ fprintf(stderr, "Error: Invalid proc param for %s: %s\n",
+ info.name.c_str(), param_val_str.c_str());
+ return 1;
+ }
+ if (comma_pos == std::string::npos)
+ break;
+ current_pos = comma_pos + 1;
+ }
+ } else {
+ info.proc_func_name = func_and_params_str;
+ }
+
+ // Validate GPU procedural function name
+ if (info.proc_func_name != "gen_noise") {
+ fprintf(stderr,
+ "Error: PROC_GPU only supports gen_noise, got: %s for asset: %s\n",
+ info.proc_func_name.c_str(), info.name.c_str());
+ return 1;
+ }
+ } else if (compression_type_str.rfind("PROC(", 0) == 0) {
+ info.is_procedural = true;
+ info.is_gpu_procedural = false;
size_t open_paren = compression_type_str.find('(');
size_t close_paren = compression_type_str.rfind(')');
if (open_paren == std::string::npos ||
@@ -500,12 +554,13 @@ int main(int argc, char* argv[]) {
for (const auto& info : asset_build_infos) {
fprintf(assets_data_cc_file, " { ");
if (info.is_procedural) {
- fprintf(assets_data_cc_file, "nullptr, 0, true, %s, %s, %zu",
+ fprintf(assets_data_cc_file, "nullptr, 0, true, %s, %s, %s, %zu",
+ info.is_gpu_procedural ? "true" : "false",
info.func_name_str_name.c_str(), info.params_array_name.c_str(),
info.proc_params.size());
} else {
fprintf(assets_data_cc_file,
- "%s, ASSET_SIZE_%s, false, nullptr, nullptr, 0",
+ "%s, ASSET_SIZE_%s, false, false, nullptr, nullptr, 0",
info.data_array_name.c_str(), info.name.c_str());
}
fprintf(assets_data_cc_file, " },\n");
generated by cgit v1.2.3 (git 2.51.0) at 2026-02-10 05:11:53 +0000