From a0b7e4c0d80ea8ae2da2d16ce5554b7590764119 Mon Sep 17 00:00:00 2001 From: Vicente Ferrari Smith Date: Tue, 28 Apr 2026 19:46:32 +0200 Subject: [PATCH] . --- metal/CMakeLists.txt | 0 metal/definition.cpp | 3 + renderer/graphics.cpp | 3 + renderer/graphics.h | 14 + renderer/graphics_private.h | 8 + renderer/init.h | 47 -- renderer/metal/AAPLMathUtilities.cpp | 3 + renderer/metal/AAPLMathUtilities.h | 8 + renderer/metal/metal.cpp | 216 ++++++ renderer/metal/metal.h | 22 + renderer/{ => metal}/renderer.cpp | 4 +- renderer/metal/renderer.h | 402 +++++++++++ renderer/metal/vertex_data.h | 8 + renderer/sprite.cpp | 2 +- renderer/texture.cpp | 37 +- renderer/texture_sheet.cpp | 2 +- renderer/vulkan/renderer.cpp | 828 +++++++++++++++++++++++ renderer/{ => vulkan}/renderer.h | 2 +- renderer/{init.cpp => vulkan/vulkan.cpp} | 58 ++ renderer/vulkan/vulkan.h | 25 + renderer/webgpu/renderer.cpp | 5 + renderer/webgpu/renderer.h | 16 + renderer/webgpu/utils_emscripten.cpp | 79 +++ renderer/webgpu/utils_emscripten.h | 79 +++ renderer/webgpu/webgpu.cpp | 5 + renderer/webgpu/webgpu.h | 16 + shaders/compute.slang | 0 shaders/compute.wgsl | 13 + shaders/shader.air | Bin 0 -> 10416 bytes shaders/shader.metal | 0 shaders/shader.wgsl | 63 ++ shaders/shaders.metallib | Bin 0 -> 21840 bytes 32 files changed, 1880 insertions(+), 88 deletions(-) create mode 100644 metal/CMakeLists.txt create mode 100644 metal/definition.cpp create mode 100644 renderer/graphics.cpp create mode 100644 renderer/graphics.h create mode 100644 renderer/graphics_private.h delete mode 100644 renderer/init.h create mode 100644 renderer/metal/AAPLMathUtilities.cpp create mode 100644 renderer/metal/AAPLMathUtilities.h create mode 100644 renderer/metal/metal.cpp create mode 100644 renderer/metal/metal.h rename renderer/{ => metal}/renderer.cpp (99%) create mode 100644 renderer/metal/renderer.h create mode 100644 renderer/metal/vertex_data.h create mode 100644 renderer/vulkan/renderer.cpp rename renderer/{ => vulkan}/renderer.h (99%) rename renderer/{init.cpp => vulkan/vulkan.cpp} (92%) create mode 100644 renderer/vulkan/vulkan.h create mode 100644 renderer/webgpu/renderer.cpp create mode 100644 renderer/webgpu/renderer.h create mode 100644 renderer/webgpu/utils_emscripten.cpp create mode 100644 renderer/webgpu/utils_emscripten.h create mode 100644 renderer/webgpu/webgpu.cpp create mode 100644 renderer/webgpu/webgpu.h create mode 100644 shaders/compute.slang create mode 100644 shaders/compute.wgsl create mode 100644 shaders/shader.air create mode 100644 shaders/shader.metal create mode 100644 shaders/shader.wgsl create mode 100644 shaders/shaders.metallib diff --git a/metal/CMakeLists.txt b/metal/CMakeLists.txt new file mode 100644 index 0000000..e69de29 diff --git a/metal/definition.cpp b/metal/definition.cpp new file mode 100644 index 0000000..24c2895 --- /dev/null +++ b/metal/definition.cpp @@ -0,0 +1,3 @@ +// +// Created by Vicente Ferrari Smith on 26.02.26. +// diff --git a/renderer/graphics.cpp b/renderer/graphics.cpp new file mode 100644 index 0000000..072d86b --- /dev/null +++ b/renderer/graphics.cpp @@ -0,0 +1,3 @@ +// +// Created by Vicente Ferrari Smith on 02.03.26. +// diff --git a/renderer/graphics.h b/renderer/graphics.h new file mode 100644 index 0000000..750ab32 --- /dev/null +++ b/renderer/graphics.h @@ -0,0 +1,14 @@ +// +// Created by Vicente Ferrari Smith on 26.02.26. +// + +#ifndef V_RENDERER_H +#define V_RENDERER_H + + +struct Graphics { + Graphics(); +}; + + +#endif //V_RENDERER_H \ No newline at end of file diff --git a/renderer/graphics_private.h b/renderer/graphics_private.h new file mode 100644 index 0000000..c272ff8 --- /dev/null +++ b/renderer/graphics_private.h @@ -0,0 +1,8 @@ +// +// Created by Vicente Ferrari Smith on 02.03.26. +// + +#ifndef V_GRAPHICS_PRIVATE_H +#define V_GRAPHICS_PRIVATE_H + +#endif //V_GRAPHICS_PRIVATE_H \ No newline at end of file diff --git a/renderer/init.h b/renderer/init.h deleted file mode 100644 index 7475a0f..0000000 --- a/renderer/init.h +++ /dev/null @@ -1,47 +0,0 @@ -// -// Created by Vicente Ferrari Smith on 12.02.26. -// - -#ifndef V_INIT_H -#define V_INIT_H - -#include -#include -#include -#include - -inline VkInstance instance{}; -inline VkPhysicalDevice physicalDevice{}; -inline VkDevice device{}; -inline VkQueue graphics_queue{}; -inline uint32_t queueFamily{}; - -inline VkSurfaceKHR surface{}; -inline VkDebugUtilsMessengerEXT debugMessenger{}; - -inline VmaAllocator allocator{}; - -inline constexpr uint32_t MAX_FRAMES_IN_FLIGHT = 2; -inline constexpr uint32_t MAX_VERTICES_PER_BATCH = 65536; - -inline VkSwapchainKHR swapchain; -inline VkExtent2D swapchain_extent; -inline VkSurfaceFormatKHR swapchain_format{ - VK_FORMAT_B8G8R8A8_UNORM, - VK_COLOR_SPACE_SRGB_NONLINEAR_KHR -}; -inline std::vector renderFinished; - -inline std::vector images; -inline std::vector imageViews; -inline std::vector imageLayouts; - -void createSwapchain(GLFWwindow* window); - -int createInstance(GLFWwindow* window); -void createSurface(GLFWwindow* window); -void pickPhysicalDevice(); -void createDevice(); - - -#endif //V_INIT_H diff --git a/renderer/metal/AAPLMathUtilities.cpp b/renderer/metal/AAPLMathUtilities.cpp new file mode 100644 index 0000000..7e3ca1c --- /dev/null +++ b/renderer/metal/AAPLMathUtilities.cpp @@ -0,0 +1,3 @@ +// +// Created by Vicente Ferrari Smith on 27.02.26. +// diff --git a/renderer/metal/AAPLMathUtilities.h b/renderer/metal/AAPLMathUtilities.h new file mode 100644 index 0000000..a098e82 --- /dev/null +++ b/renderer/metal/AAPLMathUtilities.h @@ -0,0 +1,8 @@ +// +// Created by Vicente Ferrari Smith on 27.02.26. +// + +#ifndef V_AAPLMATHUTILITIES_H +#define V_AAPLMATHUTILITIES_H + +#endif //V_AAPLMATHUTILITIES_H diff --git a/renderer/metal/metal.cpp b/renderer/metal/metal.cpp new file mode 100644 index 0000000..bc254c6 --- /dev/null +++ b/renderer/metal/metal.cpp @@ -0,0 +1,216 @@ +// +// Created by Vicente Ferrari Smith on 26.02.26. +// + +#include "init.h" +#include "../graphics.h" +#include +#include +#include +#include +#define GLFW_EXPOSE_NATIVE_COCOA +#include +#include +#include +#include "vertex_data.h" + +Device metal_device{}; +MTL::Buffer* triangle_vertex_buffer{}; +MTL::CommandQueue *queue{}; +CA::MetalLayer *metal_layer{}; +MTL::RenderPipelineState *pipeline_state{}; +CA::MetalDrawable *metal_drawable{}; +MTL::CommandBuffer* metal_command_buffer{}; + +MTL::Function *vertex_shader{}; +MTL::Function *fragment_shader{}; + +void create_window(GLFWwindow *window) { + void *ns_window = glfwGetCocoaWindow(window); + if (!ns_window) { + throw std::runtime_error("Failed to get Cocoa window from GLFWwindow"); + } + + SEL contentViewSel = sel_registerName("contentView"); + id content_view = ((id (*)(id, SEL))objc_msgSend)((id)ns_window, contentViewSel); + + SEL setWantsLayerSel = sel_registerName("setWantsLayer:"); + ((void (*)(id, SEL, bool))objc_msgSend)(content_view, setWantsLayerSel, true); + + metal_layer = CA::MetalLayer::layer()->retain(); + + SEL setLayerSel = sel_registerName("setLayer:"); + ((void (*)(id, SEL, id))objc_msgSend)(content_view, setLayerSel, (id)metal_layer); + + metal_layer->retain(); + metal_layer->setDevice(metal_device.device); + metal_layer->setPixelFormat(MTL::PixelFormatRGBA16Float); + metal_layer->setFramebufferOnly(true); + metal_layer->setDrawableSize(CGSizeMake(800, 600)); + CGColorSpaceRef p3Space = CGColorSpaceCreateWithName(kCGColorSpaceSRGB); + metal_layer->setColorspace(p3Space); +} + +void encode_render_command(MTL::RenderCommandEncoder *renderCommandEncoder) { + renderCommandEncoder->setRenderPipelineState(pipeline_state); + renderCommandEncoder->setVertexBuffer(triangle_vertex_buffer, 0, 0); + MTL::PrimitiveType typeTriangle = MTL::PrimitiveTypeTriangle; + NS::UInteger vertexStart = 0; + NS::UInteger vertexCount = 6; + renderCommandEncoder->drawPrimitives(typeTriangle, vertexStart, vertexCount); +} + +void send_render_command() { + metal_command_buffer = queue->commandBuffer(); + + MTL::RenderPassDescriptor* renderPassDescriptor = MTL::RenderPassDescriptor::alloc()->init(); + MTL::RenderPassColorAttachmentDescriptor *cd = renderPassDescriptor->colorAttachments()->object(0); + cd->setTexture(metal_drawable->texture()); + cd->setLoadAction(MTL::LoadActionClear); + cd->setClearColor(MTL::ClearColor( + 100.0f / 255.0f, + 149.0f / 255.0f, + 237.0f / 255.0f, + 1.0 + )); + cd->setStoreAction(MTL::StoreActionStore); + + MTL::RenderCommandEncoder* renderCommandEncoder = metal_command_buffer->renderCommandEncoder(renderPassDescriptor); + encode_render_command(renderCommandEncoder); + renderCommandEncoder->endEncoding(); + + metal_command_buffer->presentDrawable(metal_drawable); + metal_command_buffer->commit(); + metal_command_buffer->waitUntilCompleted(); + + renderPassDescriptor->release(); +} + +void LoadMetalShader(const std::string &shader_path, + const std::string &vertex_fn_name, + const std::string &fragment_fn_name) +{ + NS::Error *error = nullptr; + MTL::Library *library = nullptr; + + auto ends_with = [](const std::string& s, const char* suf) -> bool { + const size_t n = std::strlen(suf); + return s.size() >= n && s.compare(s.size() - n, n, suf) == 0; + }; + + if (ends_with(shader_path, ".metal")) { + // Compile from source at runtime + std::ifstream file(shader_path, std::ios::in | std::ios::binary); + if (!file.is_open()) { + throw std::runtime_error("Failed to open .metal source file"); + } + std::string src; + file.seekg(0, std::ios::end); + src.resize(static_cast(file.tellg())); + file.seekg(0, std::ios::beg); + file.read(src.data(), static_cast(src.size())); + file.close(); + + NS::String* source = NS::String::string(src.c_str(), NS::UTF8StringEncoding); + MTL::CompileOptions* opts = MTL::CompileOptions::alloc()->init(); + library = metal_device.device->newLibrary(source, opts, &error); + opts->release(); + } else { + // Load a precompiled metallib from file path + NS::String *nsPath = NS::String::string(shader_path.c_str(), NS::UTF8StringEncoding); + library = metal_device.device->newLibrary(nsPath, &error); + } + + if (error || library == nullptr) { + if (error) { + // Extract the actual compiler error message + const char* errorMessage = error->localizedDescription()->utf8String(); + std::string detailedError = "Metal Library Error: "; + detailedError += errorMessage; + + // It is good practice to release the error object if it exists + error->release(); + + throw std::runtime_error(detailedError); + } + throw std::runtime_error("Failed to create Metal library (Unknown error)"); + } + NS::String *vname = NS::String::string(vertex_fn_name.c_str(), NS::UTF8StringEncoding); + NS::String *fname = NS::String::string(fragment_fn_name.c_str(), NS::UTF8StringEncoding); + vertex_shader = library->newFunction(vname); + fragment_shader = library->newFunction(fname); + + if (vertex_shader == nullptr || fragment_shader == nullptr) { + throw std::runtime_error("Failed to create Metal shader functions"); + } + + library->release(); +} + +void create_render_pipeline() { + LoadMetalShader("shaders/shader.metal", "vertex_main", "fragment_main"); + + MTL::RenderPipelineDescriptor* renderPipelineDescriptor = MTL::RenderPipelineDescriptor::alloc()->init(); + renderPipelineDescriptor->setLabel(NS::String::string("Triangle Rendering Pipeline", NS::ASCIIStringEncoding)); + renderPipelineDescriptor->setVertexFunction(vertex_shader); + renderPipelineDescriptor->setFragmentFunction(fragment_shader); + assert(renderPipelineDescriptor); + const MTL::PixelFormat pixel_format = metal_layer->pixelFormat(); + renderPipelineDescriptor->colorAttachments()->object(0)->setPixelFormat(pixel_format); + + NS::Error* error; + pipeline_state = metal_device.device->newRenderPipelineState(renderPipelineDescriptor, &error); + renderPipelineDescriptor->release(); +} + +void create_command_queue() { + queue = metal_device.device->newCommandQueue(); +} + +void create_triangle() { + VertexData square_vertices[] = { + {{-0.5, -0.5}, {1.0, 0.0, 0.0, 1.0}}, + {{0.5, -0.5}, {0.0, 1.0, 0.0, 1.0}}, + {{0.5, 0.5}, {0.0, 0.0, 1.0, 1.0}}, + + {{0.5, 0.5}, {0.0, 0.0, 1.0, 1.0}}, + {{-0.5, 0.5}, {0.0, 1.0, 0.0, 1.0}}, + {{-0.5, -0.5}, {1.0, 0.0, 0.0, 1.0}}, + }; + + triangle_vertex_buffer = metal_device.device->newBuffer(&square_vertices, + sizeof(square_vertices), + MTL::ResourceStorageModeShared); + +} + +void graphics_init(GLFWwindow *window) { + std::println("wow, we are on macos!! crazy!!"); + + create_device(); + create_window(window); + create_triangle(); + create_command_queue(); + create_render_pipeline(); +} + +void graphics_deinit() { + +} + +void begin_frame() { + +} + +void end_frame() { + auto pPool = NS::AutoreleasePool::alloc()->init(); + metal_drawable = metal_layer->nextDrawable(); + + send_render_command(); + + pPool->release(); +} + +void create_device() { + metal_device.device = MTL::CreateSystemDefaultDevice(); +} diff --git a/renderer/metal/metal.h b/renderer/metal/metal.h new file mode 100644 index 0000000..9b44768 --- /dev/null +++ b/renderer/metal/metal.h @@ -0,0 +1,22 @@ +// +// Created by Vicente Ferrari Smith on 26.02.26. +// + +#ifndef M_INIT_H +#define M_INIT_H + +#include +#define GLFW_EXPOSE_NATIVE_COCOA +#import + +#include +#include +#include + +struct Device { + MTL::Device *device; +}; + +void create_device(); + +#endif //M_INIT_H \ No newline at end of file diff --git a/renderer/renderer.cpp b/renderer/metal/renderer.cpp similarity index 99% rename from renderer/renderer.cpp rename to renderer/metal/renderer.cpp index f3642fd..17ad577 100644 --- a/renderer/renderer.cpp +++ b/renderer/metal/renderer.cpp @@ -2,12 +2,12 @@ // Created by Vicente Ferrari Smith on 13.02.26. // -#include "renderer.h" +#include "../graphics.h" #include #include "init.h" -#include "sprite.h" +#include "../sprite.h" #include #include diff --git a/renderer/metal/renderer.h b/renderer/metal/renderer.h new file mode 100644 index 0000000..2f21359 --- /dev/null +++ b/renderer/metal/renderer.h @@ -0,0 +1,402 @@ +// +// Created by Vicente Ferrari Smith on 13.02.26. +// + +#ifndef V_RENDERER_H +#define V_RENDERER_H + +#include "init.h" +#include +#include +#define GLM_FORCE_RADIANS +#define GLM_FORCE_DEPTH_ZERO_TO_ONE +#define GLM_ENABLE_EXPERIMENTAL +#include +#include +#include "glm/gtx/string_cast.hpp" +#include +#include "../sprite.h" +#include "../texture.h" +#include +#include +#include +#include +#include + +inline Slang::ComPtr slangGlobalSession; + +enum class PROJECTION_TYPE : uint8_t { + NONE, + ORTHOGRAPHIC_WORLD, + ORTHOGRAPHIC_WINDOW, + PERSPECTIVE_WORLD, + PERSPECTIVE_WINDOW, + COUNT, +}; + +struct vertex_p2_s2_st2_col4_a1_u32 { + glm::vec2 pos; + glm::vec2 scale; + glm::vec2 uv; + glm::vec4 color; + float alpha; + uint32_t textureID; + + static VkVertexInputBindingDescription getBindingDescription() { + return {0, sizeof(vertex_p2_s2_st2_col4_a1_u32), VK_VERTEX_INPUT_RATE_VERTEX}; + } + + static std::array getAttributeDescriptions() { + return { + { + {0, 0, VK_FORMAT_R32G32_SFLOAT, offsetof(vertex_p2_s2_st2_col4_a1_u32, pos)}, + {1, 0, VK_FORMAT_R32G32_SFLOAT, offsetof(vertex_p2_s2_st2_col4_a1_u32, scale)}, + {2, 0, VK_FORMAT_R32G32_SFLOAT, offsetof(vertex_p2_s2_st2_col4_a1_u32, uv)}, + {3, 0, VK_FORMAT_R32G32B32A32_SFLOAT, offsetof(vertex_p2_s2_st2_col4_a1_u32, color)}, + {4, 0, VK_FORMAT_R32_SFLOAT, offsetof(vertex_p2_s2_st2_col4_a1_u32, alpha)}, + {5, 0, VK_FORMAT_R32_UINT, offsetof(vertex_p2_s2_st2_col4_a1_u32, textureID)}, + } + }; + } +}; + +// commands + +enum class PipelineType : uint8_t { + None, + TexturedQuad, + ColoredQuad, + Line, + Text, + Chunk +}; + +struct TexturedQuadCmd { + glm::vec2 position; + glm::vec2 size; + glm::vec2 uvMin; + glm::vec2 uvMax; + glm::vec4 color; + uint16_t textureID; +}; + +struct ColoredQuadCmd { + glm::vec2 pos; + glm::vec2 scale; + glm::vec4 color; +}; + +struct LineCmd { + glm::vec2 start; + glm::vec2 end; + glm::vec4 color; +}; + +// struct TextCmd { +// Font* font; +// std::string text; +// glm::vec2 position; +// glm::vec4 color; +// }; + +struct ChunkCmd { + VkBuffer vertexBuffer; + VkBuffer indexBuffer; + uint32_t indexCount; +}; + +struct SortKey { + uint16_t depth; // world Z or Y-sorted depth + uint16_t materialID; // texture sheet, font atlas, etc. + uint8_t pipeline; // PipelineType + + bool operator<(const SortKey& b) const; +}; + +struct RenderCommand { + + SortKey key; + PipelineType pipeline; + + union { + TexturedQuadCmd textured_quad; + ColoredQuadCmd colored_quad; + LineCmd line; + // TextCmd text; + ChunkCmd chunk; + }; +}; + +//////////////////////////////////////////////////////////////////////////////////////////////// + +struct AllocatedBuffer { + VkBuffer buffer; + VmaAllocation allocation; + VmaAllocationInfo info; +}; + +struct GPUMeshBuffers { + AllocatedBuffer indexBuffer; + AllocatedBuffer vertexBuffer; + VkDeviceAddress vertexBufferAddress; +}; + +struct Renderer { + std::vector commands{}; + + VkDescriptorSetLayout descriptor_set_layout{}; + VkPipelineLayout pipelineLayout{}; + VkPipeline textured_quad_pipeline{}; + VkPipeline colored_quad_pipeline{}; + VkPipeline line_pipeline{}; + VkPipeline text_pipeline{}; + VkPipeline chunk_pipeline{}; + VkDescriptorSet set{}; + + VkSampler defaultSampler{}; + + uint32_t nextTextureSlot = 0; + + struct Frame { + VkCommandPool commandPool{}; + VkCommandBuffer command_buffer{}; + + VkSemaphore imageAvailable{}; + VkFence in_flight_fence{}; + + AllocatedBuffer vertexBuffer{}; + }; + + std::vector frames; + uint32_t currentFrame = 0; + + VkDescriptorPool descriptorPool{}; + std::vector textureSets{}; + + void begin_frame(); + void end_frame(); + void flush(); + + void submit_sprite(glm::vec2 pos, const sprite_t &sprite); + void submit_quad(glm::vec2 pos, glm::vec2 scale); + + explicit Renderer(GLFWwindow *window); + void create_pipeline_layout(); + void createFrameResources(); + void create_default_sampler(); + void recordCommandBuffer( + VkCommandBuffer cmd, + VkImage image, + VkImageView imageView, + VkExtent2D extent, + VkImageLayout oldLayout, + const Frame &frame, + const std::vector &vertices) const; + void immediate_submit(std::function&& func) const; + void transition_image_layout(VkCommandBuffer cmd, VkImage image, VkImageLayout oldLayout, VkImageLayout newLayout) const; + VkImageView create_image_view(VkImage image, VkFormat format) const; + AllocatedBuffer create_buffer(size_t allocSize, VkBufferUsageFlags usage, VmaMemoryUsage memoryUsage); + void destroy_buffer(const AllocatedBuffer& buffer); + // GPUMeshBuffers uploadMesh(std::span indices, std::span vertices); + void upload_vertex_buffer( + VkCommandBuffer cmd, + const Frame &frame, + std::span vertices) const; + + [[nodiscard]] VkPipeline get_pipeline(PipelineType type) const; + // void bind_material(VkCommandBuffer cmd, uint16_t materialID); + void create_descriptor_pool(); + void update_bindless_slot(uint32_t slot, VkImageView view, VkSampler sampler) const; + + // Returns the resource info so the Manager can store it + void upload_texture( + int w, + int h, + const void* pixels, + VkImage *image, + VmaAllocation *allocation, + VkImageView *view, + uint32_t *descriptor_index); + + template + VkPipeline create_graphics_pipeline( + VkDevice device, + VkPipelineLayout layout, + VkFormat colorFormat, + // VkShaderModule vertShader, + // VkShaderModule fragShader, + VkPrimitiveTopology topology, + bool enableBlending) + { + + auto slangTargets{ std::to_array({ { + .format = SLANG_SPIRV, + .profile = slangGlobalSession->findProfile("spirv_1_4") + } })}; + auto slangOptions{ std::to_array({ { + slang::CompilerOptionName::EmitSpirvDirectly, + {slang::CompilerOptionValueKind::Int, 1} + } })}; + slang::SessionDesc slangSessionDesc{ + .targets = slangTargets.data(), + .targetCount = SlangInt(slangTargets.size()), + .defaultMatrixLayoutMode = SLANG_MATRIX_LAYOUT_COLUMN_MAJOR, + .compilerOptionEntries = slangOptions.data(), + .compilerOptionEntryCount = uint32_t(slangOptions.size()) + }; + Slang::ComPtr slangSession; + slangGlobalSession->createSession(slangSessionDesc, slangSession.writeRef()); + + Slang::ComPtr slangModule{ + slangSession->loadModuleFromSource("triangle", "shaders/shader.slang", nullptr, nullptr) + }; + Slang::ComPtr spirv; + slangModule->getTargetCode(0, spirv.writeRef()); + + VkShaderModuleCreateInfo shaderModuleCI{ + .sType = VK_STRUCTURE_TYPE_SHADER_MODULE_CREATE_INFO, + .codeSize = spirv->getBufferSize(), + .pCode = (uint32_t*)spirv->getBufferPointer() + }; + VkShaderModule shaderModule{}; + vkCreateShaderModule(device, &shaderModuleCI, nullptr, &shaderModule); + + auto vsCode = loadFile("shaders/triangle.vert.spv"); + auto fsCode = loadFile("shaders/triangle.frag.spv"); + + VkShaderModuleCreateInfo smci{ + .sType = VK_STRUCTURE_TYPE_SHADER_MODULE_CREATE_INFO + }; + + smci.codeSize = vsCode.size(); + smci.pCode = reinterpret_cast(vsCode.data()); + // VkShaderModule vs; + // vkCreateShaderModule(device, &smci, nullptr, &vs); + + smci.codeSize = fsCode.size(); + smci.pCode = reinterpret_cast(fsCode.data()); + // VkShaderModule fs; + // vkCreateShaderModule(device, &smci, nullptr, &fs); + + // --- Shaders --- + std::vector shaderStages{ + { .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO, + .stage = VK_SHADER_STAGE_VERTEX_BIT, + .module = shaderModule, .pName = "main"}, + { .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO, + .stage = VK_SHADER_STAGE_FRAGMENT_BIT, + .module = shaderModule, .pName = "main" } + }; + + // --- Vertex Input (Generic) --- + auto binding = T::getBindingDescription(); + auto attrs = T::getAttributeDescriptions(); + + // --- Vertex Input (Matching our vertex_p2_st2_col4 struct) --- + VkPipelineVertexInputStateCreateInfo vi{ + .sType = VK_STRUCTURE_TYPE_PIPELINE_VERTEX_INPUT_STATE_CREATE_INFO, + .vertexBindingDescriptionCount = 1, + .pVertexBindingDescriptions = &binding, + .vertexAttributeDescriptionCount = attrs.size(), + .pVertexAttributeDescriptions = attrs.data(), + }; + + // --- Input Assembly (Changes based on Topology parameter) --- + VkPipelineInputAssemblyStateCreateInfo ia{VK_STRUCTURE_TYPE_PIPELINE_INPUT_ASSEMBLY_STATE_CREATE_INFO}; + ia.topology = topology; + + // --- Blending (Changes based on enableBlending parameter) --- + VkPipelineColorBlendAttachmentState colorBlend{ + .blendEnable = enableBlending ? VK_TRUE : VK_FALSE, + .srcColorBlendFactor = VK_BLEND_FACTOR_SRC_ALPHA, + .dstColorBlendFactor = VK_BLEND_FACTOR_ONE_MINUS_SRC_ALPHA, + .colorBlendOp = VK_BLEND_OP_ADD, + .srcAlphaBlendFactor = VK_BLEND_FACTOR_ONE, + .dstAlphaBlendFactor = VK_BLEND_FACTOR_ZERO, + .alphaBlendOp = VK_BLEND_OP_ADD, + .colorWriteMask = 0xF + }; + + // --- Boilerplate (Standard 2D Defaults) --- + VkPipelineViewportStateCreateInfo vp{ + VK_STRUCTURE_TYPE_PIPELINE_VIEWPORT_STATE_CREATE_INFO, + nullptr, + 0, + 1, + nullptr, + 1, + nullptr + }; + + VkPipelineRasterizationStateCreateInfo rs{ + VK_STRUCTURE_TYPE_PIPELINE_RASTERIZATION_STATE_CREATE_INFO, + nullptr, + 0, + 0, + 0, + VK_POLYGON_MODE_FILL, + VK_CULL_MODE_NONE, + VK_FRONT_FACE_COUNTER_CLOCKWISE, + 0, + 0, + 0, + 0, + 1.0f + }; + + VkPipelineMultisampleStateCreateInfo ms{ + VK_STRUCTURE_TYPE_PIPELINE_MULTISAMPLE_STATE_CREATE_INFO, + nullptr, + 0, + VK_SAMPLE_COUNT_1_BIT + }; + + VkPipelineColorBlendStateCreateInfo cb{ + VK_STRUCTURE_TYPE_PIPELINE_COLOR_BLEND_STATE_CREATE_INFO, + nullptr, + 0, + 0, + VK_LOGIC_OP_AND, + 1, + &colorBlend + }; + + VkDynamicState dyns[] = { VK_DYNAMIC_STATE_VIEWPORT, VK_DYNAMIC_STATE_SCISSOR }; + VkPipelineDynamicStateCreateInfo ds{ + VK_STRUCTURE_TYPE_PIPELINE_DYNAMIC_STATE_CREATE_INFO, + nullptr, + 0, + 2, + dyns + }; + + VkPipelineRenderingCreateInfo rci{ + VK_STRUCTURE_TYPE_PIPELINE_RENDERING_CREATE_INFO, + nullptr, + 0, + 1, + &colorFormat + }; + + VkGraphicsPipelineCreateInfo gpci{ + .sType = VK_STRUCTURE_TYPE_GRAPHICS_PIPELINE_CREATE_INFO, + .pNext = &rci, + .stageCount = (uint32_t) shaderStages.size(), + .pStages = shaderStages.data(), + .pVertexInputState = &vi, + .pInputAssemblyState = &ia, + .pViewportState = &vp, + .pRasterizationState = &rs, + .pMultisampleState = &ms, + .pColorBlendState = &cb, + .pDynamicState = &ds, + .layout = layout + }; + + VkPipeline pipeline; + vkCreateGraphicsPipelines(device, VK_NULL_HANDLE, 1, &gpci, nullptr, &pipeline); + return pipeline; + } +}; + +#endif //V_RENDERER_H diff --git a/renderer/metal/vertex_data.h b/renderer/metal/vertex_data.h new file mode 100644 index 0000000..39d1eae --- /dev/null +++ b/renderer/metal/vertex_data.h @@ -0,0 +1,8 @@ +// +// Created by Vicente Ferrari Smith on 27.02.26. +// + +#ifndef V_VERTEX_DATA_H +#define V_VERTEX_DATA_H + +#endif //V_VERTEX_DATA_H diff --git a/renderer/sprite.cpp b/renderer/sprite.cpp index 148b117..dfab6ac 100644 --- a/renderer/sprite.cpp +++ b/renderer/sprite.cpp @@ -2,4 +2,4 @@ // Created by Vicente Ferrari Smith on 14.02.26. // -#include "sprite.h" \ No newline at end of file +#include "../sprite.h" diff --git a/renderer/texture.cpp b/renderer/texture.cpp index eac67bf..43aac3d 100644 --- a/renderer/texture.cpp +++ b/renderer/texture.cpp @@ -1,38 +1,3 @@ // -// Created by Vicente Ferrari Smith on 14.02.26. +// Created by Vicente Ferrari Smith on 01.03.26. // - -#include "texture.h" -#include -#include "renderer.h" - -TextureManager::TextureManager() { - -} - -Texture TextureManager::load(const std::string& path, Renderer &renderer) { - // Dedup: Don't load the same file twice! - // if (path_to_id.contains(path)) return path_to_id[path]; - - int w, h, ch; - unsigned char* data = stbi_load(path.c_str(), &w, &h, &ch, STBI_rgb_alpha); - - // Tell the renderer to make the GPU version - Texture res; - res.width = w; - res.height = h; - res.channels = STBI_rgb_alpha; - res.srgb = true; - renderer.upload_texture(w, h, data, &res.image, &res.allocation, &res.view, &res.descriptor_index); - - stbi_image_free(data); - - res.id = path; - res.path = path; - res.uploaded = true; - - textures[path] = res; - // path_to_id[path] = id; - - return res; // This is the textureID for your sprites -} diff --git a/renderer/texture_sheet.cpp b/renderer/texture_sheet.cpp index 54342e6..c984cd1 100644 --- a/renderer/texture_sheet.cpp +++ b/renderer/texture_sheet.cpp @@ -2,4 +2,4 @@ // Created by Vicente Ferrari Smith on 14.02.26. // -#include "texture_sheet.h" +#include "../texture_sheet.h" diff --git a/renderer/vulkan/renderer.cpp b/renderer/vulkan/renderer.cpp new file mode 100644 index 0000000..5dfbfc1 --- /dev/null +++ b/renderer/vulkan/renderer.cpp @@ -0,0 +1,828 @@ +// +// Created by Vicente Ferrari Smith on 13.02.26. +// + +#include "../Grpahics.h" + +#include + +#include "init.h" +#include "../sprite.h" +#include +#include + +extern int32_t window_width; +extern int32_t window_height; + +bool SortKey::operator<(const SortKey& b) const { + if (depth != b.depth) return depth < b.depth; + if (pipeline != b.pipeline) return pipeline < b.pipeline; + return materialID < b.materialID; +} + +Renderer::Renderer(GLFWwindow *window) { + + create_pipeline_layout(); + colored_quad_pipeline = create_graphics_pipeline( + device, + pipelineLayout, + swapchain_format.format, + VK_PRIMITIVE_TOPOLOGY_TRIANGLE_LIST, + true + ); + create_default_sampler(); + create_descriptor_pool(); + createFrameResources(); +} + +void Renderer::begin_frame() { + commands.clear(); +} + +void Renderer::flush() { + +} + +void Renderer::submit_quad(glm::vec2 pos, glm::vec2 scale) { + RenderCommand cmd {}; + cmd.pipeline = PipelineType::ColoredQuad; + cmd.key = { + (uint16_t) pos.y, + 0, + (uint8_t) PipelineType::ColoredQuad + }; + + cmd.colored_quad = { + .pos = pos, + .scale = scale, + .color = {0, 1, 1, 1}, + }; + + commands.push_back(cmd); +} + +void Renderer::submit_sprite(glm::vec2 pos, const sprite_t &sprite) { + RenderCommand cmd {}; + cmd.pipeline = PipelineType::TexturedQuad; + cmd.key = { + (uint16_t) pos.y, + 0, + (uint8_t) PipelineType::TexturedQuad + }; + + cmd.textured_quad = { + .position = pos, + .size = {0, 0}, + .uvMin = {0, 0}, + .uvMax = {0, 0}, + .color = {1, 1, 1, 1}, + .textureID = 0, + }; + + commands.push_back(cmd); + + // assert(started == true, "You can't submit without having started the renderer first."); + // renderable : Renderable; + // renderable.type = .Sprite; + // + // if sprite.window_space + // renderable.projection_type = .ORTHOGRAPHIC_WINDOW; + // else + // renderable.projection_type = .ORTHOGRAPHIC_WORLD; + // + // renderable.pos = pos; + // renderable.sprite.texture_sheet = sprite.texture_sheet; + // renderable.sprite.texture_cell = sprite.texture_cell; + // renderable.sprite.origin = sprite.origin; + // renderable.sprite.scale = sprite.scale; + // renderable.sprite.colour = sprite.colour; + // renderable.sprite.alpha = alpha; + // + // array_add(*renderer.renderable_list, renderable); +} + +void Renderer::create_pipeline_layout() { + std::array bindings = { + { + { + .binding = 0, + .descriptorType = VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER, + .descriptorCount = nextTextureSlot, + .stageFlags = VK_SHADER_STAGE_FRAGMENT_BIT + } + } + }; + + VkDescriptorBindingFlags flags[1] = { + VK_DESCRIPTOR_BINDING_PARTIALLY_BOUND_BIT | VK_DESCRIPTOR_BINDING_UPDATE_AFTER_BIND_BIT + }; + + VkDescriptorSetLayoutBindingFlagsCreateInfo layoutFlags{ + .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_BINDING_FLAGS_CREATE_INFO, + .bindingCount = 1, + .pBindingFlags = flags + }; + + VkDescriptorSetLayoutCreateInfo dslci{ + .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO, + .pNext = &layoutFlags, + // .flags = VK_DESCRIPTOR_SET_LAYOUT_CREATE_UPDATE_AFTER_BIND_POOL_BIT, + .bindingCount = bindings.size(), + .pBindings = bindings.data() + }; + + vkCreateDescriptorSetLayout(device, &dslci, nullptr, &descriptor_set_layout); + + VkPushConstantRange push_constant{ + .stageFlags = VK_SHADER_STAGE_VERTEX_BIT, + .offset = 0, + .size = sizeof(glm::mat4), + }; + + VkPipelineLayoutCreateInfo plci{ + .sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO, + .pNext = nullptr, + .flags = 0, + .setLayoutCount = 1, + .pSetLayouts = &descriptor_set_layout, + .pushConstantRangeCount = 1, + .pPushConstantRanges = &push_constant, + }; + + vkCreatePipelineLayout(device, &plci, nullptr, &pipelineLayout); +} + +void Renderer::createFrameResources() { + + const VkSemaphoreCreateInfo seci{ + .sType = VK_STRUCTURE_TYPE_SEMAPHORE_CREATE_INFO, + }; + + VkFenceCreateInfo fenceInfo{ + .sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO, + .flags = VK_FENCE_CREATE_SIGNALED_BIT, + }; + + VkCommandPoolCreateInfo cpci{ + .sType = VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO, + .flags = VK_COMMAND_POOL_CREATE_RESET_COMMAND_BUFFER_BIT, + .queueFamilyIndex = queueFamily + }; + + frames.resize(MAX_FRAMES_IN_FLIGHT); + + for (uint32_t i = 0; i < MAX_FRAMES_IN_FLIGHT; ++i) { + Frame &frame = frames[i]; + + vkCreateSemaphore(device, &seci, nullptr, &frame.imageAvailable); + + vkCreateFence(device, &fenceInfo, nullptr, &frame.in_flight_fence); + + + vkCreateCommandPool(device, &cpci, nullptr, &frame.commandPool); + + const VkCommandBufferAllocateInfo cbai{ + .sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO, + .commandPool = frame.commandPool, + .level = VK_COMMAND_BUFFER_LEVEL_PRIMARY, + .commandBufferCount = 1 + }; + + vkAllocateCommandBuffers(device, &cbai, &frame.command_buffer); + + VkBufferCreateInfo bufferInfo = { + .sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO, + .size = 1024 * 1024 * 4, + .usage = VK_BUFFER_USAGE_VERTEX_BUFFER_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT, + }; + + VmaAllocationCreateInfo allocCreateInfo = {}; + allocCreateInfo.usage = VMA_MEMORY_USAGE_AUTO; + allocCreateInfo.flags = VMA_ALLOCATION_CREATE_HOST_ACCESS_SEQUENTIAL_WRITE_BIT | + VMA_ALLOCATION_CREATE_HOST_ACCESS_ALLOW_TRANSFER_INSTEAD_BIT | + VMA_ALLOCATION_CREATE_MAPPED_BIT; + + vmaCreateBuffer( + allocator, + &bufferInfo, + &allocCreateInfo, + &frame.vertexBuffer.buffer, + &frame.vertexBuffer.allocation, + &frame.vertexBuffer.info); + } +} + +AllocatedBuffer Renderer::create_buffer(size_t allocSize, VkBufferUsageFlags usage, VmaMemoryUsage memoryUsage) { + // allocate buffer + VkBufferCreateInfo bufferInfo = {.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO}; + bufferInfo.pNext = nullptr; + bufferInfo.size = allocSize; + + bufferInfo.usage = usage; + + VmaAllocationCreateInfo vmaallocInfo = {}; + vmaallocInfo.usage = memoryUsage; + vmaallocInfo.flags = VMA_ALLOCATION_CREATE_MAPPED_BIT; + AllocatedBuffer newBuffer{}; + + // allocate the buffer + vmaCreateBuffer(allocator, &bufferInfo, &vmaallocInfo, &newBuffer.buffer, &newBuffer.allocation, &newBuffer.info); + + return newBuffer; +} + +void Renderer::destroy_buffer(const AllocatedBuffer& buffer) { + vmaDestroyBuffer(allocator, buffer.buffer, buffer.allocation); +} + +// GPUMeshBuffers Renderer::uploadMesh(std::span indices, std::span vertices) { +// const size_t vertexBufferSize = vertices.size() * sizeof(vertex_p2_st2_col4_a1_u32); +// const size_t indexBufferSize = indices.size() * sizeof(uint32_t); +// +// GPUMeshBuffers newSurface; +// +// //create vertex buffer +// newSurface.vertexBuffer = create_buffer(vertexBufferSize, VK_BUFFER_USAGE_STORAGE_BUFFER_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT | VK_BUFFER_USAGE_SHADER_DEVICE_ADDRESS_BIT, +// VMA_MEMORY_USAGE_GPU_ONLY); +// +// //find the adress of the vertex buffer +// VkBufferDeviceAddressInfo deviceAdressInfo{ .sType = VK_STRUCTURE_TYPE_BUFFER_DEVICE_ADDRESS_INFO,.buffer = newSurface.vertexBuffer.buffer }; +// newSurface.vertexBufferAddress = vkGetBufferDeviceAddress(device, &deviceAdressInfo); +// +// //create index buffer +// newSurface.indexBuffer = create_buffer(indexBufferSize, VK_BUFFER_USAGE_INDEX_BUFFER_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT, +// VMA_MEMORY_USAGE_GPU_ONLY); +// +// AllocatedBuffer staging = create_buffer(vertexBufferSize + indexBufferSize, VK_BUFFER_USAGE_TRANSFER_SRC_BIT, VMA_MEMORY_USAGE_CPU_ONLY); +// +// void* data = staging.allocation->GetMappedData(); +// +// // copy vertex buffer +// memcpy(data, vertices.data(), vertexBufferSize); +// // copy index buffer +// memcpy((char*)data + vertexBufferSize, indices.data(), indexBufferSize); +// +// immediate_submit([&](VkCommandBuffer cmd) { +// VkBufferCopy vertexCopy{ 0 }; +// vertexCopy.dstOffset = 0; +// vertexCopy.srcOffset = 0; +// vertexCopy.size = vertexBufferSize; +// +// vkCmdCopyBuffer(cmd, staging.buffer, newSurface.vertexBuffer.buffer, 1, &vertexCopy); +// +// VkBufferCopy indexCopy{ 0 }; +// indexCopy.dstOffset = 0; +// indexCopy.srcOffset = vertexBufferSize; +// indexCopy.size = indexBufferSize; +// +// vkCmdCopyBuffer(cmd, staging.buffer, newSurface.indexBuffer.buffer, 1, &indexCopy); +// }); +// +// destroy_buffer(staging); +// +// return newSurface; +// +// } + +VkPipeline Renderer::get_pipeline(PipelineType type) const { + switch (type) { + case PipelineType::TexturedQuad: return textured_quad_pipeline; + case PipelineType::ColoredQuad: return colored_quad_pipeline; + case PipelineType::Line: return line_pipeline; + default: return {}; + } +} + +// void Renderer::bind_material(VkCommandBuffer cmd, uint16_t materialID) { +// // In a real app, you'd have an array/map: std::vector textureSets; +// VkDescriptorSet set = textureSets[materialID]; +// +// vkCmdBindDescriptorSets( +// cmd, +// VK_PIPELINE_BIND_POINT_GRAPHICS, +// pipelineLayout, // Our shared layout +// 0, // Starting at Set 0 +// 1, // Binding 1 set +// &set, +// 0, nullptr +// ); +// } + +void Renderer::create_descriptor_pool() { + VkDescriptorPoolSize pool_sizes[] = { + { VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER, nextTextureSlot }, + }; + + VkDescriptorPoolCreateInfo pool_info{ + .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_POOL_CREATE_INFO, + .flags = VK_DESCRIPTOR_POOL_CREATE_UPDATE_AFTER_BIND_BIT, + .maxSets = 1, + .poolSizeCount = 1, + .pPoolSizes = pool_sizes + }; + + vkCreateDescriptorPool(device, &pool_info, nullptr, &descriptorPool); + + VkDescriptorSetAllocateInfo alloc_info{ + .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_ALLOCATE_INFO, + .descriptorPool = descriptorPool, + .descriptorSetCount = 1, + .pSetLayouts = &descriptor_set_layout + }; + + vkAllocateDescriptorSets(device, &alloc_info, &set); +} + +void Renderer::update_bindless_slot(uint32_t slot, VkImageView view, VkSampler sampler) const { + VkDescriptorImageInfo image_info{ + .sampler = sampler, + .imageView = view, + .imageLayout = VK_IMAGE_LAYOUT_SHADER_READ_ONLY_OPTIMAL + }; + + VkWriteDescriptorSet write{ + .sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET, + .dstSet = set, + .dstArrayElement = slot, + .descriptorCount = 1, + .descriptorType = VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER, + .pImageInfo = &image_info + }; + + vkUpdateDescriptorSets(device, 1, &write, 0, nullptr); +} + +void Renderer::upload_texture( + const int w, + const int h, + const void* pixels, + VkImage *image, + VmaAllocation *allocation, + VkImageView *view, + uint32_t *descriptor_index) +{ + VkDeviceSize imageSize = w * h * 4; + + // --- 1. Create Staging Buffer (CPU Visible) --- + VkBuffer stagingBuffer; + VmaAllocation stagingAlloc; + + VkBufferCreateInfo stagingBufferInfo = { .sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO }; + stagingBufferInfo.size = imageSize; + stagingBufferInfo.usage = VK_BUFFER_USAGE_TRANSFER_SRC_BIT; + + VmaAllocationCreateInfo stagingAllocCreateInfo = { + .flags = VMA_ALLOCATION_CREATE_HOST_ACCESS_SEQUENTIAL_WRITE_BIT | VMA_ALLOCATION_CREATE_MAPPED_BIT, + .usage = VMA_MEMORY_USAGE_AUTO, + }; + + VmaAllocationInfo stagingResultInfo; + vmaCreateBuffer(allocator, &stagingBufferInfo, &stagingAllocCreateInfo, &stagingBuffer, &stagingAlloc, &stagingResultInfo); + + // Copy raw pixels into the mapped memory provided by VMA + memcpy(stagingResultInfo.pMappedData, pixels, imageSize); + + // --- 2. Create GPU Image (Device Local / Tiled) --- + VkExtent3D imageExtent = { (uint32_t) w, (uint32_t) h, 1 }; + + VkImageCreateInfo imageInfo = { + .sType = VK_STRUCTURE_TYPE_IMAGE_CREATE_INFO, + .imageType = VK_IMAGE_TYPE_2D, + .format = VK_FORMAT_R8G8B8A8_UNORM, + .extent = imageExtent, + .mipLevels = 1, + .arrayLayers = 1, + .samples = VK_SAMPLE_COUNT_1_BIT, + .tiling = VK_IMAGE_TILING_OPTIMAL, + .usage = VK_IMAGE_USAGE_TRANSFER_DST_BIT | VK_IMAGE_USAGE_SAMPLED_BIT, + .initialLayout = VK_IMAGE_LAYOUT_UNDEFINED + }; + + VmaAllocationCreateInfo imageAllocCreateInfo = { + .usage = VMA_MEMORY_USAGE_AUTO, + .priority = 1.0f, + }; + + vmaCreateImage(allocator, &imageInfo, &imageAllocCreateInfo, image, allocation, nullptr); + + // --- 3. The Transfer --- + immediate_submit([&](VkCommandBuffer cmd) { + // Transition image from UNDEFINED to TRANSFER_DST + transition_image_layout(cmd, *image, VK_IMAGE_LAYOUT_UNDEFINED, VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL); + + VkBufferImageCopy copyRegion = {}; + copyRegion.bufferOffset = 0; + copyRegion.bufferRowLength = 0; + copyRegion.bufferImageHeight = 0; + copyRegion.imageSubresource.aspectMask = VK_IMAGE_ASPECT_COLOR_BIT; + copyRegion.imageSubresource.mipLevel = 0; + copyRegion.imageSubresource.baseArrayLayer = 0; + copyRegion.imageSubresource.layerCount = 1; + copyRegion.imageExtent = imageExtent; + + vkCmdCopyBufferToImage(cmd, stagingBuffer, *image, VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL, 1, ©Region); + + // Transition image from TRANSFER_DST to SHADER_READ_ONLY + transition_image_layout(cmd, *image, VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL, VK_IMAGE_LAYOUT_SHADER_READ_ONLY_OPTIMAL); + }); + + // Clean up temporary staging resources + vmaDestroyBuffer(allocator, stagingBuffer, stagingAlloc); + + // --- 4. Finalize Handles --- + *view = create_image_view(*image, imageInfo.format); + + // Register in your Bindless Array (Set 0, Binding 0, Index N) + *descriptor_index = nextTextureSlot++; + update_bindless_slot(*descriptor_index, *view, defaultSampler); +} + +void Renderer::immediate_submit(std::function&& func) const { + VkCommandBufferAllocateInfo allocInfo{ .sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO }; + allocInfo.commandPool = frames[currentFrame].commandPool; // Use a pool created with VK_COMMAND_POOL_CREATE_RESET_COMMAND_BUFFER_BIT + allocInfo.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY; + allocInfo.commandBufferCount = 1; + + VkCommandBuffer cmd; + vkAllocateCommandBuffers(device, &allocInfo, &cmd); + + VkCommandBufferBeginInfo beginInfo{ .sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO }; + beginInfo.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT; + + vkBeginCommandBuffer(cmd, &beginInfo); + + // Execute the code passed in the lambda + func(cmd); + + vkEndCommandBuffer(cmd); + + VkSubmitInfo submit{ .sType = VK_STRUCTURE_TYPE_SUBMIT_INFO }; + submit.commandBufferCount = 1; + submit.pCommandBuffers = &cmd; + + // Submit and wait + vkQueueSubmit(graphics_queue, 1, &submit, VK_NULL_HANDLE); + vkQueueWaitIdle(graphics_queue); + + vkFreeCommandBuffers(device, frames[currentFrame].commandPool, 1, &cmd); +} + +void Renderer::transition_image_layout(VkCommandBuffer cmd, VkImage image, VkImageLayout oldLayout, VkImageLayout newLayout) const { + VkImageMemoryBarrier2 barrier{ .sType = VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER_2 }; + + barrier.oldLayout = oldLayout; + barrier.newLayout = newLayout; + barrier.image = image; + barrier.subresourceRange = { VK_IMAGE_ASPECT_COLOR_BIT, 0, 1, 0, 1 }; + + // Simple synchronization: wait for all previous commands, and block all future ones + // You can optimize these masks later, but this is safe for a 2D engine + barrier.srcStageMask = VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT; + barrier.srcAccessMask = VK_ACCESS_2_MEMORY_WRITE_BIT; + barrier.dstStageMask = VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT; + barrier.dstAccessMask = VK_ACCESS_2_MEMORY_READ_BIT | VK_ACCESS_2_MEMORY_WRITE_BIT; + + VkDependencyInfo dep{ .sType = VK_STRUCTURE_TYPE_DEPENDENCY_INFO }; + dep.imageMemoryBarrierCount = 1; + dep.pImageMemoryBarriers = &barrier; + + vkCmdPipelineBarrier2(cmd, &dep); +} + +VkImageView Renderer::create_image_view(VkImage image, VkFormat format) const { + VkImageViewCreateInfo viewInfo{ + .sType = VK_STRUCTURE_TYPE_IMAGE_VIEW_CREATE_INFO, + .image = image, + .viewType = VK_IMAGE_VIEW_TYPE_2D, + .format = format, + }; + + // Default component mapping (R,G,B,A) + viewInfo.components.r = VK_COMPONENT_SWIZZLE_IDENTITY; + viewInfo.components.g = VK_COMPONENT_SWIZZLE_IDENTITY; + viewInfo.components.b = VK_COMPONENT_SWIZZLE_IDENTITY; + viewInfo.components.a = VK_COMPONENT_SWIZZLE_IDENTITY; + + // Which part of the image to look at (Mip 0, Layer 0) + viewInfo.subresourceRange.aspectMask = VK_IMAGE_ASPECT_COLOR_BIT; + viewInfo.subresourceRange.baseMipLevel = 0; + viewInfo.subresourceRange.levelCount = 1; + viewInfo.subresourceRange.baseArrayLayer = 0; + viewInfo.subresourceRange.layerCount = 1; + + VkImageView view; + vkCreateImageView(device, &viewInfo, nullptr, &view); + return view; +} + +void Renderer::create_default_sampler() { + VkSamplerCreateInfo samplerInfo{ + .sType = VK_STRUCTURE_TYPE_SAMPLER_CREATE_INFO, + // For crisp pixel art, use NEAREST. For smooth textures, use LINEAR. + .magFilter = VK_FILTER_NEAREST, + .minFilter = VK_FILTER_NEAREST, + + .mipmapMode = VK_SAMPLER_MIPMAP_MODE_NEAREST, + + // How to handle "out of bounds" UVs + .addressModeU = VK_SAMPLER_ADDRESS_MODE_CLAMP_TO_EDGE, + .addressModeV = VK_SAMPLER_ADDRESS_MODE_CLAMP_TO_EDGE, + .addressModeW = VK_SAMPLER_ADDRESS_MODE_CLAMP_TO_EDGE, + + // Optimization: turn off things we don't need for simple 2D + .anisotropyEnable = VK_FALSE, + .maxAnisotropy = 1.0f, + .compareEnable = VK_FALSE, + .compareOp = VK_COMPARE_OP_ALWAYS, + .borderColor = VK_BORDER_COLOR_INT_OPAQUE_BLACK, + .unnormalizedCoordinates = VK_FALSE, + }; + + vkCreateSampler(device, &samplerInfo, nullptr, &defaultSampler); +} + +void Renderer::end_frame() { + + Frame &frame = frames[currentFrame]; + + vkWaitForFences(device, 1, &frame.in_flight_fence, VK_TRUE, UINT64_MAX); + vkResetFences(device, 1, &frame.in_flight_fence); + + uint32_t imageIndex; + vkAcquireNextImageKHR( + device, + swapchain, + UINT64_MAX, + frame.imageAvailable, + VK_NULL_HANDLE, + &imageIndex + ); + + commands = counting_sort_descending(commands, [](const RenderCommand &cmd){ + return cmd.key.depth; + }); + + std::vector vertices; + + for (auto& cmd : commands) { + + + switch (cmd.pipeline) { + case PipelineType::ColoredQuad: { + const auto &q = cmd.colored_quad; + + // Calculate spatial corners + //float x0 = q.position.x; + //float y0 = q.position.y; + //float x1 = q.position.x + q.size.x; + //float y1 = q.position.y + q.size.y; + + // Calculate UV corners + // float u0 = q.uvMin.x; + // float v0 = q.uvMin.y; + // float u1 = q.uvMax.x; + // float v1 = q.uvMax.y; + + // Define the 4 corners of the quad + vertex_p2_s2_st2_col4_a1_u32 vTL = { q.pos, q.scale, {0, 0}, {1, 0, 0, 0}, 1, 0 }; + vertex_p2_s2_st2_col4_a1_u32 vTR = { q.pos, q.scale, {0, 0}, q.color, 1, 0 }; + vertex_p2_s2_st2_col4_a1_u32 vBL = { q.pos, q.scale, {0, 0}, q.color, 1, 0 }; + vertex_p2_s2_st2_col4_a1_u32 vBR = { q.pos, q.scale, {0, 0}, q.color, 1, 0 }; + + // vertex_p2_st2_col4_a1_u32 vTL = { {x0, y0}, {u0, v0}, q.color, 1, q.textureID }; + // vertex_p2_st2_col4_a1_u32 vTR = { {x1, y0}, {u1, v0}, q.color, 1, q.textureID }; + // vertex_p2_st2_col4_a1_u32 vBL = { {x0, y1}, {u0, v1}, q.color, 1, q.textureID }; + // vertex_p2_st2_col4_a1_u32 vBR = { {x1, y1}, {u1, v1}, q.color, 1, q.textureID }; + + vertices.push_back(vTL); + vertices.push_back(vBL); + vertices.push_back(vTR); + + vertices.push_back(vTR); + vertices.push_back(vBL); + vertices.push_back(vBR); + + break; + } + default: + break; + } + } + + VkCommandBuffer cmd = frame.command_buffer; + vkResetCommandBuffer(cmd, 0); + + VkCommandBufferBeginInfo cbBI { + .sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO, + .flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT + }; + vkBeginCommandBuffer(cmd, &cbBI); + + recordCommandBuffer( + cmd, + images[imageIndex], + imageViews[imageIndex], + swapchain_extent, + imageLayouts[imageIndex], + frame, + vertices + ); + + vkEndCommandBuffer(cmd); + + imageLayouts[imageIndex] = VK_IMAGE_LAYOUT_PRESENT_SRC_KHR; + + VkSemaphoreSubmitInfo waitBinary{ + .sType = VK_STRUCTURE_TYPE_SEMAPHORE_SUBMIT_INFO, + .semaphore = frame.imageAvailable, + .stageMask = VK_PIPELINE_STAGE_2_COLOR_ATTACHMENT_OUTPUT_BIT + }; + + VkSemaphoreSubmitInfo signalBinary{ + .sType = VK_STRUCTURE_TYPE_SEMAPHORE_SUBMIT_INFO, + .semaphore = renderFinished[imageIndex], + .stageMask = VK_PIPELINE_STAGE_2_ALL_GRAPHICS_BIT + }; + + VkCommandBufferSubmitInfo cmdInfo{ + .sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_SUBMIT_INFO, + .commandBuffer = cmd, + }; + + const VkSubmitInfo2 submit{ + .sType = VK_STRUCTURE_TYPE_SUBMIT_INFO_2, + .waitSemaphoreInfoCount = 1, + .pWaitSemaphoreInfos = &waitBinary, + .commandBufferInfoCount = 1, + .pCommandBufferInfos = &cmdInfo, + .signalSemaphoreInfoCount = 1, + .pSignalSemaphoreInfos = &signalBinary, + }; + + vkQueueSubmit2(graphics_queue, 1, &submit, frame.in_flight_fence); + + VkPresentInfoKHR present{ + .sType = VK_STRUCTURE_TYPE_PRESENT_INFO_KHR, + .waitSemaphoreCount = 1, + .pWaitSemaphores = &renderFinished[imageIndex], + .swapchainCount = 1, + .pSwapchains = &swapchain, + .pImageIndices = &imageIndex, + }; + + vkQueuePresentKHR(graphics_queue, &present); + + currentFrame = (currentFrame + 1) % MAX_FRAMES_IN_FLIGHT; +} + +void Renderer::upload_vertex_buffer( + VkCommandBuffer cmd, + const Frame &frame, + std::span vertices) const +{ + VkMemoryPropertyFlags memPropFlags; + vmaGetAllocationMemoryProperties(allocator, frame.vertexBuffer.allocation, &memPropFlags); + + if(memPropFlags & VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT) { + // The Allocation ended up in a mappable memory. + // Calling vmaCopyMemoryToAllocation() does vmaMapMemory(), memcpy(), vmaUnmapMemory(), and vmaFlushAllocation(). + VkResult result = vmaCopyMemoryToAllocation(allocator, vertices.data(), frame.vertexBuffer.allocation, 0, vertices.size() * sizeof(vertex_p2_s2_st2_col4_a1_u32)); + // Check result... + + VkBufferMemoryBarrier bufMemBarrier = { VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER }; + bufMemBarrier.srcAccessMask = VK_ACCESS_HOST_WRITE_BIT; + bufMemBarrier.dstAccessMask = VK_ACCESS_UNIFORM_READ_BIT; + bufMemBarrier.srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED; + bufMemBarrier.dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED; + bufMemBarrier.buffer = frame.vertexBuffer.buffer; + bufMemBarrier.offset = 0; + bufMemBarrier.size = VK_WHOLE_SIZE; + + // It's important to insert a buffer memory barrier here to ensure writing to the buffer has finished. + vkCmdPipelineBarrier(cmd, VK_PIPELINE_STAGE_HOST_BIT, VK_PIPELINE_STAGE_VERTEX_SHADER_BIT, + 0, 0, nullptr, 1, &bufMemBarrier, 0, nullptr); + } +} + +void Renderer::recordCommandBuffer( + VkCommandBuffer cmd, + VkImage image, + VkImageView imageView, + VkExtent2D extent, + VkImageLayout oldLayout, + const Frame &frame, + const std::vector &vertices) const +{ + + { + VkImageMemoryBarrier2 toColor{ .sType = VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER_2 }; + toColor.srcStageMask = VK_PIPELINE_STAGE_2_COLOR_ATTACHMENT_OUTPUT_BIT; + toColor.dstStageMask = VK_PIPELINE_STAGE_2_COLOR_ATTACHMENT_OUTPUT_BIT; + toColor.dstAccessMask = VK_ACCESS_2_COLOR_ATTACHMENT_WRITE_BIT; + toColor.oldLayout = VK_IMAGE_LAYOUT_UNDEFINED; + toColor.newLayout = VK_IMAGE_LAYOUT_COLOR_ATTACHMENT_OPTIMAL; + toColor.image = image; + toColor.subresourceRange = { VK_IMAGE_ASPECT_COLOR_BIT, 0, 1, 0, 1 }; + + VkDependencyInfo dep{ + .sType = VK_STRUCTURE_TYPE_DEPENDENCY_INFO, + .imageMemoryBarrierCount = 1, + .pImageMemoryBarriers = &toColor + }; + vkCmdPipelineBarrier2(cmd, &dep); + } + + VkClearValue clearColor = {{{0.1f, 0.1f, 0.2f, 1.0f}}}; + VkRenderingAttachmentInfo colorAttach{ + .sType = VK_STRUCTURE_TYPE_RENDERING_ATTACHMENT_INFO, + .imageView = imageView, + .imageLayout = VK_IMAGE_LAYOUT_COLOR_ATTACHMENT_OPTIMAL, + .loadOp = VK_ATTACHMENT_LOAD_OP_CLEAR, + .storeOp = VK_ATTACHMENT_STORE_OP_STORE, + .clearValue = clearColor + }; + + VkRenderingInfo ri{ + .sType = VK_STRUCTURE_TYPE_RENDERING_INFO, + .renderArea = {{0,0}, extent}, + .layerCount = 1, + .colorAttachmentCount = 1, + .pColorAttachments = &colorAttach + }; + + upload_vertex_buffer(cmd, frame, vertices); + + vkCmdBeginRendering(cmd, &ri); + + VkViewport vp{0.0f, 0.0f, (float)extent.width, (float)extent.height, 0.0f, 1.0f}; + VkRect2D sc{{0, 0}, extent}; + vkCmdSetViewport(cmd, 0, 1, &vp); + vkCmdSetScissor(cmd, 0, 1, &sc); + + vkCmdBindDescriptorSets(cmd, VK_PIPELINE_BIND_POINT_GRAPHICS, pipelineLayout, 0, 1, &set, 0, nullptr); + VkDeviceSize vOffset{ 0 }; + vkCmdBindVertexBuffers(cmd, 0, 1, &frame.vertexBuffer.buffer, &vOffset); + + glm::mat4 projection = glm::ortho(0.0f, (float)window_width, 0.0f, (float)window_height, -1.0f, 1.0f); + + vkCmdPushConstants( + cmd, + pipelineLayout, + VK_SHADER_STAGE_VERTEX_BIT, + 0, + sizeof(glm::mat4), + &projection + ); + + PipelineType lastPipeline = PipelineType::None; // Track current state + // uint32_t vertexOffset = currentFrame * MAX_VERTICES_PER_BATCH; + uint32_t currentBatchVertices = 0; + + for (const auto & render_command : commands) { + // Only switch pipelines if we have to + if (render_command.pipeline != lastPipeline) { + // If we were mid-batch, draw what we have before switching + if (currentBatchVertices > 0) { + vkCmdDraw(cmd, currentBatchVertices, 1, 0, 0); + // vertexOffset += currentBatchVertices; + currentBatchVertices = 0; + } + + vkCmdBindPipeline(cmd, VK_PIPELINE_BIND_POINT_GRAPHICS, get_pipeline(render_command.pipeline)); + lastPipeline = render_command.pipeline; + } + + currentBatchVertices += 6; + } + + // Draw the final batch + if (currentBatchVertices > 0) { + vkCmdDraw(cmd, currentBatchVertices, 1, 0, 0); + } + + vkCmdEndRendering(cmd); + + // 3. Transition back to Present + { + VkImageMemoryBarrier2 toPresent{ + .sType = VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER_2, + .srcStageMask = VK_PIPELINE_STAGE_2_COLOR_ATTACHMENT_OUTPUT_BIT, + .srcAccessMask = VK_ACCESS_2_COLOR_ATTACHMENT_WRITE_BIT, + .dstStageMask = VK_PIPELINE_STAGE_COLOR_ATTACHMENT_OUTPUT_BIT, + .dstAccessMask = 0, + .oldLayout = VK_IMAGE_LAYOUT_COLOR_ATTACHMENT_OPTIMAL, + .newLayout = VK_IMAGE_LAYOUT_PRESENT_SRC_KHR, + .image = image, + .subresourceRange = { VK_IMAGE_ASPECT_COLOR_BIT, 0, 1, 0, 1 }, + }; + + VkDependencyInfo dep{ + .sType = VK_STRUCTURE_TYPE_DEPENDENCY_INFO, + .imageMemoryBarrierCount = 1, + .pImageMemoryBarriers = &toPresent + }; + + vkCmdPipelineBarrier2(cmd, &dep); + } + +} diff --git a/renderer/renderer.h b/renderer/vulkan/renderer.h similarity index 99% rename from renderer/renderer.h rename to renderer/vulkan/renderer.h index 7df6462..91c455a 100644 --- a/renderer/renderer.h +++ b/renderer/vulkan/renderer.h @@ -15,7 +15,7 @@ #include #include "glm/gtx/string_cast.hpp" #include -#include "sprite.h" +#include "../sprite.h" #include "texture.h" #include #include diff --git a/renderer/init.cpp b/renderer/vulkan/vulkan.cpp similarity index 92% rename from renderer/init.cpp rename to renderer/vulkan/vulkan.cpp index 6caa8e9..cfc45de 100644 --- a/renderer/init.cpp +++ b/renderer/vulkan/vulkan.cpp @@ -2,10 +2,68 @@ // Created by Vicente Ferrari Smith on 12.02.26. // +#define VOLK_IMPLEMENTATION +#include +#define VMA_IMPLEMENTATION +#include + #include "init.h" #include #include +#include "../graphics.h" + +VkInstance instance{}; +VkPhysicalDevice physicalDevice{}; +VkDevice device{}; +VkQueue graphics_queue{}; +uint32_t queueFamily{}; + +VkSurfaceKHR surface{}; +VkDebugUtilsMessengerEXT debugMessenger{}; + +VmaAllocator allocator{}; + +constexpr uint32_t MAX_FRAMES_IN_FLIGHT = 2; +constexpr uint32_t MAX_VERTICES_PER_BATCH = 65536; + +VkSwapchainKHR swapchain; +VkExtent2D swapchain_extent; +VkSurfaceFormatKHR swapchain_format{ + VK_FORMAT_B8G8R8A8_UNORM, + VK_COLOR_SPACE_SRGB_NONLINEAR_KHR +}; +std::vector renderFinished; + +std::vector images; +std::vector imageViews; +std::vector imageLayouts; + +void graphics_init() { + createInstance(window); + createSurface(window); + createDevice(); + + createSwapchain(window); + + slang::createGlobalSession(slangGlobalSession.writeRef()); + + Renderer renderer(window); + texture_manager.load("assets/boy.png", renderer); +} + +void graphics_deinit() { + vkDeviceWaitIdle(device); +} + +void begin_frame() { + +} + +void end_frame() { + +} + VKAPI_ATTR VkBool32 VKAPI_CALL debugCallback( VkDebugUtilsMessageSeverityFlagBitsEXT severity, VkDebugUtilsMessageTypeFlagsEXT type, diff --git a/renderer/vulkan/vulkan.h b/renderer/vulkan/vulkan.h new file mode 100644 index 0000000..bc280d3 --- /dev/null +++ b/renderer/vulkan/vulkan.h @@ -0,0 +1,25 @@ +// +// Created by Vicente Ferrari Smith on 12.02.26. +// + +#ifndef V_INIT_H +#define V_INIT_H + +#include +#include +#include +#include + +struct Device { + VkDevice device; +}; + +void createSwapchain(GLFWwindow* window); + +int createInstance(GLFWwindow* window); +void createSurface(GLFWwindow* window); +void pickPhysicalDevice(); +void createDevice(); + + +#endif //V_INIT_H diff --git a/renderer/webgpu/renderer.cpp b/renderer/webgpu/renderer.cpp new file mode 100644 index 0000000..e9348e2 --- /dev/null +++ b/renderer/webgpu/renderer.cpp @@ -0,0 +1,5 @@ +// +// Created by Vicente Ferrari Smith on 06.03.26. +// + +#include "renderer.h" diff --git a/renderer/webgpu/renderer.h b/renderer/webgpu/renderer.h new file mode 100644 index 0000000..be8cbcf --- /dev/null +++ b/renderer/webgpu/renderer.h @@ -0,0 +1,16 @@ +// +// Created by Vicente Ferrari Smith on 06.03.26. +// + +#ifndef V_RENDERER_H +#define V_RENDERER_H + + + +class renderer { + +}; + + + +#endif //V_RENDERER_H diff --git a/renderer/webgpu/utils_emscripten.cpp b/renderer/webgpu/utils_emscripten.cpp new file mode 100644 index 0000000..a6e37d9 --- /dev/null +++ b/renderer/webgpu/utils_emscripten.cpp @@ -0,0 +1,79 @@ +// +// Created by Vicente Ferrari Smith on 06.03.26. +// + +// Copyright 2025 The Dawn & Tint Authors +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// +// 1. Redistributions of source code must retain the above copyright notice, this +// list of conditions and the following disclaimer. +// +// 2. Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// 3. Neither the name of the copyright holder nor the names of its +// contributors may be used to endorse or promote products derived from +// this software without specific prior written permission. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE +// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL +// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR +// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER +// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, +// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + +#if !defined(EMSCRIPTEN) +#error "utils_emscripten.cpp: This file requires EMSCRIPTEN to be defined." +#endif // !defined(EMSCRIPTEN) + +#include + +#include "GLFW/glfw3.h" +#include "emscripten/emscripten.h" +#include "webgpu/webgpu_glfw.h" + +WGPU_GLFW_EXPORT WGPUSurface wgpuGlfwCreateSurfaceForWindow(const WGPUInstance instance, + GLFWwindow* window) { + wgpu::Surface s = wgpu::glfw::CreateSurfaceForWindow(instance, window); + return s.MoveToCHandle(); +} + +namespace wgpu::glfw { + +wgpu::Surface CreateSurfaceForWindow(const wgpu::Instance& instance, GLFWwindow* window) { + auto chainedDescriptor = SetupWindowAndGetSurfaceDescriptor(window); + + wgpu::SurfaceDescriptor descriptor; + descriptor.nextInChain = chainedDescriptor.get(); + wgpu::Surface surface = instance.CreateSurface(&descriptor); + + return surface; +} + +std::unique_ptr +SetupWindowAndGetSurfaceDescriptor(GLFWwindow* window) { + if (glfwGetWindowAttrib(window, GLFW_CLIENT_API) != GLFW_NO_API) { + emscripten_log(EM_LOG_ERROR, + "GL context was created on the window. Disable context creation by " + "setting the GLFW_CLIENT_API hint to GLFW_NO_API."); + return {nullptr, [](wgpu::ChainedStruct*) {}}; + } + + wgpu::EmscriptenSurfaceSourceCanvasHTMLSelector* desc = + new wgpu::EmscriptenSurfaceSourceCanvasHTMLSelector(); + // Map "!canvas" CSS selector to the canvas held in the Module.canvas object. + EM_ASM({self.specialHTMLTargets && (specialHTMLTargets["!canvas"] = Module.canvas)}); + desc->selector = "!canvas"; + return {desc, [](wgpu::ChainedStruct* desc) { + delete reinterpret_cast(desc); + }}; +} + +} // namespace wgpu::glfw diff --git a/renderer/webgpu/utils_emscripten.h b/renderer/webgpu/utils_emscripten.h new file mode 100644 index 0000000..a6e37d9 --- /dev/null +++ b/renderer/webgpu/utils_emscripten.h @@ -0,0 +1,79 @@ +// +// Created by Vicente Ferrari Smith on 06.03.26. +// + +// Copyright 2025 The Dawn & Tint Authors +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// +// 1. Redistributions of source code must retain the above copyright notice, this +// list of conditions and the following disclaimer. +// +// 2. Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// 3. Neither the name of the copyright holder nor the names of its +// contributors may be used to endorse or promote products derived from +// this software without specific prior written permission. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE +// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL +// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR +// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER +// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, +// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + +#if !defined(EMSCRIPTEN) +#error "utils_emscripten.cpp: This file requires EMSCRIPTEN to be defined." +#endif // !defined(EMSCRIPTEN) + +#include + +#include "GLFW/glfw3.h" +#include "emscripten/emscripten.h" +#include "webgpu/webgpu_glfw.h" + +WGPU_GLFW_EXPORT WGPUSurface wgpuGlfwCreateSurfaceForWindow(const WGPUInstance instance, + GLFWwindow* window) { + wgpu::Surface s = wgpu::glfw::CreateSurfaceForWindow(instance, window); + return s.MoveToCHandle(); +} + +namespace wgpu::glfw { + +wgpu::Surface CreateSurfaceForWindow(const wgpu::Instance& instance, GLFWwindow* window) { + auto chainedDescriptor = SetupWindowAndGetSurfaceDescriptor(window); + + wgpu::SurfaceDescriptor descriptor; + descriptor.nextInChain = chainedDescriptor.get(); + wgpu::Surface surface = instance.CreateSurface(&descriptor); + + return surface; +} + +std::unique_ptr +SetupWindowAndGetSurfaceDescriptor(GLFWwindow* window) { + if (glfwGetWindowAttrib(window, GLFW_CLIENT_API) != GLFW_NO_API) { + emscripten_log(EM_LOG_ERROR, + "GL context was created on the window. Disable context creation by " + "setting the GLFW_CLIENT_API hint to GLFW_NO_API."); + return {nullptr, [](wgpu::ChainedStruct*) {}}; + } + + wgpu::EmscriptenSurfaceSourceCanvasHTMLSelector* desc = + new wgpu::EmscriptenSurfaceSourceCanvasHTMLSelector(); + // Map "!canvas" CSS selector to the canvas held in the Module.canvas object. + EM_ASM({self.specialHTMLTargets && (specialHTMLTargets["!canvas"] = Module.canvas)}); + desc->selector = "!canvas"; + return {desc, [](wgpu::ChainedStruct* desc) { + delete reinterpret_cast(desc); + }}; +} + +} // namespace wgpu::glfw diff --git a/renderer/webgpu/webgpu.cpp b/renderer/webgpu/webgpu.cpp new file mode 100644 index 0000000..d7c02a4 --- /dev/null +++ b/renderer/webgpu/webgpu.cpp @@ -0,0 +1,5 @@ +// +// Created by Vicente Ferrari Smith on 06.03.26. +// + +#include "webgpu.h" diff --git a/renderer/webgpu/webgpu.h b/renderer/webgpu/webgpu.h new file mode 100644 index 0000000..733cee9 --- /dev/null +++ b/renderer/webgpu/webgpu.h @@ -0,0 +1,16 @@ +// +// Created by Vicente Ferrari Smith on 06.03.26. +// + +#ifndef V_WEBGPU_H +#define V_WEBGPU_H + + + +class webgpu { + +}; + + + +#endif //V_WEBGPU_H diff --git a/shaders/compute.slang b/shaders/compute.slang new file mode 100644 index 0000000..e69de29 diff --git a/shaders/compute.wgsl b/shaders/compute.wgsl new file mode 100644 index 0000000..87cd138 --- /dev/null +++ b/shaders/compute.wgsl @@ -0,0 +1,13 @@ +@binding(1) @group(0) var output_buffer_0 : array; + +@binding(0) @group(0) var input_buffer_0 : array; + +@compute +@workgroup_size(32, 1, 1) +fn main(@builtin(global_invocation_id) thread_id_0 : vec3) +{ + var _S1 : u32 = thread_id_0.x; + output_buffer_0[_S1] = 2.0f * input_buffer_0[_S1]; + return; +} + diff --git a/shaders/shader.air b/shaders/shader.air new file mode 100644 index 0000000000000000000000000000000000000000..3ce250fdd1a86edf18a4c9a07d7512d0c9e92fa1 GIT binary patch literal 10416 zcmdTqe|S?>mN!3|q-mO_P-(Ggd8ME!G|6jPN=sX4nuCGjI+AuzL!7? zIJ$rA_szU7@7;ILJ?Gr>N`0>_I-d)0mT^S7{Dnm|{ zEN`{k(AwMN8!0MooE7a;8FvY7W3q!eglT0j@f(4mwI^?pW>Ku{9|kr(MU0LC-<KIpL%EA@vHg!b{_fa*IzCC@@I>_ z%)K;k|L$zDKuAVpPsGJzYd)-KTs0BNKV+(&FjWmi3eOvBdPH?&re&v1@}C&R zO=RP~s_a!wtJXKB@BaF}ktOT0|M4d=OGmd=U?$?kYw+RKM5XDd=r({TB4kzgqkS z_KAQ@OJt!c^c9LB`FYWN^=U!j)R}@ydo#!P7M`TkClcz@doxcIsm}`v(AHRVt410k zrCqXUP!p{TMLU|L9q}j=lM?F5Z~@v^aD1@fM5y4{8TC-O;6$SE7^OacreJDs=4V9( zW8uOHL9eJfA+q?Q9S%|Z{;1y~-3lD~r6CN7_vwP<@h+5w&d zX-O#ynmQs~O;Q*YC=2a_QC&DRs6KWpD^=!61f})=B3SknnE8k31Zoi_a?lF2ATIo~ zB@z}Vgd6hv8qu8iiG&7wVg+)e0|G*D+k_NiY8Av*8kDnp>+^ck8>lsT$_3AnnkJmn zW*_l2NxKF_%f3k42G(>$d?9J41^fX6J=zX~H6g7Y5G@VCFOvDjI;9(BZM8`KacQt# zieDxqe^5!chVaiFQ+CFs_{RZ*acMh*kxCkbDk){+QkIZ{-V9V{ayU1$_1w(XtH5Cz ztE&X={Ti#QGik5ZV*#$ocJ&5UF^(01aoXWIApzzL2tV!Vwq&)(u(V7wM#fJ-BcuV) zW!e+a0zeg)UVn1je4KdgNXpb>wzxVW`HbrVxPjRBWXZC386WM-nJcJh=zrmv!Izim z-~GM*t*)=;Kl7>n%PmhTcfWi2Ph7!#|E)DyuXi+|@=}UaAtVNW2}BoT{l94$u@F#! zODM^^fBWNNR1fb2+z2I*Y0(l#X%d|S0nWbGGBhFy4;_KY?Lh+F2_Qou z%rx0+Vtwg%+&wBbEn11>-FxP1q79nJEX_}aIhtPKT!AJoT%17i(OLB>@fkBRs(OBP zP>jSz)z{NNKcvK#QUL(J!dPDJ^=o+jBmc@l&3UVVdSpdNoO2)}t<)6$dGkc&?Gu-d zyn1%IVRY905B3;U;;;WhBvY@R5Tk=IexC&y%Vi<;$QiXgM^M*ms_Gf;YvAIF?Qd6{ zLs_4ylvhH^azz6~5rm@JPNs2ak^FHyiPiOq>c%_!AjNExNg=#_kZXLRU}w}vMCXoY zD=)^BXUHrMnRWJ(@=~+1Ts2Z~3i1w@2kqp}nzti06Q=4zqT073wUGG6FXaR+(rpgu zQk=Y`OG8l~Br=OMsF4Q0NCUPVp{N-mC@;$NN4xsNrJr3=es)^8|-uR5LKF__Z_YYl$+x zh~I%j)t?vj$D?!KU!V0E%DUXFoQf&W)@Pj|lovq@LQ`;ZuyDv&cpUOwxbWm@HB6!x z1cmmzbyWkVWiLc(#+w%;6|qD+I;EZc>(*R}Wnl{xKFm#%MFpn@3m$>l7X&UUxI{@d z?1)s257!Ks>cAj>8LFK!)$~GjxbLVOG>vwEJTOPXgx4S5bXKLjm{5Mwto)oT9)y`@ z>P+tBnSxUwK-1isfk@q1*zJ(>^b*$qsohx ztWQr%E)iLukXfG)Sy!8t9+mPesf6n1rzL+lt(hYGh!TX@>R*2E#r^in-l#|i9sZO(v4mz#0xiw(71G^N| z>h-XG%HzKg?sqUWWou)+0pCqcbf-VS(AQXCL~{q@^wVbl4NWZN_PL-p{r6b>4BAC|tqYtZ4H11T{G0n7v;Bm z88@q6OIhz*-=gy+0eHM6aUhaX&4x;yLC4y*@W`nejtQTab>jZ;JH@ndqL@7pPIi}< z@>it57y_uc3Ih=ajM*FT`vU$p+5-{ir)?TTAlMi*nC>A=IunEwOrCUZycv0j=EY zG&W9K0HT)XxvCL2@Bxe8TwdN5W~awScVY+*JzP>Y8w4b4W_^^E#sJpO1gw5udMD%v z)=If}6)2a_L2)!G?Qq7J0iayYHV>CblDhc}X4YU1fI8L+GgSr*Ib(EIcp1OLi)k4e zL}JW-Fd0M#0Ic93KKFRMX7CW}qpf~t@LER!Pk^O4*&$iAxd9ur0s~^&8_<(U1-J0| zBndFsYB@Pd2_C$hv*iXp1r;a5iKv|(K+EOV;Yr@uwywZ@@b|< zF6T{&lPvU0Y?Bjk@jfMB1k_Bg2ZA;7?Ld-BswpDe735xaW94R|F3DosHZ7@JQc4i; z++Ip-BXwjcQKloeZ!hIorrI&qwRS+d-i|4+(*x->`XqyIM!-6J*A7wwa2+HCz)H@r z%Y6a=bkkH|(-0dsrlM){=9}zv&5RWU=QrMRyi<~P!%UN2azm!oaBAW3NJ3XNLKjH{ zr)~3Q-f-89pKQ5d@Gc^po*NKx>RGldHKvJ0+^8;*^R9k?_j?c=hZ%BimwT*UISuW5ZIR zjfOx4TN5`#;d0)M$?yZ*WV?ne*}SETo zxtt!FVoC|O({lq1Ou*@)*AqM?;djs}UWDIEcv~TpcnFFYio@mlPHD**BuQE~L-J&f z(ImxB3RVjC1T!a1uh%CN0L!$sVwcwuEbW2>rzz7Jstw>}4x1llAmcY&5K&8(0J6{9 zq2UQ|0?-?D6(xig`+LcfZ)3rc5z8?$SzcC7RwX4&I(eF!{0PzOlP$?nm{cG+x_}!; zp`(D!cKJ;UGCs@0bjz=0_%JZ@8C0*AgQg~!kzn>CTDyn`AXH!uAXd^0LorTbz1!(` z5U^r1EY*<)9Zb4FAeR^WaQ>&D!Aeid&fIbUKW>?ual`c#l$7I!K%btE3^%1}F3@6*BHl8oq=a3^K?s7IIr%h|%zYae%ERJi2e(9NgqdV|rxl$Qgd5f@E(j&O8k^P| zZVEc-j$5}(oBysMK*r^!Rj=Cz3sIeo_Tbfy+2&-aRu?y`!eW+dvQw-dlR3<^=cYEi zG2j;zW;^Y-PRC*S5Hv%K!xH&#W~toj^zg%t&}wb4ey15NHfkL=1PQI3p{-uVre(bW z24>j0E}A8@Z8&pi{aAZe>-D(02<utTC?Swf;6r2KG6OpRA@@X)A4C+S+RWHVVJP$=>TV;>Etv?*~u+bPkL`va+X6#7Icn?54CSN_4B%wDSjOx16aW9)H$v-nQ&unAnG!)7nIRnraCm=Lv~7ElPn; znFtw$C?u{&NU;IUPD7s;p4!?{}r za~~0?hm7jcqTI^Lw<>cD2g{#5_{(*|O-&0+*S>%B)iqDPs_^bX^LC}L-p!Wn&Zyb4 zn0+ySA=-HW^L*#R`Qwx4Hhp$#>Z9W)-yb_ZwKBb7{z71M_!rB5S#h}f**D)<_U6IY z-drc#HC6)r=!>sp>YjMz;(um7sr&Rjw`b4Z1#3QDaK9XQKbNs3VqLuZ%OAa%@#5m= z^B30d+`vhFeC*iC^DEQupI?OKd-mWPhYyoW4Tme9M4c9?WfzW#KN&iOOJW zUe2gl{N-zlcR#;+8=Jq-vh$^(<5Qc?ow_hJdhDI?gbncRVKkGX`rp9c0MH-d+sW^soPYONY+Yq@CZ>P) zbN@3ta^(dr?IDQ=>Nr!wkrCcudEFkT4<=!l&G7NbwazmpQ&R#TSw}cu>{&v=L-+J`PulAA)UwpWS*FPIKCjsUlY_q(K8D>-) zKHBP)d|9Pf;(Lj zQ6qd>V+802Fg00j3ryX0FX|HP6R+vf2rQ_#n~;4XL~p}Kh{U1`V;Mh2mm3u6{dp$g zj9$=>rL)FV^Gz!o{!uV`ANr9ZeaCB68QOi~j0fX8zW3o%xix~RwBA`{_C9e)K*+m^ zJsIfIERbHiuSWl;jA{oB(k#5~1AK!Vtb@=q`6~^3r;lW8&-_f`bf0;x- z%UNP=4u$hVX~@<*(q|M%R4u4}UhmT8$m!=vq>+vF2qvWSm6H}Sk}e8$ht}kw^xMj; z&3%Kn1I-nCsJo)SOf!nI5<(t36fXLzTyiw;CHV70z4y_bRgt-oR>dYkDGGJRfx1e3 zdIct*2ju?V{TJ#pFIYz7O-DmTj~_ys`eS{@qJ850J!uc2h{9CyUS7`pdwLrljHiXd z1n}L_J1gEi$06|dj$KY<4$O_rzfJUL>7+1@O1MT>e9jhu!k~T{)aum~$sayA2g$=< z{3~V?>VD`zTIepsvw0X}`gxJg zEFsE3(A#WD!X;XY$|Y&Hi(xtOdgkD5pd1tg6!~xq%CM(5icjBpSYa~eZhUm#{>Qb_ zS>s2k+_BBB>SruEx|?+f zmNE#PuZ`6|s+wz)&HDwbqscUva*hO>#I^)lg7YxCQJ^|LDleRqfx zJ%WKj7@5+vwWz)xCM!&^QEa&b)8VU~;Q!lZP^#wl2*1jzQfcw@g=oNC3gs-(o&gNN z|9`t5Zam+ufg8hG;Kmm&;qxe(z!kWD0Jj?Mt#IS{AJgC~)sEjWKECF`jsJ9)D)B#5 z&;o=O0W1^lRGETfneyRIl^8Aqe0-()!6k>U;$gFsZMfv%P~vMdATEXybXV)6tr, i32(4)>, +}; + +struct Uniforms_std140_0 +{ + @align(16) proj_0 : _MatrixStorage_float4x4_ColMajorstd140_0, +}; + +@binding(1) @group(0) var uniforms_0 : Uniforms_std140_0; +@binding(0) @group(0) var colorTexture_0 : texture_2d; + +@binding(2) @group(0) var samplerState_0 : sampler; + +const square_0 : array, i32(6)> = array, i32(6)>( vec2(-0.5f, -0.5f), vec2(-0.5f, 0.5f), vec2(0.5f, -0.5f), vec2(0.5f, -0.5f), vec2(-0.5f, 0.5f), vec2(0.5f, 0.5f) ); +struct VSOutput_0 +{ + @builtin(position) pos_0 : vec4, + uv_0 : vec2, + color_0 : vec4, + alpha_0 : f32, +}; + +struct vertexInput_0 +{ + @location(0) pos_1 : vec2, + @location(1) scale_0 : vec2, + @location(2) uv_1 : vec2, + @location(3) color_1 : vec4, + @location(4) alpha_1 : f32, +}; + +@vertex +fn vs_main( _S1 : vertexInput_0, @builtin(vertex_index) vertex_index_0 : u32) -> VSOutput_0 +{ + var output_0 : VSOutput_0; + output_0.pos_0 = (((vec4(square_0[vertex_index_0 % u32(6)] * _S1.scale_0 + _S1.pos_1, 0.0f, 1.0f)) * (mat4x4(uniforms_0.proj_0.data_0[i32(0)][i32(0)], uniforms_0.proj_0.data_0[i32(1)][i32(0)], uniforms_0.proj_0.data_0[i32(2)][i32(0)], uniforms_0.proj_0.data_0[i32(3)][i32(0)], uniforms_0.proj_0.data_0[i32(0)][i32(1)], uniforms_0.proj_0.data_0[i32(1)][i32(1)], uniforms_0.proj_0.data_0[i32(2)][i32(1)], uniforms_0.proj_0.data_0[i32(3)][i32(1)], uniforms_0.proj_0.data_0[i32(0)][i32(2)], uniforms_0.proj_0.data_0[i32(1)][i32(2)], uniforms_0.proj_0.data_0[i32(2)][i32(2)], uniforms_0.proj_0.data_0[i32(3)][i32(2)], uniforms_0.proj_0.data_0[i32(0)][i32(3)], uniforms_0.proj_0.data_0[i32(1)][i32(3)], uniforms_0.proj_0.data_0[i32(2)][i32(3)], uniforms_0.proj_0.data_0[i32(3)][i32(3)])))); + output_0.uv_0 = _S1.uv_1; + output_0.color_0 = _S1.color_1; + output_0.alpha_0 = _S1.alpha_1; + return output_0; +} + +struct pixelOutput_0 +{ + @location(0) output_1 : vec4, +}; + +struct pixelInput_0 +{ + @location(0) uv_2 : vec2, + @location(1) color_2 : vec4, + @location(2) alpha_2 : f32, +}; + +@fragment +fn fs_main( _S2 : pixelInput_0, @builtin(position) pos_2 : vec4) -> pixelOutput_0 +{ + var _S3 : pixelOutput_0 = pixelOutput_0( (textureSample((colorTexture_0), (samplerState_0), (_S2.uv_2))) * vec4(_S2.color_2.xyz, _S2.alpha_2) ); + return _S3; +} + diff --git a/shaders/shaders.metallib b/shaders/shaders.metallib new file mode 100644 index 0000000000000000000000000000000000000000..ac97b15eab13b3c75e266fb729378718930855bf GIT binary patch literal 21840 zcmeHv3s_Ufw*Q2L?0_MJhZ-Qk4JZcW5dtcRXhML3LXCnJtThQqFi{902~V{R0ZeSM z21_lrwha$kU;h@Zt+utN37|!c7GFKqS_>BIqiC@{jy*?@cP3Gbw&(WT|M$Q5fBSve zU-q7vHEY(aSu?ZtnjiZkCoemdpmzfv0MrHo0E7$S177*g1mFZ`X5)jD{(%0@gL2P< z@^Uw4#v0*FQE@%IpXTSJQ$dAXqmx&ON+e1Z$SV-0!^WBX+)NxyK0G*ZYoGroL$SAB zp4|S`FOhv)g|lxj2o$DxO!5Efe0EM+?qUi6fjq4HRi4%Sz-b-DrgfPKCa_+I#$YA@`8YaQn__W&OuD7PJ?P z-SSe_^j`uiKX3nz+IjD5LD1{6+TEGs{2O@W#%qOql3sOXS47JmP4wd0gO(U z!g5-AYE~o&g1Oj*TDe9WTT!HqO_Nu|R>W!*5*ea6)0h&uPEri{g&7n7kXNGAIC}cE zhsgh?9?n_EkAb3rKfDLV0dOB)sRHY{sgOW~v_dLpqMq*Fi|RwD0O;HS7~>WKClmn4 z4zSv=CVnp|i%jPxZ6MTBC%IDsf=MqG>iyH*Qa6lA?Vxl06Df?jp4xRImwOF z62@yLNG9j_Ov-*fJiw+5e;Fwa<-1Fc?XLYS8bEKp-M z*03=>2dkzSE7-{UR3gT)uR$XPRRuHv% z<1~d)ZNYTl50{jAcS>xkokltQ?Lq6GKQKqWBV-kPjAs?IFzL{NOV@<`o;_hXz3- zf<_BrbxDHtUaUAUhUdwfS}&N`1rI}+e)45HR-;G15;LvJf}x*-SYg4+;PlY33RvV|8Vjbyu^O0x#m6Q5@ObSH zkJsLY5=M!-jicTliMsuG+uJ;(z$4vmKOxHO5QQV2GWhH#L*W{z{HSM%*t^V(wABa^ zGkQWc94m)o8TEv0<*=0n`}OINrgo@T1~3f|iABYJ@-;^nP#RQwjyJ`}m%pabKY);s z%Ba|8g4_9Rg1N{@;&h+)f?Eu}H+WVch%BlG@Pz^#uM_|nyuv!3j=?F0k?9+O z9;clDZWJlD5j0MM#I6a5YaPm-N7;WmvQFw4sRwL*2Q%{|3;egIhs^Phd8hW^(Z$Q2 z){X3$8?U5b)tN@M8mmzgE2@lY+!$c@p$(g9gB))a$NT#c z+K7-APw!!!3wf{yia2YN_tz!&3zGH_XYQ|?2^YP6#J@s}tyEys6>vdg)AdHR*{Bv{ z6>O~Hjyn`vsW*z?DrXoqEyn7W+Nf(Iv}*&jTUOdIhgjH)wwZIYF`*;t3t{%uvicElRLYqUnq_0e4|EPr&FNqq+=L$7GuEC zT<>cDx?kFmnFc-ZAWj>Gqrlm$v$o(ae(>=S=Gofdvja@%K(7#jix~40I|VZ~*CpG9 z6CVsjY^SA*>MzIj^BE(cYchDM*0^x!W~! zNRVuT#qP!fo^YhbN~jO?kI>n*)Gqm+P8+t;E(mEiIAJ#EWrhaFpB`kLg9>B@BzM-O z?5PW??@nwKB-y)X?z1KJcEd5wv}f)qqhL_0GHAejW572DX~Sags{`Z_-1`E@`vUHL zTS!yUY2R~buy|vDe0hL2B&6NpgqB#irrACyrf=tp<8URvM^R!EKx0UmXZ3gN|m&@Tqbw&%|xTtMJrXaoce2ZGT7$# zV#z11=HxGt=rqczKhc0=waOA`bcfo`*@{ubqWM-zB|13I+Sq()sZ5ThD^^gd(!mpg zBf^#!m#XEOSYfe5S5&Gg(Z=RTr1Ns~W7H1_C<=eTlbbg8VKsMJ3^ztATj6AWSb_}r zBImM7tCVPp9u?5q!?RRHr4K2UaOyQ8t+GZAk1HkeXmzO)s-%GddrXO>N+gq4D5Y{y zVR=yz6!(xSg`?0(;K}9*x=x{yOJvJ5rRD11RLQG!DlV1i5+0Hf1+t>zgF|L%xlUcKTP9aQTj}I7Hn+SYt%55m zif}Z9rC+P2N=u7NHE9(Iu-1_sB}OSLhxT^L%t}KQPDN3!xYW5|L`s!RUWF*gXegD) zWYCmakyb5{I?SNeY09NKr}j!XyIQHF*f|7Av05QG}n(2B*n^QD#sFdFg7(; zq~(gr;c%o0B1s$+=rBfALa9clCbio5Y4kg%Oc&Z1SBRqFH1#rY>7!Cn=AR~^| zhXCs;ir!*~Bv~T?kqm(YObB}v5^xV-Hv&rJo>Z{U&C$f-Y+`kDHgSv6@uEOHb?1oZ zZceuudpX$78Mbv}@ziQ?%SNr6khY4`-H9!5kZ(frO5QyiwL3WP1wS~;Ha^neMCU2HpD5l{MtN|-J z{GP(h>>ha$=)}>l`K@H^LVY|hX&$kBZ;H%Q3 zuVLMh_l`7V58LHCRm}tb+(PLGE40d$VLPN#v~kig$!`S7H@Pp~>i+CL_a(jV&z*K( zdZ{7r^y%R0S9&{ep&`Aph87NK~m+-xqc2Nk!FYh4~ zgn3^S(r$L@Yq4s!BRZ@y!pH{2M)~T^-HCe$$sKhm&E0M*V4%r_2fnj7b4Y@omuiE% zB+UL+X0J1#k=O{{@Dx2L^oB7mR0$k|E1C>b%s#WDE_pKy5{W66x|uLwGswI-s|zL|3YL_Br%|fU$(ALJk8OjYzF2l_9T*>^0Nj z;-c^b$Bz?*$3@}XC_H}F!`6|bIcEJ=CH=b%{$0K(9H}@9t^>dN%bW1*{>UrphDxwR z^3qywE18b8O51podY}f4G?iUFY|0xMxL=Cy_eY5k$8JezO{pQ32c(^bQa6lTL zUN0m7F?7Y^vE@TODIkf^mX=-5CWwK_fKytbUxi0iYvh$)k5xdEN?qH+5V$?o3wRUr z?WKnbGE&zQGG>F((R(o2~w^$Sln`TYWhNG9JnpEi=7fO{J7dQYf1)IPec z-T+y~%{N?_xS$L5uo`IE($=|)NH)(MKX34gB*pGLjg@13wK1g!PqXHp~eHBce&E_13> zsIO(*Ap#jkde8!hy%T`NyU=SNpoUX{r`skD_&a{a8gx|L1uVGPgNU~V#FGy&cGk`Y zyl061TWyW}IC9m1H{kJcKQbV%x$`B%Uuq`*6M*Nu{A$UFd7oi60Pd~X43J0~m@O4H z+9tpbZ)W5z>x6F*sj#EaYU;@X{>@ye(8RWp!z8{|yqlgAPi`>zwsCyJ31ZxIzmXr( zMq)SCXMxXrz7E`PgPNm;^>co0$ZkRn=NC8|>U|>!7W%`6mb5t_-_gB^-Aaa>ywZ9M z({&zU_-|uyLa+M*_bqmY|KvHmDTBU%MnL13JSPzrc+b=^vR6Yrn`0m~YudK1jmCWKJrcj(oZ_U7 z;&9pj|w0Kyh4hi0h*FgfFp~jeI8p!w|L)$silvh?Wm;bfFZkd|*Nac!j~632(%U zfMk@S>tQ|0BYt#^=0Us&BLx#3Rgc;s9tRyceRw?w325a*-CjUt(cNtsUi= zM<4v4lb)6!s^BI@J5G4fB@$_=whD&&oEZ3`c;Xt3e_NozQQ!X={=E=x?f)tKI|SM5W!FCnZ^Nw=la8NlOxQ8u6JK1@Dfq6*Ox4axGOhkmM|o(XAaL!+>3BVm1oSY^gEx#N+}7lYJJWd%T-)g}{qS9U%g9&) z;EkhrK^BN>(?{XjHBaK&JV{i@tLY39f8>6O{!+dHxqTP*?82V^H)GG=!{FkNuxB<_ ztZs*V{~CKvYqGI@;}ehHWEd7kZ7{_7xNv6|?)(Qpy9;;z6XDJpImC1|A{_)=;NM^n zr27z_eH?rqjZaGbYG~CE$52*00&_=!)JR&ns$45~z{QU~L{h4LAG1aZ zj1u^7Lf8**d(`;1P_+}pf55_(Jw#8OyqxKIVe-Zqvf=&*8cn1X#2wxt7 zz;biEVT2w=nHOgmVT4h>gb5ByT*C+(mS6&bc=Jiyhls4dZbQH_2V2>X#nU&jZt^FL0+w&_Csch$<;U91HwldT zT1M4(3^PB#ZXT!&*vw~GYF{HuE#!E%R!m+&Co2Y4ORZ!%r%r>F43JmrvC;vXjnQMH z_QdMp->SPI+)4MDwbDAQjXwD0j>KK$}xqx>=f z^S*8M4kbV=`nS-q_wPZ&1@4O!?$1=Y=bPLMS|CjPI5b@R1T_5Yzk-Ix|2{N~@a6^? zf`;uwT}RJ2;O2uNTOi#0+Q*&K52#q$PrdN57Hf2F)a;2!(g3)@%YDwe6a+8#5odsFR`S(;2)pUo6a$?NKr(D}X{MRCh91E6 zJ>#{wpU_D-okoH5I?n-w^gzsqe1vWg(+e{6{=^W9y^(*N+&E(=k@GqoKwt|3b`4Db z%Mdm~7#2FoJWB#B)DB%`_-G93n&TS;L`N7D`rZ%XJYuMv;~a}rL3mI$yb*0Q28GVK z7I>rUDR?8z5Ke^7Nhn2^#qo4L5I?#`^B`Vyibfa|4)f?M&W9S9Gw9KRqu zx8MsKu?LWD-T&mAnFQv%aQ}Yn>@WM^gkAvz=qh+{dr-t=k5Vj zZ&BBgDLrC0Nynk)WAnn_;ntlDB((G3%2dEz7U$pp`ubeK$zan>JxBoP<8Cn6`bv%$ zHzA?7m&;Ay68$oAfULKN(-fKzmTkzHcc^K5&D+gyZ%@fd;c>mZsCE2y%B@-V4{u_9 z*Hxc>=`gWaY5^<0J+ZgMbg$p$pD7ijX1fhey0ZKISo*hrnIh@ca_UZv%Wgj1b1>rC z=A;h}XZ_=~RIl4@>-`o71}28oN01C76uNif%{$sJr>4w^-Sdp!-u%7|b4oYH#O~O= zy!@4fb&B^_?4A0_tkrM7e4S-mw(e=^Uv!F&g&iiJ*{sRFlJ>33v#;g6oSYjoB+(pu zP42$$PAKElfYs;hj(CycyHAvbU9I6E_T%kaaw&Q7VA+`Jp9&W9M3tVKMrew2ubtg- zq-v4;yU&yFPM^_&kDWABdErP&Nr?X9YG5F2O$GL6Cj-v|J6`X-JDGl}Z)YHLUP-3q zSU~D7tU$0Tc3Xq0s&mnYGn!w_|K-p@g1>|{oRT>5UU}1&+)H1WLPJAkM@lx~2A`t+ zIVut};D<=_-(FSrvzx2GKRmIo!smo1Vaz`)w-YN3tskGhyA{~nKl1qe%8uJlT|HIu zQOD6k`)5>;4BH2z6Hji`$pyZmc~WB2v;)RZi9tcKqd&7f!qxGIwS^6|r8OM1@e1{M|^do6>} zXByV_`v%GP>=^mVIBwzfW=|?v#dtrpLBWm4@Hu733Le=!^7_VeC$F-u-|Jw#6u)t` z>8Dk5#vb!bUrFB+(Sq?r{Zp5`;@Muff8&w5{o6}P!^iv?E7VNSYg4|t!aA_1?_A7y zul_HW_rIzuT#-hkOF~~*^6F3QGw~-r{pP8hlm#)-n-pL!-RPCFQX8mz`F!k1y zn#TNd5#-Dl{X=(kuh{$JI!n{rcw^i3+2M@E?~ugTc>C|Y`^|bRAz69!!tmKS9~3;V z9CuidK>g)_c9MwKi=l1qC(#IUvhwATJ4lZk5IL^?$sQ!Fd?~&%W{0iDm%xhh;^s`UDupKSM795RhxUl;Ars8+w zc2`_s^#?{bHk*7u{z)~k^snzWeYjy>+pCMVlu2J&p4&LBJbkZt;9v>qr4@m--gAj1 z1on1(rug%flR}?~_>42YgEGaQ^w;;L>6@-+ja{w@sfjxzk1tpHuMKbgC_3A?{!}RM z+@d$PQxHR;D`9qSGI<#TjP zpD7!;v~7(4bE~GD-Dzm8IWr85>D5~wQt(8)%74il%*9=A#Dpi_|3B8I7SravBue}t zA5QwWw6(j~e6gU3S+u_}t@fqcOX_3|jrO#KGdL5e>`$(|S*)sBw;*rK z$A{V!Me+|tOU2fNA6EXdXzbsf(MMfXUy8rq`*YnV6BlPlO_4Ga7TWt(Q&Qr%lIe+w zYnJAvtX$aiTF7SOr2O~8LwD+TG vWdWB3To!Oyz-0lK1zZ+zS-@ohmjzrFa9O}*0ha|_7I0bMzhnXQA3Og8MKwI< literal 0 HcmV?d00001