From b70317d85da341a0a1f90457d7cf2bb08035a9c9 Mon Sep 17 00:00:00 2001 From: Kbz-8 Date: Mon, 17 Mar 2025 22:41:04 +0100 Subject: [PATCH] yes --- Examples/WebGPU/main.c | 2 +- Includes/Pulse.h | 2 +- Sources/Backends/Software/Soft.h | 4 + Sources/Backends/Vulkan/VulkanLoader.c | 7 +- Sources/Backends/WebGPU/WebGPU.c | 197 ++++++++++++++++++ Sources/Backends/WebGPU/WebGPU.h | 9 + Sources/Backends/WebGPU/WebGPUBuffer.c | 47 +++++ Sources/Backends/WebGPU/WebGPUComputePass.c | 7 + Sources/Backends/WebGPU/WebGPUDevice.c | 8 +- Sources/Backends/WebGPU/WebGPUImage.c | 69 ++++++ Sources/Backends/WebGPU/WebGPUImage.h | 1 + Tests/Backend.c | 72 +++++++ Tests/{Vulkan => }/Buffer.c | 35 ++-- Tests/{Vulkan => }/Common.c | 27 ++- Tests/{Vulkan => }/Common.h | 0 Tests/{Vulkan => }/Device.c | 12 +- Tests/{Vulkan => }/Image.c | 0 Tests/{Vulkan => }/Pipeline.c | 52 +++-- Tests/Shaders/.gitignore | 2 + .../Vulkan}/BufferCopy.nzsl | 2 +- .../Vulkan}/ReadOnlyBindings.nzsl | 2 +- .../Vulkan}/ReadWriteBindings.nzsl | 2 +- .../Shaders => Shaders/Vulkan}/Simple.nzsl | 0 .../Vulkan}/SimpleBufferWrite.nzsl | 2 +- .../Vulkan}/WriteOnlyBindings.nzsl | 4 +- Tests/Shaders/WebGPU/BufferCopy.wgsl | 8 + Tests/Shaders/WebGPU/ReadOnlyBindings.wgsl | 7 + Tests/Shaders/WebGPU/ReadWriteBindings.wgsl | 9 + Tests/Shaders/WebGPU/Simple.wgsl | 4 + Tests/Shaders/WebGPU/SimpleBufferWrite.wgsl | 7 + Tests/Shaders/WebGPU/WriteOnlyBindings.wgsl | 7 + Tests/Vulkan/Backend.c | 52 ----- Tests/Vulkan/Shaders/.gitignore | 1 - Tests/Vulkan/xmake.lua | 114 ---------- Tests/{Vulkan => }/main.c | 0 Tests/xmake.lua | 148 +++++++++++++ xmake.lua | 5 +- 37 files changed, 698 insertions(+), 229 deletions(-) create mode 100644 Tests/Backend.c rename Tests/{Vulkan => }/Buffer.c (96%) rename Tests/{Vulkan => }/Common.c (78%) rename Tests/{Vulkan => }/Common.h (100%) rename Tests/{Vulkan => }/Device.c (78%) rename Tests/{Vulkan => }/Image.c (100%) rename Tests/{Vulkan => }/Pipeline.c (88%) create mode 100644 Tests/Shaders/.gitignore rename Tests/{Vulkan/Shaders => Shaders/Vulkan}/BufferCopy.nzsl (95%) rename Tests/{Vulkan/Shaders => Shaders/Vulkan}/ReadOnlyBindings.nzsl (93%) rename Tests/{Vulkan/Shaders => Shaders/Vulkan}/ReadWriteBindings.nzsl (95%) rename Tests/{Vulkan/Shaders => Shaders/Vulkan}/Simple.nzsl (100%) rename Tests/{Vulkan/Shaders => Shaders/Vulkan}/SimpleBufferWrite.nzsl (93%) rename Tests/{Vulkan/Shaders => Shaders/Vulkan}/WriteOnlyBindings.nzsl (74%) create mode 100644 Tests/Shaders/WebGPU/BufferCopy.wgsl create mode 100644 Tests/Shaders/WebGPU/ReadOnlyBindings.wgsl create mode 100644 Tests/Shaders/WebGPU/ReadWriteBindings.wgsl create mode 100644 Tests/Shaders/WebGPU/Simple.wgsl create mode 100644 Tests/Shaders/WebGPU/SimpleBufferWrite.wgsl create mode 100644 Tests/Shaders/WebGPU/WriteOnlyBindings.wgsl delete mode 100644 Tests/Vulkan/Backend.c delete mode 100644 Tests/Vulkan/Shaders/.gitignore delete mode 100644 Tests/Vulkan/xmake.lua rename Tests/{Vulkan => }/main.c (100%) create mode 100644 Tests/xmake.lua diff --git a/Examples/WebGPU/main.c b/Examples/WebGPU/main.c index 45bb3a1..b4a4b65 100644 --- a/Examples/WebGPU/main.c +++ b/Examples/WebGPU/main.c @@ -27,7 +27,7 @@ const char* wgsl_source = WGSL_SOURCE( @compute @workgroup_size(16, 16, 1) fn main(@builtin(global_invocation_id) grid: vec3) { - ssbo[grid.x * grid.y] = i32(grid.x * grid.y); + ssbo[grid.x * grid.y] = i32(grid.x * grid.y); } ); diff --git a/Includes/Pulse.h b/Includes/Pulse.h index 8b8d36b..8f7c101 100644 --- a/Includes/Pulse.h +++ b/Includes/Pulse.h @@ -49,7 +49,7 @@ typedef enum PulseBufferUsageBits PULSE_BUFFER_USAGE_TRANSFER_DOWNLOAD = PULSE_BIT(2), PULSE_BUFFER_USAGE_STORAGE_READ = PULSE_BIT(3), PULSE_BUFFER_USAGE_STORAGE_WRITE = PULSE_BIT(4), -} PulseShaderFormatBits; +} PulseBufferUsageBits; typedef PulseFlags PulseBufferUsageFlags; /** diff --git a/Sources/Backends/Software/Soft.h b/Sources/Backends/Software/Soft.h index 4735da3..a3a9577 100644 --- a/Sources/Backends/Software/Soft.h +++ b/Sources/Backends/Software/Soft.h @@ -9,6 +9,10 @@ #ifndef PULSE_SOFTWARE_H_ #define PULSE_SOFTWARE_H_ +#ifdef __STDC_NO_ATOMICS__ + #error "Atomic support is not present" +#endif + #define SOFT_RETRIEVE_DRIVER_DATA_AS(handle, cast) ((cast)handle->driver_data) PulseBackendFlags SoftCheckSupport(PulseBackendFlags candidates, PulseShaderFormatsFlags shader_formats_used); // Return PULSE_BACKEND_SOFTWARE in case of success and PULSE_BACKEND_INVALID otherwise diff --git a/Sources/Backends/Vulkan/VulkanLoader.c b/Sources/Backends/Vulkan/VulkanLoader.c index 476ae05..e7a4c01 100644 --- a/Sources/Backends/Vulkan/VulkanLoader.c +++ b/Sources/Backends/Vulkan/VulkanLoader.c @@ -28,12 +28,7 @@ #endif #endif -#ifdef PULSE_PLAT_WINDOWS - __declspec(dllimport) HMODULE __stdcall LoadLibraryA(LPCSTR); - __declspec(dllimport) FARPROC __stdcall GetProcAddress(HMODULE, LPCSTR); - __declspec(dllimport) int __stdcall FreeLibrary(HMODULE); - typedef HMODULE LibModule; -#else +#ifndef PULSE_PLAT_WINDOWS #include typedef void* LibModule; #endif diff --git a/Sources/Backends/WebGPU/WebGPU.c b/Sources/Backends/WebGPU/WebGPU.c index bb7a487..afbf3dc 100644 --- a/Sources/Backends/WebGPU/WebGPU.c +++ b/Sources/Backends/WebGPU/WebGPU.c @@ -59,3 +59,200 @@ PulseBackendHandler WebGPUDriver = { .supported_shader_formats = PULSE_SHADER_FORMAT_WGSL_BIT, .driver_data = PULSE_NULLPTR }; + +int32_t WebGPUGetImageBlockWidth(PulseImageFormat format) +{ + switch(format) + { + case PULSE_IMAGE_FORMAT_BC1_RGBA_UNORM: + case PULSE_IMAGE_FORMAT_BC2_RGBA_UNORM: + case PULSE_IMAGE_FORMAT_BC3_RGBA_UNORM: + case PULSE_IMAGE_FORMAT_BC4_R_UNORM: + case PULSE_IMAGE_FORMAT_BC5_RG_UNORM: + case PULSE_IMAGE_FORMAT_BC7_RGBA_UNORM: + case PULSE_IMAGE_FORMAT_BC6H_RGB_FLOAT: + case PULSE_IMAGE_FORMAT_BC6H_RGB_UFLOAT: return 4; + + case PULSE_IMAGE_FORMAT_R8G8B8A8_UNORM: + case PULSE_IMAGE_FORMAT_B8G8R8A8_UNORM: + case PULSE_IMAGE_FORMAT_B5G6R5_UNORM: + case PULSE_IMAGE_FORMAT_B5G5R5A1_UNORM: + case PULSE_IMAGE_FORMAT_B4G4R4A4_UNORM: + case PULSE_IMAGE_FORMAT_R10G10B10A2_UNORM: + case PULSE_IMAGE_FORMAT_R8G8_UNORM: + case PULSE_IMAGE_FORMAT_R16G16_UNORM: + case PULSE_IMAGE_FORMAT_R16G16B16A16_UNORM: + case PULSE_IMAGE_FORMAT_R8_UNORM: + case PULSE_IMAGE_FORMAT_R16_UNORM: + case PULSE_IMAGE_FORMAT_A8_UNORM: + case PULSE_IMAGE_FORMAT_R8_SNORM: + case PULSE_IMAGE_FORMAT_R8G8_SNORM: + case PULSE_IMAGE_FORMAT_R8G8B8A8_SNORM: + case PULSE_IMAGE_FORMAT_R16_SNORM: + case PULSE_IMAGE_FORMAT_R16G16_SNORM: + case PULSE_IMAGE_FORMAT_R16G16B16A16_SNORM: + case PULSE_IMAGE_FORMAT_R16_FLOAT: + case PULSE_IMAGE_FORMAT_R16G16_FLOAT: + case PULSE_IMAGE_FORMAT_R16G16B16A16_FLOAT: + case PULSE_IMAGE_FORMAT_R32_FLOAT: + case PULSE_IMAGE_FORMAT_R32G32_FLOAT: + case PULSE_IMAGE_FORMAT_R32G32B32A32_FLOAT: + case PULSE_IMAGE_FORMAT_R11G11B10_UFLOAT: + case PULSE_IMAGE_FORMAT_R8_UINT: + case PULSE_IMAGE_FORMAT_R8G8_UINT: + case PULSE_IMAGE_FORMAT_R8G8B8A8_UINT: + case PULSE_IMAGE_FORMAT_R16_UINT: + case PULSE_IMAGE_FORMAT_R16G16_UINT: + case PULSE_IMAGE_FORMAT_R16G16B16A16_UINT: + case PULSE_IMAGE_FORMAT_R32_UINT: + case PULSE_IMAGE_FORMAT_R32G32_UINT: + case PULSE_IMAGE_FORMAT_R32G32B32A32_UINT: + case PULSE_IMAGE_FORMAT_R8_INT: + case PULSE_IMAGE_FORMAT_R8G8_INT: + case PULSE_IMAGE_FORMAT_R8G8B8A8_INT: + case PULSE_IMAGE_FORMAT_R16_INT: + case PULSE_IMAGE_FORMAT_R16G16_INT: + case PULSE_IMAGE_FORMAT_R16G16B16A16_INT: + case PULSE_IMAGE_FORMAT_R32_INT: + case PULSE_IMAGE_FORMAT_R32G32_INT: + case PULSE_IMAGE_FORMAT_R32G32B32A32_INT: return 1; + + default: return 0; + } +} + +int32_t WebGPUGetImageBlockHeight(PulseImageFormat format) +{ + switch(format) + { + case PULSE_IMAGE_FORMAT_BC1_RGBA_UNORM: + case PULSE_IMAGE_FORMAT_BC2_RGBA_UNORM: + case PULSE_IMAGE_FORMAT_BC3_RGBA_UNORM: + case PULSE_IMAGE_FORMAT_BC4_R_UNORM: + case PULSE_IMAGE_FORMAT_BC5_RG_UNORM: + case PULSE_IMAGE_FORMAT_BC7_RGBA_UNORM: + case PULSE_IMAGE_FORMAT_BC6H_RGB_FLOAT: + case PULSE_IMAGE_FORMAT_BC6H_RGB_UFLOAT: return 4; + + case PULSE_IMAGE_FORMAT_R8G8B8A8_UNORM: + case PULSE_IMAGE_FORMAT_B8G8R8A8_UNORM: + case PULSE_IMAGE_FORMAT_B5G6R5_UNORM: + case PULSE_IMAGE_FORMAT_B5G5R5A1_UNORM: + case PULSE_IMAGE_FORMAT_B4G4R4A4_UNORM: + case PULSE_IMAGE_FORMAT_R10G10B10A2_UNORM: + case PULSE_IMAGE_FORMAT_R8G8_UNORM: + case PULSE_IMAGE_FORMAT_R16G16_UNORM: + case PULSE_IMAGE_FORMAT_R16G16B16A16_UNORM: + case PULSE_IMAGE_FORMAT_R8_UNORM: + case PULSE_IMAGE_FORMAT_R16_UNORM: + case PULSE_IMAGE_FORMAT_A8_UNORM: + case PULSE_IMAGE_FORMAT_R8_SNORM: + case PULSE_IMAGE_FORMAT_R8G8_SNORM: + case PULSE_IMAGE_FORMAT_R8G8B8A8_SNORM: + case PULSE_IMAGE_FORMAT_R16_SNORM: + case PULSE_IMAGE_FORMAT_R16G16_SNORM: + case PULSE_IMAGE_FORMAT_R16G16B16A16_SNORM: + case PULSE_IMAGE_FORMAT_R16_FLOAT: + case PULSE_IMAGE_FORMAT_R16G16_FLOAT: + case PULSE_IMAGE_FORMAT_R16G16B16A16_FLOAT: + case PULSE_IMAGE_FORMAT_R32_FLOAT: + case PULSE_IMAGE_FORMAT_R32G32_FLOAT: + case PULSE_IMAGE_FORMAT_R32G32B32A32_FLOAT: + case PULSE_IMAGE_FORMAT_R11G11B10_UFLOAT: + case PULSE_IMAGE_FORMAT_R8_UINT: + case PULSE_IMAGE_FORMAT_R8G8_UINT: + case PULSE_IMAGE_FORMAT_R8G8B8A8_UINT: + case PULSE_IMAGE_FORMAT_R16_UINT: + case PULSE_IMAGE_FORMAT_R16G16_UINT: + case PULSE_IMAGE_FORMAT_R16G16B16A16_UINT: + case PULSE_IMAGE_FORMAT_R32_UINT: + case PULSE_IMAGE_FORMAT_R32G32_UINT: + case PULSE_IMAGE_FORMAT_R32G32B32A32_UINT: + case PULSE_IMAGE_FORMAT_R8_INT: + case PULSE_IMAGE_FORMAT_R8G8_INT: + case PULSE_IMAGE_FORMAT_R8G8B8A8_INT: + case PULSE_IMAGE_FORMAT_R16_INT: + case PULSE_IMAGE_FORMAT_R16G16_INT: + case PULSE_IMAGE_FORMAT_R16G16B16A16_INT: + case PULSE_IMAGE_FORMAT_R32_INT: + case PULSE_IMAGE_FORMAT_R32G32_INT: + case PULSE_IMAGE_FORMAT_R32G32B32A32_INT: return 1; + + default: return 0; + } +} + +uint32_t WebGPUImageFormatTexelBlockSize(PulseImageFormat format) +{ + switch(format) + { + case PULSE_IMAGE_FORMAT_BC1_RGBA_UNORM: + case PULSE_IMAGE_FORMAT_BC4_R_UNORM: return 8; + + case PULSE_IMAGE_FORMAT_BC2_RGBA_UNORM: + case PULSE_IMAGE_FORMAT_BC3_RGBA_UNORM: + case PULSE_IMAGE_FORMAT_BC5_RG_UNORM: + case PULSE_IMAGE_FORMAT_BC7_RGBA_UNORM: + case PULSE_IMAGE_FORMAT_BC6H_RGB_FLOAT: + case PULSE_IMAGE_FORMAT_BC6H_RGB_UFLOAT: return 16; + + case PULSE_IMAGE_FORMAT_R8_UNORM: + case PULSE_IMAGE_FORMAT_R8_SNORM: + case PULSE_IMAGE_FORMAT_A8_UNORM: + case PULSE_IMAGE_FORMAT_R8_UINT: + case PULSE_IMAGE_FORMAT_R8_INT: return 1; + + case PULSE_IMAGE_FORMAT_B5G6R5_UNORM: + case PULSE_IMAGE_FORMAT_B4G4R4A4_UNORM: + case PULSE_IMAGE_FORMAT_B5G5R5A1_UNORM: + case PULSE_IMAGE_FORMAT_R16_FLOAT: + case PULSE_IMAGE_FORMAT_R8G8_SNORM: + case PULSE_IMAGE_FORMAT_R8G8_UNORM: + case PULSE_IMAGE_FORMAT_R8G8_UINT: + case PULSE_IMAGE_FORMAT_R8G8_INT: + case PULSE_IMAGE_FORMAT_R16_UNORM: + case PULSE_IMAGE_FORMAT_R16_SNORM: + case PULSE_IMAGE_FORMAT_R16_UINT: + case PULSE_IMAGE_FORMAT_R16_INT: return 2; + + case PULSE_IMAGE_FORMAT_R8G8B8A8_UNORM: + case PULSE_IMAGE_FORMAT_B8G8R8A8_UNORM: + case PULSE_IMAGE_FORMAT_R32_FLOAT: + case PULSE_IMAGE_FORMAT_R16G16_FLOAT: + case PULSE_IMAGE_FORMAT_R11G11B10_UFLOAT: + case PULSE_IMAGE_FORMAT_R8G8B8A8_SNORM: + case PULSE_IMAGE_FORMAT_R10G10B10A2_UNORM: + case PULSE_IMAGE_FORMAT_R8G8B8A8_UINT: + case PULSE_IMAGE_FORMAT_R8G8B8A8_INT: + case PULSE_IMAGE_FORMAT_R16G16_UINT: + case PULSE_IMAGE_FORMAT_R16G16_INT: + case PULSE_IMAGE_FORMAT_R16G16_UNORM: + case PULSE_IMAGE_FORMAT_R16G16_SNORM: + case PULSE_IMAGE_FORMAT_R32_UINT: + case PULSE_IMAGE_FORMAT_R32_INT: return 4; + + case PULSE_IMAGE_FORMAT_R16G16B16A16_FLOAT: + case PULSE_IMAGE_FORMAT_R16G16B16A16_UNORM: + case PULSE_IMAGE_FORMAT_R16G16B16A16_SNORM: + case PULSE_IMAGE_FORMAT_R16G16B16A16_UINT: + case PULSE_IMAGE_FORMAT_R16G16B16A16_INT: + case PULSE_IMAGE_FORMAT_R32G32_FLOAT: + case PULSE_IMAGE_FORMAT_R32G32_UINT: + case PULSE_IMAGE_FORMAT_R32G32_INT: return 8; + + case PULSE_IMAGE_FORMAT_R32G32B32A32_FLOAT: + case PULSE_IMAGE_FORMAT_R32G32B32A32_INT: + case PULSE_IMAGE_FORMAT_R32G32B32A32_UINT: return 16; + + default: return 0; + } +} + +uint32_t WebGPUBytesPerRow(int32_t width, PulseImageFormat format) +{ + uint32_t block_width = WebGPUGetImageBlockWidth(format); + if(block_width == 0) + return 0; + uint32_t blocks_per_row = (width + block_width - 1) / block_width; + return blocks_per_row * WebGPUImageFormatTexelBlockSize(format); +} diff --git a/Sources/Backends/WebGPU/WebGPU.h b/Sources/Backends/WebGPU/WebGPU.h index 867b7cc..8941a71 100644 --- a/Sources/Backends/WebGPU/WebGPU.h +++ b/Sources/Backends/WebGPU/WebGPU.h @@ -9,6 +9,10 @@ #ifndef PULSE_WEBGPU_H_ #define PULSE_WEBGPU_H_ +#ifdef __STDC_NO_ATOMICS__ + #error "Atomic support is not present" +#endif + #include #define WEBGPU_RETRIEVE_DRIVER_DATA_AS(handle, cast) ((cast)handle->driver_data) @@ -21,6 +25,11 @@ typedef struct WebGPUDriverData PulseBackendFlags WebGPUCheckSupport(PulseBackendFlags candidates, PulseShaderFormatsFlags shader_formats_used); // Return PULSE_BACKEND_WEBGPU in case of success and PULSE_BACKEND_INVALID otherwise void WebGPUDeviceTick(PulseDevice device); +uint32_t WebGPUBytesPerRow(int32_t width, PulseImageFormat format); +uint32_t WebGPUImageFormatTexelBlockSize(PulseImageFormat format); +int32_t WebGPUGetImageBlockHeight(PulseImageFormat format); +int32_t WebGPUGetImageBlockWidth(PulseImageFormat format); + #endif // PULSE_WEBGPU_H_ #endif // PULSE_ENABLE_WEBGPU_BACKEND diff --git a/Sources/Backends/WebGPU/WebGPUBuffer.c b/Sources/Backends/WebGPU/WebGPUBuffer.c index 9eeb3f2..9dfdd48 100644 --- a/Sources/Backends/WebGPU/WebGPUBuffer.c +++ b/Sources/Backends/WebGPU/WebGPUBuffer.c @@ -10,6 +10,7 @@ #include "WebGPU.h" #include "WebGPUDevice.h" #include "WebGPUBuffer.h" +#include "WebGPUImage.h" #include "WebGPUCommandList.h" PulseBuffer WebGPUCreateBuffer(PulseDevice device, const PulseBufferCreateInfo* create_infos) @@ -156,6 +157,52 @@ bool WebGPUCopyBufferToBuffer(PulseCommandList cmd, const PulseBufferRegion* src bool WebGPUCopyBufferToImage(PulseCommandList cmd, const PulseBufferRegion* src, const PulseImageRegion* dst) { + WebGPUDevice* webgpu_device = WEBGPU_RETRIEVE_DRIVER_DATA_AS(cmd->device, WebGPUDevice*); + WebGPUBuffer* webgpu_src_buffer = WEBGPU_RETRIEVE_DRIVER_DATA_AS(src->buffer, WebGPUBuffer*); + WebGPUImage* webgpu_dst_image = WEBGPU_RETRIEVE_DRIVER_DATA_AS(dst->image, WebGPUImage*); + WebGPUCommandList* webgpu_cmd = WEBGPU_RETRIEVE_DRIVER_DATA_AS(cmd, WebGPUCommandList*); + + PulseImageFormat format = dst->image->format; + uint32_t block_height = WebGPUGetImageBlockHeight(format) > 1 ? WebGPUGetImageBlockHeight(format) : 1; + uint32_t blocks_per_column = (dst->image->height + block_height - 1) / block_height; + uint32_t bytes_per_row = WebGPUBytesPerRow(dst->image->width, dst->image->format); + + if(bytes_per_row == 0) + { + if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(cmd->device->backend)) + PulseLogError(cmd->device->backend, "(WebGPU) unsupported image format"); + PulseSetInternalError(PULSE_ERROR_INVALID_IMAGE_FORMAT); + return false; + } + + WGPUTexelCopyBufferLayout layout = { 0 }; + layout.bytesPerRow = bytes_per_row; + layout.rowsPerImage = blocks_per_column; + + WGPUTexelCopyTextureInfo texture_copy_info = { 0 }; + texture_copy_info.texture = webgpu_dst_image->texture; + texture_copy_info.mipLevel = 1; + texture_copy_info.aspect = WGPUTextureAspect_All; + texture_copy_info.origin.x = dst->x; + texture_copy_info.origin.y = dst->y; + texture_copy_info.origin.z = dst->z; + + WGPUExtent3D extent = { 0 }; + extent.width = dst->width; + extent.height = dst->height; + extent.depthOrArrayLayers = dst->depth; + + if(bytes_per_row >= 256 && bytes_per_row % 256 == 0) + { + WGPUTexelCopyBufferInfo buffer_copy_info = { 0 }; + buffer_copy_info.buffer = webgpu_src_buffer->buffer; + buffer_copy_info.layout = layout; + wgpuCommandEncoderCopyBufferToTexture(webgpu_cmd->encoder, &buffer_copy_info, &texture_copy_info, &extent); + } + else + wgpuQueueWriteTexture(webgpu_device->queue, &texture_copy_info, webgpu_src_buffer->buffer, src->size, &layout, &extent); + + return true; } void WebGPUDestroyBuffer(PulseDevice device, PulseBuffer buffer) diff --git a/Sources/Backends/WebGPU/WebGPUComputePass.c b/Sources/Backends/WebGPU/WebGPUComputePass.c index 59f2cc2..b8d882b 100644 --- a/Sources/Backends/WebGPU/WebGPUComputePass.c +++ b/Sources/Backends/WebGPU/WebGPUComputePass.c @@ -8,6 +8,7 @@ #include "../../PulseInternal.h" #include "WebGPU.h" #include "WebGPUDevice.h" +#include "WebGPUImage.h" #include "WebGPUComputePass.h" #include "WebGPUComputePipeline.h" @@ -151,9 +152,12 @@ static void WebGPUBindBindGroups(PulseComputePass pass) uint32_t entry_index = 0; for(uint32_t i = 0; i < pass->current_pipeline->num_readonly_storage_images; i++, entry_index++) { + WebGPUImage* webgpu_image = WEBGPU_RETRIEVE_DRIVER_DATA_AS(pass->readonly_images[i], WebGPUImage*); + WGPUBindGroupEntry* entry = &read_only_entries[entry_index]; memset(entry, 0, sizeof(WGPUBindGroupEntry)); entry->binding = i; + entry->textureView = webgpu_image->view; } for(uint32_t i = 0; i < pass->current_pipeline->num_readonly_storage_buffers; i++, entry_index++) @@ -180,9 +184,12 @@ static void WebGPUBindBindGroups(PulseComputePass pass) uint32_t entry_index = 0; for(uint32_t i = 0; i < pass->current_pipeline->num_readwrite_storage_images; i++, entry_index++) { + WebGPUImage* webgpu_image = WEBGPU_RETRIEVE_DRIVER_DATA_AS(pass->readwrite_images[i], WebGPUImage*); + WGPUBindGroupEntry* entry = &read_write_entries[entry_index]; memset(entry, 0, sizeof(WGPUBindGroupEntry)); entry->binding = i; + entry->textureView = webgpu_image->view; } for(uint32_t i = 0; i < pass->current_pipeline->num_readwrite_storage_buffers; i++, entry_index++) diff --git a/Sources/Backends/WebGPU/WebGPUDevice.c b/Sources/Backends/WebGPU/WebGPUDevice.c index 9335741..f425351 100644 --- a/Sources/Backends/WebGPU/WebGPUDevice.c +++ b/Sources/Backends/WebGPU/WebGPUDevice.c @@ -38,9 +38,9 @@ static uint64_t WebGPUScoreAdapter(WGPUAdapter adapter) { uint64_t score = 0; - WGPUAdapterInfo infos; + WGPUAdapterInfo infos = { 0 }; wgpuAdapterGetInfo(adapter, &infos); - WGPULimits limits; + WGPULimits limits = { 0 }; wgpuAdapterGetLimits(adapter, &limits); if(infos.adapterType == WGPUAdapterType_DiscreteGPU) @@ -106,7 +106,7 @@ static void WebGPUDeviceLostCallback(const WGPUDevice* _, WGPUDeviceLostReason r "creation failed", }; if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(backend)) - PulseLogErrorFmt(backend, "(WebGPU) device %.*s lost because %s. %.*s", device->infos.device.length, device->infos.device.data, reasons[reason], message.length, message.data); + PulseLogErrorFmt(backend, "(WebGPU) device %.*s lost because %s. %.*s", device->infos.device.length, device->infos.device.data, reasons[reason - 1], message.length, message.data); } static void WebGPUDeviceUncapturedErrorCallback(const WGPUDevice* _, WGPUErrorType type, WGPUStringView message, void* userdata1, void* userdata2) @@ -121,7 +121,7 @@ static void WebGPUDeviceUncapturedErrorCallback(const WGPUDevice* _, WGPUErrorTy "has recieved an unknown error", }; if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(backend)) - PulseLogErrorFmt(backend, "(WebGPU) device %.*s %s. %.*s", device->infos.device.length, device->infos.device.data, types[type], message.length, message.data); + PulseLogErrorFmt(backend, "(WebGPU) device %.*s %s. %.*s", device->infos.device.length, device->infos.device.data, types[type - 1], message.length, message.data); } PulseDevice WebGPUCreateDevice(PulseBackend backend, PulseDevice* forbiden_devices, uint32_t forbiden_devices_count) diff --git a/Sources/Backends/WebGPU/WebGPUImage.c b/Sources/Backends/WebGPU/WebGPUImage.c index 016b365..4abbfba 100644 --- a/Sources/Backends/WebGPU/WebGPUImage.c +++ b/Sources/Backends/WebGPU/WebGPUImage.c @@ -74,6 +74,15 @@ static WGPUTextureDimension PulseImageTypeToWGPUTextureDimension[] = { }; PULSE_STATIC_ASSERT(PulseImageTypeToWGPUTextureDimension, (sizeof(PulseImageTypeToWGPUTextureDimension) / sizeof(WGPUTextureDimension)) == PULSE_IMAGE_TYPE_MAX_ENUM); +static WGPUTextureViewDimension PulseImageTypeToWGPUTextureViewDimension[] = { + WGPUTextureViewDimension_2D, //PULSE_IMAGE_TYPE_2D + WGPUTextureViewDimension_2DArray, //PULSE_IMAGE_TYPE_2D_ARRAY + WGPUTextureViewDimension_3D, //PULSE_IMAGE_TYPE_3D + WGPUTextureViewDimension_Cube, //PULSE_IMAGE_TYPE_CUBE + WGPUTextureViewDimension_CubeArray, //PULSE_IMAGE_TYPE_CUBE_ARRAY +}; +PULSE_STATIC_ASSERT(PulseImageTypeToWGPUTextureViewDimension, (sizeof(PulseImageTypeToWGPUTextureViewDimension) / sizeof(WGPUTextureViewDimension)) == PULSE_IMAGE_TYPE_MAX_ENUM); + PulseImage WebGPUCreateImage(PulseDevice device, const PulseImageCreateInfo* create_infos) { WebGPUDevice* webgpu_device = WEBGPU_RETRIEVE_DRIVER_DATA_AS(device, WebGPUDevice*); @@ -121,6 +130,24 @@ PulseImage WebGPUCreateImage(PulseDevice device, const PulseImageCreateInfo* cre webgpu_image->texture = wgpuDeviceCreateTexture(webgpu_device->device, &descriptor); if(webgpu_image->texture == PULSE_NULLPTR) { + PulseSetInternalError(PULSE_ERROR_INVALID_IMAGE_FORMAT); + free(webgpu_image); + free(image); + return PULSE_NULL_HANDLE; + } + + WGPUTextureViewDescriptor view_descriptor = { 0 }; + view_descriptor.format = descriptor.format; + view_descriptor.dimension = PulseImageTypeToWGPUTextureViewDimension[create_infos->type]; + view_descriptor.baseMipLevel = 0; + view_descriptor.mipLevelCount = 1; + view_descriptor.baseArrayLayer = 0; + view_descriptor.arrayLayerCount = view_descriptor.dimension == WGPUTextureViewDimension_3D ? 1 : create_infos->layer_count_or_depth; + webgpu_image->view = wgpuTextureCreateView(webgpu_image->texture, &view_descriptor); + if(webgpu_image->texture == PULSE_NULLPTR) + { + PulseSetInternalError(PULSE_ERROR_INVALID_IMAGE_FORMAT); + wgpuTextureRelease(webgpu_image->texture); free(webgpu_image); free(image); return PULSE_NULL_HANDLE; @@ -173,6 +200,47 @@ bool WebGPUIsImageFormatValid(PulseDevice device, PulseImageFormat format, Pulse bool WebGPUCopyImageToBuffer(PulseCommandList cmd, const PulseImageRegion* src, const PulseBufferRegion* dst) { + WebGPUImage* webgpu_src_image = WEBGPU_RETRIEVE_DRIVER_DATA_AS(src->image, WebGPUImage*); + WebGPUBuffer* webgpu_dst_buffer = WEBGPU_RETRIEVE_DRIVER_DATA_AS(dst->buffer, WebGPUBuffer*); + WebGPUCommandList* webgpu_cmd = WEBGPU_RETRIEVE_DRIVER_DATA_AS(cmd, WebGPUCommandList*); + + PulseImageFormat format = src->image->format; + uint32_t block_height = WebGPUGetImageBlockHeight(format) > 1 ? WebGPUGetImageBlockHeight(format) : 1; + uint32_t blocks_per_column = (src->image->height + block_height - 1) / block_height; + uint32_t bytes_per_row = WebGPUBytesPerRow(src->image->width, src->image->format); + + if(bytes_per_row == 0) + { + if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(cmd->device->backend)) + PulseLogError(cmd->device->backend, "(WebGPU) unsupported image format"); + PulseSetInternalError(PULSE_ERROR_INVALID_IMAGE_FORMAT); + return false; + } + + WGPUTexelCopyBufferLayout layout = { 0 }; + layout.bytesPerRow = bytes_per_row; + layout.rowsPerImage = blocks_per_column; + + WGPUTexelCopyTextureInfo texture_copy_info = { 0 }; + texture_copy_info.texture = webgpu_src_image->texture; + texture_copy_info.mipLevel = 1; + texture_copy_info.aspect = WGPUTextureAspect_All; + texture_copy_info.origin.x = src->x; + texture_copy_info.origin.y = src->y; + texture_copy_info.origin.z = src->z; + + WGPUTexelCopyBufferInfo buffer_copy_info = { 0 }; + buffer_copy_info.buffer = webgpu_dst_buffer->buffer; + buffer_copy_info.layout = layout; + + WGPUExtent3D extent = { 0 }; + extent.width = src->width; + extent.height = src->height; + extent.depthOrArrayLayers = src->depth; + + wgpuCommandEncoderCopyTextureToBuffer(webgpu_cmd->encoder, &texture_copy_info, &buffer_copy_info, &extent); + + return true; } bool WebGPUBlitImage(PulseCommandList cmd, const PulseImageRegion* src, const PulseImageRegion* dst) @@ -183,6 +251,7 @@ void WebGPUDestroyImage(PulseDevice device, PulseImage image) { PULSE_UNUSED(device); WebGPUImage* webgpu_image = WEBGPU_RETRIEVE_DRIVER_DATA_AS(image, WebGPUImage*); + wgpuTextureViewRelease(webgpu_image->view); wgpuTextureRelease(webgpu_image->texture); free(webgpu_image); free(image); diff --git a/Sources/Backends/WebGPU/WebGPUImage.h b/Sources/Backends/WebGPU/WebGPUImage.h index 6f94d53..100ab6a 100644 --- a/Sources/Backends/WebGPU/WebGPUImage.h +++ b/Sources/Backends/WebGPU/WebGPUImage.h @@ -15,6 +15,7 @@ typedef struct WebGPUImage { WGPUTexture texture; + WGPUTextureView view; } WebGPUImage; PulseImage WebGPUCreateImage(PulseDevice device, const PulseImageCreateInfo* create_infos); diff --git a/Tests/Backend.c b/Tests/Backend.c new file mode 100644 index 0000000..1fd6ec1 --- /dev/null +++ b/Tests/Backend.c @@ -0,0 +1,72 @@ +#include "Common.h" + +#include +#include + +void DumbDebugCallBack(PulseDebugMessageSeverity severity, const char* message) +{ + (void)severity; + (void)message; +} + +void TestVulkanSupport() +{ + #if defined(VULKAN_ENABLED) + if(!PulseSupportsBackend(PULSE_BACKEND_VULKAN, PULSE_SHADER_FORMAT_SPIRV_BIT)) + #elif defined(WEBGPU_ENABLED) + if(!PulseSupportsBackend(PULSE_BACKEND_WEBGPU, PULSE_SHADER_FORMAT_WGSL_BIT)) + #endif + { + TEST_MESSAGE("Vulkan is not supported"); + exit(0); + } +} + +void TestBackendSetup() +{ + #if defined(VULKAN_ENABLED) + PulseBackend backend = PulseLoadBackend(PULSE_BACKEND_VULKAN, PULSE_SHADER_FORMAT_SPIRV_BIT, PULSE_HIGH_DEBUG); + #elif defined(WEBGPU_ENABLED) + PulseBackend backend = PulseLoadBackend(PULSE_BACKEND_WEBGPU, PULSE_SHADER_FORMAT_WGSL_BIT, PULSE_HIGH_DEBUG); + #endif + TEST_ASSERT_NOT_EQUAL_MESSAGE(backend, PULSE_NULL_HANDLE, PulseVerbaliseErrorType(PulseGetLastErrorType())); + PulseSetDebugCallback(backend, DumbDebugCallBack); + PulseUnloadBackend(backend); +} + +void TestBackendAnySetup() +{ + #if defined(VULKAN_ENABLED) + PulseBackend backend = PulseLoadBackend(PULSE_BACKEND_ANY, PULSE_SHADER_FORMAT_SPIRV_BIT, PULSE_HIGH_DEBUG); + #elif defined(WEBGPU_ENABLED) + PulseBackend backend = PulseLoadBackend(PULSE_BACKEND_ANY, PULSE_SHADER_FORMAT_WGSL_BIT, PULSE_HIGH_DEBUG); + #endif + TEST_ASSERT_NOT_EQUAL_MESSAGE(backend, PULSE_NULL_HANDLE, PulseVerbaliseErrorType(PulseGetLastErrorType())); + #if defined(VULKAN_ENABLED) + TEST_ASSERT_EQUAL(PulseGetBackendType(backend), PULSE_BACKEND_VULKAN); + #elif defined(WEBGPU_ENABLED) + TEST_ASSERT_EQUAL(PulseGetBackendType(backend), PULSE_BACKEND_WEBGPU); + #endif + PulseSetDebugCallback(backend, DumbDebugCallBack); + PulseUnloadBackend(backend); +} + +void TestWrongBackendSetup() +{ + #if defined(VULKAN_ENABLED) + PulseBackend backend = PulseLoadBackend(PULSE_BACKEND_VULKAN, PULSE_SHADER_FORMAT_MSL_BIT, PULSE_HIGH_DEBUG); + #elif defined(WEBGPU_ENABLED) + PulseBackend backend = PulseLoadBackend(PULSE_BACKEND_WEBGPU, PULSE_SHADER_FORMAT_MSL_BIT, PULSE_HIGH_DEBUG); + #endif + TEST_ASSERT_EQUAL(backend, PULSE_NULL_HANDLE); + PulseSetDebugCallback(backend, DumbDebugCallBack); + PulseUnloadBackend(backend); +} + +void TestBackend() +{ + RUN_TEST(TestVulkanSupport); + RUN_TEST(TestBackendSetup); + RUN_TEST(TestBackendAnySetup); + RUN_TEST(TestWrongBackendSetup); +} diff --git a/Tests/Vulkan/Buffer.c b/Tests/Buffer.c similarity index 96% rename from Tests/Vulkan/Buffer.c rename to Tests/Buffer.c index 567caf9..fcd7621 100644 --- a/Tests/Vulkan/Buffer.c +++ b/Tests/Buffer.c @@ -54,15 +54,6 @@ void TestBufferCreation() TEST_ASSERT_NOT_EQUAL_MESSAGE(buffer, PULSE_NULL_HANDLE, PulseVerbaliseErrorType(PulseGetLastErrorType())); PulseDestroyBuffer(device, buffer); - DISABLE_ERRORS; - buffer_create_info.size = -1; - buffer_create_info.usage = PULSE_BUFFER_USAGE_STORAGE_READ; - buffer = PulseCreateBuffer(device, &buffer_create_info); - TEST_ASSERT_EQUAL(buffer, PULSE_NULL_HANDLE); - PulseGetLastErrorType(); // Just to clear the error code - PulseDestroyBuffer(device, buffer); - ENABLE_ERRORS; - CleanupDevice(device); CleanupPulse(backend); } @@ -278,9 +269,14 @@ void TestBufferComputeWrite() PulseDevice device; SetupDevice(backend, &device); - const uint8_t shader_bytecode[] = { - #include "Shaders/SimpleBufferWrite.spv.h" - }; + #if defined(VULKAN_ENABLED) + const uint8_t shader_bytecode[] = { + #include "Shaders/Vulkan/SimpleBufferWrite.spv.h" + }; + #elif defined(WEBGPU_ENABLED) + #define SHADER_NAME shader_bytecode + #include "Shaders/WebGPU/SimpleBufferWrite.wgsl.h" + #endif PulseBufferCreateInfo buffer_create_info = { 0 }; buffer_create_info.size = 256 * sizeof(int32_t); @@ -300,7 +296,7 @@ void TestBufferComputeWrite() TEST_ASSERT_NOT_EQUAL_MESSAGE(pass, PULSE_NULL_HANDLE, PulseVerbaliseErrorType(PulseGetLastErrorType())); PulseBindStorageBuffers(pass, &buffer, 1); PulseBindComputePipeline(pass, pipeline); - PulseDispatchComputations(pass, 32, 32, 1); + PulseDispatchComputations(pass, 16, 1, 1); PulseEndComputePass(pass); TEST_ASSERT_TRUE_MESSAGE(PulseSubmitCommandList(device, cmd, fence), PulseVerbaliseErrorType(PulseGetLastErrorType())); @@ -340,9 +336,14 @@ void TestBufferComputeCopy() PulseDevice device; SetupDevice(backend, &device); - const uint8_t shader_bytecode[] = { - #include "Shaders/BufferCopy.spv.h" - }; + #if defined(VULKAN_ENABLED) + const uint8_t shader_bytecode[] = { + #include "Shaders/Vulkan/BufferCopy.spv.h" + }; + #elif defined(WEBGPU_ENABLED) + #define SHADER_NAME shader_bytecode + #include "Shaders/WebGPU/BufferCopy.wgsl.h" + #endif uint32_t data[256]; memset(data, 0xFF, 256 * sizeof(uint32_t)); @@ -386,7 +387,7 @@ void TestBufferComputeCopy() PulseBindStorageBuffers(pass, &read_buffer, 1); PulseBindStorageBuffers(pass, &write_buffer, 1); PulseBindComputePipeline(pass, pipeline); - PulseDispatchComputations(pass, 32, 32, 1); + PulseDispatchComputations(pass, 16, 1, 1); PulseEndComputePass(pass); TEST_ASSERT_TRUE_MESSAGE(PulseSubmitCommandList(device, cmd, fence), PulseVerbaliseErrorType(PulseGetLastErrorType())); diff --git a/Tests/Vulkan/Common.c b/Tests/Common.c similarity index 78% rename from Tests/Vulkan/Common.c rename to Tests/Common.c index 043513b..78b53ab 100644 --- a/Tests/Vulkan/Common.c +++ b/Tests/Common.c @@ -1,5 +1,7 @@ #include "Common.h" #include +#include +#include bool errors_enabled = true; bool has_recieved_error = false; @@ -7,7 +9,10 @@ bool has_recieved_error = false; void DebugCallBack(PulseDebugMessageSeverity severity, const char* message) { if(errors_enabled && severity == PULSE_DEBUG_MESSAGE_SEVERITY_ERROR) - TEST_FAIL_MESSAGE(message); + { + fprintf(stderr, "%s", message); + TEST_FAIL(); + } has_recieved_error = true; } @@ -15,11 +20,15 @@ void DebugCallBack(PulseDebugMessageSeverity severity, const char* message) void SetupPulse(PulseBackend* backend) { - *backend = PulseLoadBackend(PULSE_BACKEND_VULKAN, PULSE_SHADER_FORMAT_SPIRV_BIT, PULSE_PARANOID_DEBUG); + #if defined(VULKAN_ENABLED) + *backend = PulseLoadBackend(PULSE_BACKEND_VULKAN, PULSE_SHADER_FORMAT_SPIRV_BIT, PULSE_PARANOID_DEBUG); + #elif defined(WEBGPU_ENABLED) + *backend = PulseLoadBackend(PULSE_BACKEND_WEBGPU, PULSE_SHADER_FORMAT_WGSL_BIT, PULSE_PARANOID_DEBUG); + #endif if(*backend == PULSE_NULL_HANDLE) { char complete_message[LOG_MESSAGE_MAX_LENGTH] = { 0 }; - snprintf(complete_message, LOG_MESSAGE_MAX_LENGTH, "Fatal Error: could not load Pulse using Vulkan due to %s", PulseVerbaliseErrorType(PulseGetLastErrorType())); + snprintf(complete_message, LOG_MESSAGE_MAX_LENGTH, "Fatal Error: could not load Pulse backend due to %s", PulseVerbaliseErrorType(PulseGetLastErrorType())); TEST_FAIL_MESSAGE(complete_message); TEST_ABORT(); } @@ -50,10 +59,18 @@ void LoadComputePipeline(PulseDevice device, PulseComputePipeline* pipeline, con uint32_t num_uniform_buffers) { PulseComputePipelineCreateInfo info = { 0 }; - info.code_size = code_size; + #if defined(WEBGPU_ENABLED) + info.code_size = strlen(code); + #else + info.code_size = code_size; + #endif info.code = code; info.entrypoint = "main"; - info.format = PULSE_SHADER_FORMAT_SPIRV_BIT; + #if defined(VULKAN_ENABLED) + info.format = PULSE_SHADER_FORMAT_SPIRV_BIT; + #elif defined(WEBGPU_ENABLED) + info.format = PULSE_SHADER_FORMAT_WGSL_BIT; + #endif info.num_readonly_storage_images = num_readonly_storage_images; info.num_readonly_storage_buffers = num_readonly_storage_buffers; info.num_readwrite_storage_buffers = num_readwrite_storage_buffers; diff --git a/Tests/Vulkan/Common.h b/Tests/Common.h similarity index 100% rename from Tests/Vulkan/Common.h rename to Tests/Common.h diff --git a/Tests/Vulkan/Device.c b/Tests/Device.c similarity index 78% rename from Tests/Vulkan/Device.c rename to Tests/Device.c index d2d3fa0..1df1b9e 100644 --- a/Tests/Vulkan/Device.c +++ b/Tests/Device.c @@ -46,7 +46,11 @@ void TestBackendInUse() PulseDevice device = PulseCreateDevice(backend, NULL, 0); TEST_ASSERT_NOT_EQUAL_MESSAGE(device, PULSE_NULL_HANDLE, PulseVerbaliseErrorType(PulseGetLastErrorType())); - TEST_ASSERT_EQUAL(PulseGetBackendInUseByDevice(device), PULSE_BACKEND_VULKAN); + #if defined(VULKAN_ENABLED) + TEST_ASSERT_EQUAL(PulseGetBackendInUseByDevice(device), PULSE_BACKEND_VULKAN); + #elif defined(WEBGPU_ENABLED) + TEST_ASSERT_EQUAL(PulseGetBackendInUseByDevice(device), PULSE_BACKEND_WEBGPU); + #endif PulseDestroyDevice(device); CleanupPulse(backend); @@ -59,7 +63,11 @@ void TestShaderFormatSupport() PulseDevice device = PulseCreateDevice(backend, NULL, 0); TEST_ASSERT_NOT_EQUAL_MESSAGE(device, PULSE_NULL_HANDLE, PulseVerbaliseErrorType(PulseGetLastErrorType())); - TEST_ASSERT_TRUE(PulseDeviceSupportsShaderFormats(device, PULSE_SHADER_FORMAT_SPIRV_BIT)); + #if defined(VULKAN_ENABLED) + TEST_ASSERT_TRUE(PulseDeviceSupportsShaderFormats(device, PULSE_SHADER_FORMAT_SPIRV_BIT)); + #elif defined(WEBGPU_ENABLED) + TEST_ASSERT_TRUE(PulseDeviceSupportsShaderFormats(device, PULSE_SHADER_FORMAT_WGSL_BIT)); + #endif PulseDestroyDevice(device); CleanupPulse(backend); diff --git a/Tests/Vulkan/Image.c b/Tests/Image.c similarity index 100% rename from Tests/Vulkan/Image.c rename to Tests/Image.c diff --git a/Tests/Vulkan/Pipeline.c b/Tests/Pipeline.c similarity index 88% rename from Tests/Vulkan/Pipeline.c rename to Tests/Pipeline.c index 48ce4d5..e370e5f 100644 --- a/Tests/Vulkan/Pipeline.c +++ b/Tests/Pipeline.c @@ -10,9 +10,14 @@ void TestPipelineSetup() PulseDevice device; SetupDevice(backend, &device); - const uint8_t shader_bytecode[] = { - #include "Shaders/Simple.spv.h" - }; + #if defined(VULKAN_ENABLED) + const uint8_t shader_bytecode[] = { + #include "Shaders/Vulkan/Simple.spv.h" + }; + #elif defined(WEBGPU_ENABLED) + #define SHADER_NAME shader_bytecode + #include "Shaders/WebGPU/Simple.wgsl.h" + #endif PulseComputePipeline pipeline; LoadComputePipeline(device, &pipeline, shader_bytecode, sizeof(shader_bytecode), 0, 0, 0, 0, 0); @@ -25,7 +30,7 @@ void TestPipelineSetup() PulseComputePass pass = PulseBeginComputePass(cmd); TEST_ASSERT_NOT_EQUAL_MESSAGE(pass, PULSE_NULL_HANDLE, PulseVerbaliseErrorType(PulseGetLastErrorType())); PulseBindComputePipeline(pass, pipeline); - PulseDispatchComputations(pass, 32, 32, 1); + PulseDispatchComputations(pass, 16, 1, 1); PulseEndComputePass(pass); TEST_ASSERT_TRUE_MESSAGE(PulseSubmitCommandList(device, cmd, fence), PulseVerbaliseErrorType(PulseGetLastErrorType())); @@ -46,9 +51,14 @@ void TestPipelineReadOnlyBindings() PulseDevice device; SetupDevice(backend, &device); - const uint8_t shader_bytecode[] = { - #include "Shaders/ReadOnlyBindings.spv.h" - }; + #if defined(VULKAN_ENABLED) + const uint8_t shader_bytecode[] = { + #include "Shaders/Vulkan/ReadOnlyBindings.spv.h" + }; + #elif defined(WEBGPU_ENABLED) + #define SHADER_NAME shader_bytecode + #include "Shaders/WebGPU/ReadOnlyBindings.wgsl.h" + #endif PulseBufferCreateInfo buffer_create_info = { 0 }; buffer_create_info.size = 256 * sizeof(int32_t); @@ -79,7 +89,7 @@ void TestPipelineReadOnlyBindings() PulseBindStorageImages(pass, &image, 1); PulseBindStorageBuffers(pass, &buffer, 1); PulseBindComputePipeline(pass, pipeline); - PulseDispatchComputations(pass, 32, 32, 1); + PulseDispatchComputations(pass, 16, 1, 1); PulseEndComputePass(pass); TEST_ASSERT_TRUE_MESSAGE(PulseSubmitCommandList(device, cmd, fence), PulseVerbaliseErrorType(PulseGetLastErrorType())); @@ -102,9 +112,14 @@ void TestPipelineWriteOnlyBindings() PulseDevice device; SetupDevice(backend, &device); - const uint8_t shader_bytecode[] = { - #include "Shaders/ReadOnlyBindings.spv.h" - }; + #if defined(VULKAN_ENABLED) + const uint8_t shader_bytecode[] = { + #include "Shaders/Vulkan/WriteOnlyBindings.spv.h" + }; + #elif defined(WEBGPU_ENABLED) + #define SHADER_NAME shader_bytecode + #include "Shaders/WebGPU/WriteOnlyBindings.wgsl.h" + #endif PulseBufferCreateInfo buffer_create_info = { 0 }; buffer_create_info.size = 256 * sizeof(int32_t); @@ -135,7 +150,7 @@ void TestPipelineWriteOnlyBindings() PulseBindStorageImages(pass, &image, 1); PulseBindStorageBuffers(pass, &buffer, 1); PulseBindComputePipeline(pass, pipeline); - PulseDispatchComputations(pass, 32, 32, 1); + PulseDispatchComputations(pass, 16, 1, 1); PulseEndComputePass(pass); TEST_ASSERT_TRUE_MESSAGE(PulseSubmitCommandList(device, cmd, fence), PulseVerbaliseErrorType(PulseGetLastErrorType())); @@ -158,9 +173,14 @@ void TestPipelineReadWriteBindings() PulseDevice device; SetupDevice(backend, &device); - const uint8_t shader_bytecode[] = { - #include "Shaders/ReadOnlyBindings.spv.h" - }; + #if defined(VULKAN_ENABLED) + const uint8_t shader_bytecode[] = { + #include "Shaders/Vulkan/ReadWriteBindings.spv.h" + }; + #elif defined(WEBGPU_ENABLED) + #define SHADER_NAME shader_bytecode + #include "Shaders/WebGPU/ReadWriteBindings.wgsl.h" + #endif PulseBufferCreateInfo buffer_create_info = { 0 }; buffer_create_info.size = 256 * sizeof(int32_t); @@ -201,7 +221,7 @@ void TestPipelineReadWriteBindings() PulseBindStorageBuffers(pass, &read_buffer, 1); PulseBindStorageBuffers(pass, &write_buffer, 1); PulseBindComputePipeline(pass, pipeline); - PulseDispatchComputations(pass, 32, 32, 1); + PulseDispatchComputations(pass, 16, 1, 1); PulseEndComputePass(pass); TEST_ASSERT_TRUE_MESSAGE(PulseSubmitCommandList(device, cmd, fence), PulseVerbaliseErrorType(PulseGetLastErrorType())); diff --git a/Tests/Shaders/.gitignore b/Tests/Shaders/.gitignore new file mode 100644 index 0000000..0cbdb90 --- /dev/null +++ b/Tests/Shaders/.gitignore @@ -0,0 +1,2 @@ +*.spv.h +*.wgsl.h diff --git a/Tests/Vulkan/Shaders/BufferCopy.nzsl b/Tests/Shaders/Vulkan/BufferCopy.nzsl similarity index 95% rename from Tests/Vulkan/Shaders/BufferCopy.nzsl rename to Tests/Shaders/Vulkan/BufferCopy.nzsl index 9c9d75f..49143de 100644 --- a/Tests/Vulkan/Shaders/BufferCopy.nzsl +++ b/Tests/Shaders/Vulkan/BufferCopy.nzsl @@ -19,7 +19,7 @@ external } [entry(compute)] -[workgroup(32, 32, 1)] +[workgroup(16, 16, 1)] fn main(input: Input) { write_ssbo.data[input.indices.x * input.indices.y] = read_ssbo.data[input.indices.x * input.indices.y]; diff --git a/Tests/Vulkan/Shaders/ReadOnlyBindings.nzsl b/Tests/Shaders/Vulkan/ReadOnlyBindings.nzsl similarity index 93% rename from Tests/Vulkan/Shaders/ReadOnlyBindings.nzsl rename to Tests/Shaders/Vulkan/ReadOnlyBindings.nzsl index 47ffc01..603e328 100644 --- a/Tests/Vulkan/Shaders/ReadOnlyBindings.nzsl +++ b/Tests/Shaders/Vulkan/ReadOnlyBindings.nzsl @@ -19,7 +19,7 @@ external } [entry(compute)] -[workgroup(32, 32, 1)] +[workgroup(16, 16, 1)] fn main(input: Input) { } diff --git a/Tests/Vulkan/Shaders/ReadWriteBindings.nzsl b/Tests/Shaders/Vulkan/ReadWriteBindings.nzsl similarity index 95% rename from Tests/Vulkan/Shaders/ReadWriteBindings.nzsl rename to Tests/Shaders/Vulkan/ReadWriteBindings.nzsl index 7bbf5ba..76af418 100644 --- a/Tests/Vulkan/Shaders/ReadWriteBindings.nzsl +++ b/Tests/Shaders/Vulkan/ReadWriteBindings.nzsl @@ -21,7 +21,7 @@ external } [entry(compute)] -[workgroup(32, 32, 1)] +[workgroup(16, 16, 1)] fn main(input: Input) { } diff --git a/Tests/Vulkan/Shaders/Simple.nzsl b/Tests/Shaders/Vulkan/Simple.nzsl similarity index 100% rename from Tests/Vulkan/Shaders/Simple.nzsl rename to Tests/Shaders/Vulkan/Simple.nzsl diff --git a/Tests/Vulkan/Shaders/SimpleBufferWrite.nzsl b/Tests/Shaders/Vulkan/SimpleBufferWrite.nzsl similarity index 93% rename from Tests/Vulkan/Shaders/SimpleBufferWrite.nzsl rename to Tests/Shaders/Vulkan/SimpleBufferWrite.nzsl index a5b302d..fb7bdc3 100644 --- a/Tests/Vulkan/Shaders/SimpleBufferWrite.nzsl +++ b/Tests/Shaders/Vulkan/SimpleBufferWrite.nzsl @@ -18,7 +18,7 @@ external } [entry(compute)] -[workgroup(32, 32, 1)] +[workgroup(16, 16, 1)] fn main(input: Input) { ssbo.data[input.indices.x * input.indices.y] = u32(0xFFFFFFFF); diff --git a/Tests/Vulkan/Shaders/WriteOnlyBindings.nzsl b/Tests/Shaders/Vulkan/WriteOnlyBindings.nzsl similarity index 74% rename from Tests/Vulkan/Shaders/WriteOnlyBindings.nzsl rename to Tests/Shaders/Vulkan/WriteOnlyBindings.nzsl index c67a8c4..3284f86 100644 --- a/Tests/Vulkan/Shaders/WriteOnlyBindings.nzsl +++ b/Tests/Shaders/Vulkan/WriteOnlyBindings.nzsl @@ -14,12 +14,12 @@ struct SSBO external { - [set(1), binding(0)] write_texture: texture2D[f32, readonly, rgba8], + [set(1), binding(0)] write_texture: texture2D[f32, readwrite, rgba8], [set(1), binding(1)] write_ssbo: storage[SSBO], } [entry(compute)] -[workgroup(32, 32, 1)] +[workgroup(16, 16, 1)] fn main(input: Input) { } diff --git a/Tests/Shaders/WebGPU/BufferCopy.wgsl b/Tests/Shaders/WebGPU/BufferCopy.wgsl new file mode 100644 index 0000000..a273145 --- /dev/null +++ b/Tests/Shaders/WebGPU/BufferCopy.wgsl @@ -0,0 +1,8 @@ +@group(0) @binding(0) var read_ssbo: array; +@group(1) @binding(0) var write_ssbo: array; + +@compute @workgroup_size(16, 16, 1) +fn main(@builtin(global_invocation_id) grid: vec3) +{ + write_ssbo[grid.x * grid.y] = read_ssbo[grid.x * grid.y]; +} diff --git a/Tests/Shaders/WebGPU/ReadOnlyBindings.wgsl b/Tests/Shaders/WebGPU/ReadOnlyBindings.wgsl new file mode 100644 index 0000000..3495545 --- /dev/null +++ b/Tests/Shaders/WebGPU/ReadOnlyBindings.wgsl @@ -0,0 +1,7 @@ +@group(0) @binding(0) var read_ssbo: array; +@group(0) @binding(1) var read_texture: texture_storage_2d; + +@compute @workgroup_size(16, 16, 1) +fn main(@builtin(global_invocation_id) grid: vec3) +{ +} diff --git a/Tests/Shaders/WebGPU/ReadWriteBindings.wgsl b/Tests/Shaders/WebGPU/ReadWriteBindings.wgsl new file mode 100644 index 0000000..31b1913 --- /dev/null +++ b/Tests/Shaders/WebGPU/ReadWriteBindings.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var read_ssbo: array; +@group(0) @binding(1) var read_texture: texture_storage_2d; +@group(1) @binding(0) var write_ssbo: array; +@group(1) @binding(1) var write_texture: texture_storage_2d; + +@compute @workgroup_size(16, 16, 1) +fn main(@builtin(global_invocation_id) grid: vec3) +{ +} diff --git a/Tests/Shaders/WebGPU/Simple.wgsl b/Tests/Shaders/WebGPU/Simple.wgsl new file mode 100644 index 0000000..0a3e377 --- /dev/null +++ b/Tests/Shaders/WebGPU/Simple.wgsl @@ -0,0 +1,4 @@ +@compute @workgroup_size(16, 16, 1) +fn main(@builtin(global_invocation_id) grid: vec3) +{ +} diff --git a/Tests/Shaders/WebGPU/SimpleBufferWrite.wgsl b/Tests/Shaders/WebGPU/SimpleBufferWrite.wgsl new file mode 100644 index 0000000..716bd1b --- /dev/null +++ b/Tests/Shaders/WebGPU/SimpleBufferWrite.wgsl @@ -0,0 +1,7 @@ +@group(1) @binding(0) var write_ssbo: array; + +@compute @workgroup_size(16, 16, 1) +fn main(@builtin(global_invocation_id) grid: vec3) +{ + write_ssbo[grid.x * grid.y] = u32(0xFFFFFFFF); +} diff --git a/Tests/Shaders/WebGPU/WriteOnlyBindings.wgsl b/Tests/Shaders/WebGPU/WriteOnlyBindings.wgsl new file mode 100644 index 0000000..96decc5 --- /dev/null +++ b/Tests/Shaders/WebGPU/WriteOnlyBindings.wgsl @@ -0,0 +1,7 @@ +@group(1) @binding(0) var write_ssbo: array; +@group(1) @binding(1) var write_texture: texture_storage_2d; + +@compute @workgroup_size(16, 16, 1) +fn main(@builtin(global_invocation_id) grid: vec3) +{ +} diff --git a/Tests/Vulkan/Backend.c b/Tests/Vulkan/Backend.c deleted file mode 100644 index 6f6dbbc..0000000 --- a/Tests/Vulkan/Backend.c +++ /dev/null @@ -1,52 +0,0 @@ -#include "Common.h" - -#include -#include - -void DumbDebugCallBack(PulseDebugMessageSeverity severity, const char* message) -{ - (void)severity; - (void)message; -} - -void TestVulkanSupport() -{ - if(!PulseSupportsBackend(PULSE_BACKEND_VULKAN, PULSE_SHADER_FORMAT_SPIRV_BIT)) - { - TEST_MESSAGE("Vulkan is not supported"); - exit(0); - } -} - -void TestBackendSetup() -{ - PulseBackend backend = PulseLoadBackend(PULSE_BACKEND_VULKAN, PULSE_SHADER_FORMAT_SPIRV_BIT, PULSE_HIGH_DEBUG); - TEST_ASSERT_NOT_EQUAL_MESSAGE(backend, PULSE_NULL_HANDLE, PulseVerbaliseErrorType(PulseGetLastErrorType())); - PulseSetDebugCallback(backend, DumbDebugCallBack); - PulseUnloadBackend(backend); -} - -void TestBackendAnySetup() -{ - PulseBackend backend = PulseLoadBackend(PULSE_BACKEND_ANY, PULSE_SHADER_FORMAT_SPIRV_BIT, PULSE_HIGH_DEBUG); - TEST_ASSERT_NOT_EQUAL_MESSAGE(backend, PULSE_NULL_HANDLE, PulseVerbaliseErrorType(PulseGetLastErrorType())); - TEST_ASSERT_EQUAL(PulseGetBackendType(backend), PULSE_BACKEND_VULKAN); - PulseSetDebugCallback(backend, DumbDebugCallBack); - PulseUnloadBackend(backend); -} - -void TestWrongBackendSetup() -{ - PulseBackend backend = PulseLoadBackend(PULSE_BACKEND_VULKAN, PULSE_SHADER_FORMAT_MSL_BIT, PULSE_HIGH_DEBUG); - TEST_ASSERT_EQUAL(backend, PULSE_NULL_HANDLE); - PulseSetDebugCallback(backend, DumbDebugCallBack); - PulseUnloadBackend(backend); -} - -void TestBackend() -{ - RUN_TEST(TestVulkanSupport); - RUN_TEST(TestBackendSetup); - RUN_TEST(TestBackendAnySetup); - RUN_TEST(TestWrongBackendSetup); -} diff --git a/Tests/Vulkan/Shaders/.gitignore b/Tests/Vulkan/Shaders/.gitignore deleted file mode 100644 index 289a601..0000000 --- a/Tests/Vulkan/Shaders/.gitignore +++ /dev/null @@ -1 +0,0 @@ -*.spv.h diff --git a/Tests/Vulkan/xmake.lua b/Tests/Vulkan/xmake.lua deleted file mode 100644 index c565f47..0000000 --- a/Tests/Vulkan/xmake.lua +++ /dev/null @@ -1,114 +0,0 @@ -option("vulkan-tests", { description = "Build Vulkan tests", default = false }) - -add_repositories("nazara-engine-repo https://github.com/NazaraEngine/xmake-repo") - -add_requires("nzsl >=2023.12.31", { configs = { shared = false, nzslc = true } }) - -if is_cross() then - add_requires("nzsl~host", { kind = "binary", host = true }) -end - --- Yoinked from NZSL xmake repo -rule("find_nzsl") - on_config(function(target) - import("core.project.project") - import("core.tool.toolchain") - import("lib.detect.find_tool") - - local envs - if is_plat("windows") then - local msvc = target:toolchain("msvc") - if msvc and msvc:check() then - envs = msvc:runenvs() - end - elseif is_plat("mingw") then - local mingw = target:toolchain("mingw") - if mingw and mingw:check() then - envs = mingw:runenvs() - end - end - target:data_set("nzsl_envs", envs) - - local nzsl = project.required_package("nzsl~host") or project.required_package("nzsl") - local nzsldir - if nzsl then - nzsldir = path.join(nzsl:installdir(), "bin") - local osenvs = os.getenvs() - envs = envs or {} - for env, values in pairs(nzsl:get("envs")) do - local flatval = path.joinenv(values) - local oldenv = envs[env] or osenvs[env] - if not oldenv or oldenv == "" then - envs[env] = flatval - elseif not oldenv:startswith(flatval) then - envs[env] = flatval .. path.envsep() .. oldenv - end - end - end - - local nzsla = find_tool("nzsla", { version = true, paths = nzsldir, envs = envs }) - local nzslc = find_tool("nzslc", { version = true, paths = nzsldir, envs = envs }) - - target:data_set("nzsla", nzsla) - target:data_set("nzslc", nzslc) - target:data_set("nzsl_runenv", envs) - end) -rule_end() - -rule("compile_shaders") - set_extensions(".nzsl") - add_deps("find_nzsl") - - before_buildcmd_file(function(target, batchcmds, shaderfile, opt) - local outputdir = target:data("nzsl_includedirs") - local nzslc = target:data("nzslc") - local runenvs = target:data("nzsl_runenv") - assert(nzslc, "nzslc not found! please install nzsl package with nzslc enabled") - - batchcmds:show_progress(opt.progress, "${color.build.object}compiling.shader %s", shaderfile) - local argv = { "--compile=spv-header", "--optimize" } - if outputdir then - batchcmds:mkdir(outputdir) - table.insert(argv, "--output=" .. outputdir) - end - - local kind = target:data("plugin.project.kind") or "" - if kind:match("vs") then - table.insert(argv, "--log-format=vs") - end - - table.insert(argv, shaderfile) - - batchcmds:vrunv(nzslc.program, argv, { curdir = ".", envs = runenvs }) - - local outputfile = path.join(outputdir or path.directory(shaderfile), path.basename(shaderfile) .. ".spv.h") - - batchcmds:add_depfiles(shaderfile) - batchcmds:add_depvalues(nzslc.version) - batchcmds:set_depmtime(os.mtime(outputfile)) - batchcmds:set_depcache(target:dependfile(outputfile)) - end) -rule_end() - -if has_config("vulkan-tests") then - set_group("VulkanTests") - add_requires("unity_test") - - if is_plat("linux") then - add_requires("libbacktrace") - end - - target("VulkanUnitTests") - set_kind("binary") - add_deps("pulse_gpu") - add_rules("compile_shaders") - add_files("**.c") - add_files("**.nzsl") - add_packages("unity_test") - if is_plat("linux") then - add_packages("libbacktrace") - set_extension(".x86_64") - add_defines("BACKTRACE") - end - target_end() -end diff --git a/Tests/Vulkan/main.c b/Tests/main.c similarity index 100% rename from Tests/Vulkan/main.c rename to Tests/main.c diff --git a/Tests/xmake.lua b/Tests/xmake.lua new file mode 100644 index 0000000..32efadc --- /dev/null +++ b/Tests/xmake.lua @@ -0,0 +1,148 @@ +local tests = { + Vulkan = { + option = "vulkan", + packages = { "nzsl" }, + global_custom = function() + add_repositories("nazara-engine-repo https://github.com/NazaraEngine/xmake-repo") + add_requires("nzsl >=2023.12.31", { configs = { shared = false, nzslc = true } }) + if is_cross() then + add_requires("nzsl~host", { kind = "binary", host = true }) + end + -- Yoinked from NZSL xmake repo + rule("find_nzsl") + on_config(function(target) + import("core.project.project") + import("core.tool.toolchain") + import("lib.detect.find_tool") + local envs + if is_plat("windows") then + local msvc = target:toolchain("msvc") + if msvc and msvc:check() then + envs = msvc:runenvs() + end + elseif is_plat("mingw") then + local mingw = target:toolchain("mingw") + if mingw and mingw:check() then + envs = mingw:runenvs() + end + end + target:data_set("nzsl_envs", envs) + local nzsl = project.required_package("nzsl~host") or project.required_package("nzsl") + local nzsldir + if nzsl then + nzsldir = path.join(nzsl:installdir(), "bin") + local osenvs = os.getenvs() + envs = envs or {} + for env, values in pairs(nzsl:get("envs")) do + local flatval = path.joinenv(values) + local oldenv = envs[env] or osenvs[env] + if not oldenv or oldenv == "" then + envs[env] = flatval + elseif not oldenv:startswith(flatval) then + envs[env] = flatval .. path.envsep() .. oldenv + end + end + end + local nzsla = find_tool("nzsla", { version = true, paths = nzsldir, envs = envs }) + local nzslc = find_tool("nzslc", { version = true, paths = nzsldir, envs = envs }) + target:data_set("nzsla", nzsla) + target:data_set("nzslc", nzslc) + target:data_set("nzsl_runenv", envs) + end) + rule_end() + rule("nzsl_compile_shaders") + set_extensions(".nzsl") + add_deps("find_nzsl") + before_buildcmd_file(function(target, batchcmds, shaderfile, opt) + local outputdir = target:data("nzsl_includedirs") + local nzslc = target:data("nzslc") + local runenvs = target:data("nzsl_runenv") + assert(nzslc, "nzslc not found! please install nzsl package with nzslc enabled") + batchcmds:show_progress(opt.progress, "${color.build.object}compiling.shader %s", shaderfile) + local argv = { "--compile=spv-header", "--optimize" } + if outputdir then + batchcmds:mkdir(outputdir) + table.insert(argv, "--output=" .. outputdir) + end + local kind = target:data("plugin.project.kind") or "" + if kind:match("vs") then + table.insert(argv, "--log-format=vs") + end + table.insert(argv, shaderfile) + batchcmds:vrunv(nzslc.program, argv, { curdir = ".", envs = runenvs }) + local outputfile = path.join(outputdir or path.directory(shaderfile), path.basename(shaderfile) .. ".spv.h") + batchcmds:add_depfiles(shaderfile) + batchcmds:add_depvalues(nzslc.version) + batchcmds:set_depmtime(os.mtime(outputfile)) + batchcmds:set_depcache(target:dependfile(outputfile)) + end) + rule_end() + end, + custom = function() + add_rules("nzsl_compile_shaders") + add_files("**.nzsl") + end + }, + WebGPU = { + option = "webgpu", + global_custom = function() + rule("wgsl_compile_shaders") + set_extensions(".wgsl") + before_buildcmd_file(function(target, batchcmds, shaderfile, opt) + batchcmds:show_progress(opt.progress, "${color.build.object}compiling.shader %s", shaderfile) + io.writefile(shaderfile .. ".h", [[ +// Generated File +#undef WGSL_SOURCE +#define WGSL_SOURCE(...) #__VA_ARGS__ +const uint8_t SHADER_NAME[] = WGSL_SOURCE( + +]] .. io.readfile(shaderfile) .. [[ + +);]]) + end) + rule_end() + end, + custom = function() + add_rules("wgsl_compile_shaders") + add_files("**.wgsl") + end + } +} + +if is_plat("linux") then + add_requires("libbacktrace") +end +add_requires("unity_test") + +for name, module in table.orderpairs(tests) do + if module.option then + option(module.option .. "-tests", { description = "Enables " .. name .. " tests", default = false }) + end +end + +for name, module in pairs(tests) do + if has_config(module.option) then + if module.global_custom then + module.global_custom() + end + if module.packages then + add_requires(table.unpack(module.packages)) + end + + target(name .. "UnitTests") + set_kind("binary") + add_deps("pulse_gpu") + add_packages("unity_test") + add_files("**.c") + add_defines(string.upper(name) .. "_ENABLED") + if module.custom then + module.custom() + end + if is_plat("linux") then + add_packages("libbacktrace") + set_extension(".x86_64") + add_defines("BACKTRACE") + end + target_end() + end +end diff --git a/xmake.lua b/xmake.lua index a88b46c..39bb6bd 100644 --- a/xmake.lua +++ b/xmake.lua @@ -137,8 +137,5 @@ target("pulse_gpu") end) target_end() +includes("Tests/xmake.lua") includes("Examples/*.lua") - -if not is_plat("wasm") then - includes("Tests/Vulkan/*.lua") -end