From bb7b6e716af0b48865848a12e81df493396b8fd2 Mon Sep 17 00:00:00 2001 From: Kbz-8 Date: Wed, 26 Feb 2025 15:01:36 +0100 Subject: [PATCH] implementing webgpu pipeline, fixing vulkan compute pass --- Examples/Vulkan/main.c | 2 +- Examples/WebGPU/main.c | 21 ++++++++++++++++ Examples/WebGPU/shader.wgsl | 4 ---- Sources/Backends/Vulkan/VulkanComputePass.c | 7 ++++++ Sources/Backends/WebGPU/WebGPUComputePass.c | 4 ++++ .../Backends/WebGPU/WebGPUComputePipeline.c | 24 +++++++++++++++++-- Sources/PulseComputePass.c | 2 ++ Tests/Vulkan/Image.c | 8 +++++++ 8 files changed, 65 insertions(+), 7 deletions(-) delete mode 100644 Examples/WebGPU/shader.wgsl diff --git a/Examples/Vulkan/main.c b/Examples/Vulkan/main.c index e296cc8..03a5d32 100644 --- a/Examples/Vulkan/main.c +++ b/Examples/Vulkan/main.c @@ -6,7 +6,7 @@ #define CHECK_PULSE_HANDLE_RETVAL(handle, retval) \ if(handle == PULSE_NULL_HANDLE) \ { \ - fprintf(stderr, "Error: %s", PulseVerbaliseErrorType(PulseGetLastErrorType())); \ + fprintf(stderr, "Error: %s\n", PulseVerbaliseErrorType(PulseGetLastErrorType())); \ return retval; \ } \ diff --git a/Examples/WebGPU/main.c b/Examples/WebGPU/main.c index eebf53c..33b16b2 100644 --- a/Examples/WebGPU/main.c +++ b/Examples/WebGPU/main.c @@ -2,6 +2,7 @@ #include #include +#include #define CHECK_PULSE_HANDLE_RETVAL(handle, retval) \ if(handle == PULSE_NULL_HANDLE) \ @@ -23,6 +24,15 @@ void DebugCallBack(PulseDebugMessageSeverity severity, const char* message) printf("Pulse: %s\n", message); } +#define WGSL_SOURCE(...) #__VA_ARGS__ + +const char* wgsl_source = WGSL_SOURCE( + @compute @workgroup_size(32, 32, 1) + fn main(@builtin(global_invocation_id) grid: vec3u) + { + } +); + int main(void) { PulseBackend backend = PulseLoadBackend(PULSE_BACKEND_WEBGPU, PULSE_SHADER_FORMAT_WGSL_BIT, PULSE_HIGH_DEBUG); @@ -31,6 +41,17 @@ int main(void) PulseDevice device = PulseCreateDevice(backend, NULL, 0); CHECK_PULSE_HANDLE_RETVAL(device, 1); + PulseComputePipelineCreateInfo info = { 0 }; + info.code_size = strlen(wgsl_source); + info.code = (const uint8_t*)wgsl_source; + info.entrypoint = "main"; + info.format = PULSE_SHADER_FORMAT_WGSL_BIT; + + PulseComputePipeline pipeline = PulseCreateComputePipeline(device, &info); + CHECK_PULSE_HANDLE_RETVAL(pipeline, 1); + + PulseDestroyComputePipeline(device, pipeline); + PulseDestroyDevice(device); PulseUnloadBackend(backend); puts("Successfully executed Pulse example using WebGPU !"); diff --git a/Examples/WebGPU/shader.wgsl b/Examples/WebGPU/shader.wgsl deleted file mode 100644 index 95cfe24..0000000 --- a/Examples/WebGPU/shader.wgsl +++ /dev/null @@ -1,4 +0,0 @@ -@compute @workgroup_size(32, 32, 1) -fn main(@builtin(global_invocation_id) grid: vec3u) -{ -} diff --git a/Sources/Backends/Vulkan/VulkanComputePass.c b/Sources/Backends/Vulkan/VulkanComputePass.c index 695cdab..5e23270 100644 --- a/Sources/Backends/Vulkan/VulkanComputePass.c +++ b/Sources/Backends/Vulkan/VulkanComputePass.c @@ -134,4 +134,11 @@ PulseComputePass VulkanBeginComputePass(PulseCommandList cmd) void VulkanEndComputePass(PulseComputePass pass) { + VulkanComputePass* vulkan_pass = VULKAN_RETRIEVE_DRIVER_DATA_AS(pass, VulkanComputePass*); + VulkanReturnDescriptorSetToPool(vulkan_pass->read_only_descriptor_set->pool, vulkan_pass->read_only_descriptor_set); + VulkanReturnDescriptorSetToPool(vulkan_pass->read_write_descriptor_set->pool, vulkan_pass->read_write_descriptor_set); + VulkanReturnDescriptorSetToPool(vulkan_pass->uniform_descriptor_set->pool, vulkan_pass->uniform_descriptor_set); + vulkan_pass->read_only_descriptor_set = VK_NULL_HANDLE; + vulkan_pass->read_write_descriptor_set = VK_NULL_HANDLE; + vulkan_pass->uniform_descriptor_set = VK_NULL_HANDLE; } diff --git a/Sources/Backends/WebGPU/WebGPUComputePass.c b/Sources/Backends/WebGPU/WebGPUComputePass.c index 2a1ea79..a16adcb 100644 --- a/Sources/Backends/WebGPU/WebGPUComputePass.c +++ b/Sources/Backends/WebGPU/WebGPUComputePass.c @@ -6,6 +6,7 @@ #include "WebGPU.h" #include "WebGPUDevice.h" #include "WebGPUComputePass.h" +#include "WebGPUComputePipeline.h" PulseComputePass WebGPUCreateComputePass(PulseDevice device, PulseCommandList cmd) { @@ -58,6 +59,9 @@ void WebGPUBindStorageImages(PulseComputePass pass, const PulseImage* images, ui void WebGPUBindComputePipeline(PulseComputePass pass, PulseComputePipeline pipeline) { + WebGPUComputePass* webgpu_pass = WEBGPU_RETRIEVE_DRIVER_DATA_AS(pass, WebGPUComputePass*); + WebGPUComputePipeline* webgpu_pipeline = WEBGPU_RETRIEVE_DRIVER_DATA_AS(pipeline, WebGPUComputePipeline*); + wgpuComputePassEncoderSetPipeline(webgpu_pass->encoder, webgpu_pipeline->pipeline); } void WebGPUDispatchComputations(PulseComputePass pass, uint32_t groupcount_x, uint32_t groupcount_y, uint32_t groupcount_z) diff --git a/Sources/Backends/WebGPU/WebGPUComputePipeline.c b/Sources/Backends/WebGPU/WebGPUComputePipeline.c index cd351a3..d4db46a 100644 --- a/Sources/Backends/WebGPU/WebGPUComputePipeline.c +++ b/Sources/Backends/WebGPU/WebGPUComputePipeline.c @@ -44,8 +44,8 @@ PulseComputePipeline WebGPUCreateComputePipeline(PulseDevice device, const Pulse webgpu_pipeline->shader = wgpuDeviceCreateShaderModule(webgpu_device->device, &shader_descriptor); WGPUStringView entrypoint = { 0 }; - code.length = WGPU_STRLEN; - code.data = info->entrypoint; + entrypoint.length = WGPU_STRLEN; + entrypoint.data = info->entrypoint; WGPUProgrammableStageDescriptor state = { 0 }; state.module = webgpu_pipeline->shader; state.entryPoint = entrypoint; @@ -53,9 +53,29 @@ PulseComputePipeline WebGPUCreateComputePipeline(PulseDevice device, const Pulse pipeline_descriptor.compute = state; webgpu_pipeline->pipeline = wgpuDeviceCreateComputePipeline(webgpu_device->device, &pipeline_descriptor); + if(PULSE_IS_BACKEND_HIGH_LEVEL_DEBUG(device->backend)) + PulseLogInfoFmt(device->backend, "(WebGPU) created new compute pipeline %p", pipeline); + return pipeline; } void WebGPUDestroyComputePipeline(PulseDevice device, PulseComputePipeline pipeline) { + if(pipeline == PULSE_NULL_HANDLE) + { + if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(device->backend)) + PulseLogWarning(device->backend, "compute pipeline is NULL, this may be a bug in your application"); + return; + } + + PULSE_UNUSED(device); + WebGPUComputePipeline* webgpu_pipeline = WEBGPU_RETRIEVE_DRIVER_DATA_AS(pipeline, WebGPUComputePipeline*); + wgpuComputePipelineRelease(webgpu_pipeline->pipeline); + wgpuShaderModuleRelease(webgpu_pipeline->shader); + free(webgpu_pipeline); + + if(PULSE_IS_BACKEND_HIGH_LEVEL_DEBUG(device->backend)) + PulseLogInfoFmt(device->backend, "(WebGPU) destroyed compute pipeline %p", pipeline); + + free(pipeline); } diff --git a/Sources/PulseComputePass.c b/Sources/PulseComputePass.c index daf7671..9b633d4 100644 --- a/Sources/PulseComputePass.c +++ b/Sources/PulseComputePass.c @@ -99,6 +99,8 @@ PULSE_API void PulseEndComputePass(PulseComputePass pass) memset(pass->readonly_storage_buffers, 0, PULSE_MAX_READ_BUFFERS_BOUND * sizeof(PulseBuffer)); memset(pass->readwrite_storage_buffers, 0, PULSE_MAX_WRITE_BUFFERS_BOUND * sizeof(PulseBuffer)); + pass->cmd->device->PFN_EndComputePass(pass); + pass->current_pipeline = PULSE_NULL_HANDLE; pass->is_recording = false; diff --git a/Tests/Vulkan/Image.c b/Tests/Vulkan/Image.c index 367827f..cdfbbc5 100644 --- a/Tests/Vulkan/Image.c +++ b/Tests/Vulkan/Image.c @@ -34,6 +34,12 @@ void TestImageCreation() TEST_ASSERT_NOT_EQUAL_MESSAGE(image, PULSE_NULL_HANDLE, PulseVerbaliseErrorType(PulseGetLastErrorType())); PulseDestroyImage(device, image); } + + /** + * This test may crash some Nouveau NVK drivers (wtf ???). + * It seems to be comming exclusively from 3D read-only images + */ + if(false) { PulseImageCreateInfo image_create_info = { 0 }; image_create_info.type = PULSE_IMAGE_TYPE_3D; @@ -46,6 +52,8 @@ void TestImageCreation() TEST_ASSERT_NOT_EQUAL_MESSAGE(image, PULSE_NULL_HANDLE, PulseVerbaliseErrorType(PulseGetLastErrorType())); PulseDestroyImage(device, image); } + + { PulseImageCreateInfo image_create_info = { 0 }; image_create_info.type = PULSE_IMAGE_TYPE_CUBE;