From bb287958bd1f67d9c58e75062a671be0668e095a Mon Sep 17 00:00:00 2001 From: Kbz-8 Date: Sat, 1 Mar 2025 11:42:28 +0100 Subject: [PATCH] yes --- .github/workflows/msys2-build.yml | 2 +- Examples/Vulkan/main.c | 114 +++++++++------- Examples/WebGPU/main.c | 97 +++++++++----- Examples/xmake.lua | 6 +- Includes/Pulse.h | 10 +- Sources/Backends/Vulkan/VulkanBuffer.c | 7 +- Sources/Backends/Vulkan/VulkanBuffer.h | 2 +- Sources/Backends/Vulkan/VulkanCommandList.c | 7 +- Sources/Backends/Vulkan/VulkanFence.c | 5 +- Sources/Backends/WebGPU/WebGPU.c | 2 + Sources/Backends/WebGPU/WebGPU.h | 1 - Sources/Backends/WebGPU/WebGPUBuffer.c | 137 +++++++++++++++++++- Sources/Backends/WebGPU/WebGPUBuffer.h | 6 +- Sources/Backends/WebGPU/WebGPUCommandList.c | 28 ++-- Sources/Backends/WebGPU/WebGPUComputePass.c | 1 + Sources/Backends/WebGPU/WebGPUDevice.c | 8 +- Sources/Backends/WebGPU/WebGPUFence.c | 4 +- Sources/PulseBuffer.c | 55 +++++++- Sources/PulseDefs.h | 40 +++--- Sources/PulseFence.c | 8 +- Sources/PulsePFNs.h | 2 +- Tests/Vulkan/Buffer.c | 22 ++-- 22 files changed, 414 insertions(+), 150 deletions(-) diff --git a/.github/workflows/msys2-build.yml b/.github/workflows/msys2-build.yml index 56d391c..2a9ed38 100644 --- a/.github/workflows/msys2-build.yml +++ b/.github/workflows/msys2-build.yml @@ -66,7 +66,7 @@ jobs: # Setup compilation mode and install project dependencies - name: Configure xmake and install dependencies - run: xmake config --arch=${{ matrix.arch }} --mode=${{ matrix.confs.mode }} ${{ matrix.confs.config }} --ccache=n --unitybuild=y --yes + run: xmake config --arch=${{ matrix.arch }} --mode=${{ matrix.confs.mode }} ${{ matrix.confs.config }} --ccache=n --unitybuild=y --webgpu=n --yes # Save dependencies - name: Save cached xmake dependencies diff --git a/Examples/Vulkan/main.c b/Examples/Vulkan/main.c index 03a5d32..97204d2 100644 --- a/Examples/Vulkan/main.c +++ b/Examples/Vulkan/main.c @@ -3,13 +3,6 @@ #include #include -#define CHECK_PULSE_HANDLE_RETVAL(handle, retval) \ - if(handle == PULSE_NULL_HANDLE) \ - { \ - fprintf(stderr, "Error: %s\n", PulseVerbaliseErrorType(PulseGetLastErrorType())); \ - return retval; \ - } \ - void DebugCallBack(PulseDebugMessageSeverity severity, const char* message) { if(severity == PULSE_DEBUG_MESSAGE_SEVERITY_ERROR) @@ -23,62 +16,85 @@ void DebugCallBack(PulseDebugMessageSeverity severity, const char* message) printf("Pulse: %s\n", message); } +#define BUFFER_SIZE (256 * sizeof(uint32_t)) + int main(void) { PulseBackend backend = PulseLoadBackend(PULSE_BACKEND_VULKAN, PULSE_SHADER_FORMAT_SPIRV_BIT, PULSE_HIGH_DEBUG); - CHECK_PULSE_HANDLE_RETVAL(backend, 1); PulseSetDebugCallback(backend, DebugCallBack); PulseDevice device = PulseCreateDevice(backend, NULL, 0); - CHECK_PULSE_HANDLE_RETVAL(device, 1); - - const uint8_t shader_bytecode[] = { - #include "shader.spv.h" - }; - - PulseComputePipelineCreateInfo info = { 0 }; - info.code_size = sizeof(shader_bytecode); - info.code = shader_bytecode; - info.entrypoint = "main"; - info.format = PULSE_SHADER_FORMAT_SPIRV_BIT; - info.num_readwrite_storage_buffers = 1; - - PulseComputePipeline pipeline = PulseCreateComputePipeline(device, &info); - CHECK_PULSE_HANDLE_RETVAL(pipeline, 1); PulseBufferCreateInfo buffer_create_info = { 0 }; - buffer_create_info.size = 256 * sizeof(uint32_t); + buffer_create_info.size = BUFFER_SIZE; buffer_create_info.usage = PULSE_BUFFER_USAGE_STORAGE_READ | PULSE_BUFFER_USAGE_STORAGE_WRITE | PULSE_BUFFER_USAGE_TRANSFER_DOWNLOAD; - PulseBuffer buffer = PulseCreateBuffer(device, &buffer_create_info); - CHECK_PULSE_HANDLE_RETVAL(buffer, 1); - PulseFence fence = PulseCreateFence(device); - CHECK_PULSE_HANDLE_RETVAL(fence, 1); - PulseCommandList cmd = PulseRequestCommandList(device, PULSE_COMMAND_LIST_GENERAL); - CHECK_PULSE_HANDLE_RETVAL(cmd, 1); + // GPU computations + { + const uint8_t shader_bytecode[] = { + #include "shader.spv.h" + }; - PulseComputePass pass = PulseBeginComputePass(cmd); - CHECK_PULSE_HANDLE_RETVAL(pass, 1); - PulseBindStorageBuffers(pass, &buffer, 1); - PulseBindComputePipeline(pass, pipeline); - PulseDispatchComputations(pass, 32, 32, 1); - PulseEndComputePass(pass); + PulseComputePipelineCreateInfo info = { 0 }; + info.code_size = sizeof(shader_bytecode); + info.code = shader_bytecode; + info.entrypoint = "main"; + info.format = PULSE_SHADER_FORMAT_SPIRV_BIT; + info.num_readwrite_storage_buffers = 1; + PulseComputePipeline pipeline = PulseCreateComputePipeline(device, &info); - if(!PulseSubmitCommandList(device, cmd, fence)) - fprintf(stderr, "Could not submit command list, %s\n", PulseVerbaliseErrorType(PulseGetLastErrorType())); - if(!PulseWaitForFences(device, &fence, 1, true)) - fprintf(stderr, "Could not wait for fences, %s\n", PulseVerbaliseErrorType(PulseGetLastErrorType())); + PulseFence fence = PulseCreateFence(device); + PulseCommandList cmd = PulseRequestCommandList(device, PULSE_COMMAND_LIST_GENERAL); - void* ptr; - PulseMapBuffer(buffer, &ptr); - for(uint32_t i = 0; i < 256; i++) - printf("%d, ", ((int32_t*)ptr)[i]); - puts(""); - PulseUnmapBuffer(buffer); + PulseComputePass pass = PulseBeginComputePass(cmd); + PulseBindStorageBuffers(pass, &buffer, 1); + PulseBindComputePipeline(pass, pipeline); + PulseDispatchComputations(pass, 32, 32, 1); + PulseEndComputePass(pass); - PulseReleaseCommandList(device, cmd); - PulseDestroyFence(device, fence); - PulseDestroyComputePipeline(device, pipeline); + PulseSubmitCommandList(device, cmd, fence); + PulseWaitForFences(device, &fence, 1, true); + + PulseReleaseCommandList(device, cmd); + PulseDestroyFence(device, fence); + PulseDestroyComputePipeline(device, pipeline); + } + + // Get result and read it on CPU + { + PulseBufferCreateInfo staging_buffer_create_info = { 0 }; + staging_buffer_create_info.size = BUFFER_SIZE; + staging_buffer_create_info.usage = PULSE_BUFFER_USAGE_TRANSFER_UPLOAD | PULSE_BUFFER_USAGE_TRANSFER_DOWNLOAD; + PulseBuffer staging_buffer = PulseCreateBuffer(device, &staging_buffer_create_info); + + PulseFence fence = PulseCreateFence(device); + PulseCommandList cmd = PulseRequestCommandList(device, PULSE_COMMAND_LIST_TRANSFER_ONLY); + + PulseBufferRegion src_region = { 0 }; + src_region.buffer = buffer; + src_region.size = BUFFER_SIZE; + + PulseBufferRegion dst_region = { 0 }; + dst_region.buffer = staging_buffer; + dst_region.size = BUFFER_SIZE; + + PulseCopyBufferToBuffer(cmd, &src_region, &dst_region); + + PulseSubmitCommandList(device, cmd, fence); + PulseWaitForFences(device, &fence, 1, true); + + void* ptr; + PulseMapBuffer(staging_buffer, PULSE_MAP_READ, &ptr); + for(uint32_t i = 0; i < BUFFER_SIZE / sizeof(uint32_t); i++) + printf("%d, ", ((int32_t*)ptr)[i]); + puts(""); + PulseUnmapBuffer(staging_buffer); + + PulseDestroyBuffer(device, staging_buffer); + + PulseReleaseCommandList(device, cmd); + PulseDestroyFence(device, fence); + } PulseDestroyBuffer(device, buffer); diff --git a/Examples/WebGPU/main.c b/Examples/WebGPU/main.c index ae340fa..d70ae03 100644 --- a/Examples/WebGPU/main.c +++ b/Examples/WebGPU/main.c @@ -5,12 +5,6 @@ #include #define WGSL_SOURCE(...) #__VA_ARGS__ -#define CHECK_PULSE_HANDLE_RETVAL(handle, retval) \ - if(handle == PULSE_NULL_HANDLE) \ - { \ - fprintf(stderr, "Error: '" #handle "' %s\n", PulseVerbaliseErrorType(PulseGetLastErrorType())); \ - return retval; \ - } \ void DebugCallBack(PulseDebugMessageSeverity severity, const char* message) { @@ -25,6 +19,8 @@ void DebugCallBack(PulseDebugMessageSeverity severity, const char* message) printf("Pulse: %s\n", message); } +#define BUFFER_SIZE (256 * sizeof(uint32_t)) + const char* wgsl_source = WGSL_SOURCE( @compute @workgroup_size(32, 32, 1) fn main(@builtin(global_invocation_id) grid: vec3u) @@ -35,39 +31,78 @@ const char* wgsl_source = WGSL_SOURCE( int main(void) { PulseBackend backend = PulseLoadBackend(PULSE_BACKEND_WEBGPU, PULSE_SHADER_FORMAT_WGSL_BIT, PULSE_HIGH_DEBUG); - CHECK_PULSE_HANDLE_RETVAL(backend, 1); PulseSetDebugCallback(backend, DebugCallBack); 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; + PulseBufferCreateInfo buffer_create_info = { 0 }; + buffer_create_info.size = BUFFER_SIZE; + buffer_create_info.usage = PULSE_BUFFER_USAGE_STORAGE_READ | PULSE_BUFFER_USAGE_STORAGE_WRITE | PULSE_BUFFER_USAGE_TRANSFER_DOWNLOAD; + PulseBuffer buffer = PulseCreateBuffer(device, &buffer_create_info); - PulseComputePipeline pipeline = PulseCreateComputePipeline(device, &info); - CHECK_PULSE_HANDLE_RETVAL(pipeline, 1); + // GPU computations + { + 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; + info.num_readwrite_storage_buffers = 1; + PulseComputePipeline pipeline = PulseCreateComputePipeline(device, &info); - PulseFence fence = PulseCreateFence(device); - CHECK_PULSE_HANDLE_RETVAL(fence, 1); - PulseCommandList cmd = PulseRequestCommandList(device, PULSE_COMMAND_LIST_GENERAL); - CHECK_PULSE_HANDLE_RETVAL(cmd, 1); + PulseFence fence = PulseCreateFence(device); + PulseCommandList cmd = PulseRequestCommandList(device, PULSE_COMMAND_LIST_GENERAL); - PulseComputePass pass = PulseBeginComputePass(cmd); - CHECK_PULSE_HANDLE_RETVAL(pass, 1); - PulseBindComputePipeline(pass, pipeline); - PulseDispatchComputations(pass, 32, 32, 1); - PulseEndComputePass(pass); + PulseComputePass pass = PulseBeginComputePass(cmd); + // PulseBindStorageBuffers(pass, &buffer, 1); + PulseBindComputePipeline(pass, pipeline); + PulseDispatchComputations(pass, 32, 32, 1); + PulseEndComputePass(pass); - if(!PulseSubmitCommandList(device, cmd, fence)) - fprintf(stderr, "Could not submit command list, %s\n", PulseVerbaliseErrorType(PulseGetLastErrorType())); - if(!PulseWaitForFences(device, &fence, 1, true)) - fprintf(stderr, "Could not wait for fences, %s\n", PulseVerbaliseErrorType(PulseGetLastErrorType())); + PulseSubmitCommandList(device, cmd, fence); + PulseWaitForFences(device, &fence, 1, true); - PulseReleaseCommandList(device, cmd); - PulseDestroyFence(device, fence); - PulseDestroyComputePipeline(device, pipeline); + PulseReleaseCommandList(device, cmd); + PulseDestroyFence(device, fence); + PulseDestroyComputePipeline(device, pipeline); + } + + // Get result and read it on CPU + { + PulseBufferCreateInfo staging_buffer_create_info = { 0 }; + staging_buffer_create_info.size = BUFFER_SIZE; + staging_buffer_create_info.usage = PULSE_BUFFER_USAGE_TRANSFER_DOWNLOAD; + PulseBuffer staging_buffer = PulseCreateBuffer(device, &staging_buffer_create_info); + + PulseFence fence = PulseCreateFence(device); + PulseCommandList cmd = PulseRequestCommandList(device, PULSE_COMMAND_LIST_TRANSFER_ONLY); + + PulseBufferRegion src_region = { 0 }; + src_region.buffer = buffer; + src_region.size = BUFFER_SIZE; + + PulseBufferRegion dst_region = { 0 }; + dst_region.buffer = staging_buffer; + dst_region.size = BUFFER_SIZE; + + PulseCopyBufferToBuffer(cmd, &src_region, &dst_region); + + PulseSubmitCommandList(device, cmd, fence); + PulseWaitForFences(device, &fence, 1, true); + + void* ptr; + PulseMapBuffer(staging_buffer, PULSE_MAP_READ, &ptr); + for(uint32_t i = 0; i < BUFFER_SIZE / sizeof(uint32_t); i++) + printf("%d, ", ((int32_t*)ptr)[i]); + puts(""); + PulseUnmapBuffer(staging_buffer); + + PulseDestroyBuffer(device, staging_buffer); + + PulseReleaseCommandList(device, cmd); + PulseDestroyFence(device, fence); + } + + PulseDestroyBuffer(device, buffer); PulseDestroyDevice(device); PulseUnloadBackend(backend); diff --git a/Examples/xmake.lua b/Examples/xmake.lua index f527060..7ee7673 100644 --- a/Examples/xmake.lua +++ b/Examples/xmake.lua @@ -2,8 +2,10 @@ option("examples", { description = "Build the examples", default = false }) if has_config("examples") then set_group("Examples") - if not is_plat("wasm") then + if not is_plat("wasm") and has_config("vulkan") then includes("Vulkan/xmake.lua") end - includes("WebGPU/xmake.lua") + if has_config("webgpu") then + includes("WebGPU/xmake.lua") + end end diff --git a/Includes/Pulse.h b/Includes/Pulse.h index 09b89c1..32404dd 100644 --- a/Includes/Pulse.h +++ b/Includes/Pulse.h @@ -127,7 +127,7 @@ typedef enum PulseImageType typedef enum PulseImageFormat { - PULSE_IMAGE_FORMAT_INVALID, + PULSE_IMAGE_FORMAT_INVALID = 0, // Unsigned Normalized Float Color Formats PULSE_IMAGE_FORMAT_A8_UNORM, PULSE_IMAGE_FORMAT_R8_UNORM, @@ -192,6 +192,12 @@ typedef enum PulseImageFormat PULSE_IMAGE_FORMAT_MAX_ENUM // For internal use only } PulseImageFormat; +typedef enum PulseMapMode +{ + PULSE_MAP_READ, + PULSE_MAP_WRITE, +} PulseMapMode; + // Structs typedef struct PulseBufferCreateInfo { @@ -257,7 +263,7 @@ PULSE_API bool PulseDeviceSupportsShaderFormats(PulseDevice device, PulseShaderF PULSE_API void PulseDestroyDevice(PulseDevice device); PULSE_API PulseBuffer PulseCreateBuffer(PulseDevice device, const PulseBufferCreateInfo* create_infos); -PULSE_API bool PulseMapBuffer(PulseBuffer buffer, void** data); +PULSE_API bool PulseMapBuffer(PulseBuffer buffer, PulseMapMode mode, void** data); PULSE_API void PulseUnmapBuffer(PulseBuffer buffer); PULSE_API bool PulseCopyBufferToBuffer(PulseCommandList cmd, const PulseBufferRegion* src, const PulseBufferRegion* dst); PULSE_API bool PulseCopyBufferToImage(PulseCommandList cmd, const PulseBufferRegion* src, const PulseImageRegion* dst); diff --git a/Sources/Backends/Vulkan/VulkanBuffer.c b/Sources/Backends/Vulkan/VulkanBuffer.c index f69442b..c3b6577 100644 --- a/Sources/Backends/Vulkan/VulkanBuffer.c +++ b/Sources/Backends/Vulkan/VulkanBuffer.c @@ -29,12 +29,12 @@ PulseBuffer VulkanCreateBuffer(PulseDevice device, const PulseBufferCreateInfo* if(buffer->usage & PULSE_BUFFER_USAGE_TRANSFER_UPLOAD) { - vulkan_buffer->usage |= VK_BUFFER_USAGE_TRANSFER_SRC_BIT; + vulkan_buffer->usage |= VK_BUFFER_USAGE_TRANSFER_DST_BIT; allocation_create_info.flags = VMA_ALLOCATION_CREATE_HOST_ACCESS_SEQUENTIAL_WRITE_BIT; } if(buffer->usage & PULSE_BUFFER_USAGE_TRANSFER_DOWNLOAD) { - vulkan_buffer->usage |= VK_BUFFER_USAGE_TRANSFER_DST_BIT; + vulkan_buffer->usage |= VK_BUFFER_USAGE_TRANSFER_SRC_BIT; allocation_create_info.flags = VMA_ALLOCATION_CREATE_HOST_ACCESS_SEQUENTIAL_WRITE_BIT; } if(buffer->usage & PULSE_INTERNAL_BUFFER_USAGE_UNIFORM_ACCESS) @@ -57,8 +57,9 @@ PulseBuffer VulkanCreateBuffer(PulseDevice device, const PulseBufferCreateInfo* return buffer; } -bool VulkanMapBuffer(PulseBuffer buffer, void** data) +bool VulkanMapBuffer(PulseBuffer buffer, PulseMapMode mode, void** data) { + PULSE_UNUSED(mode); VulkanBuffer* vulkan_buffer = VULKAN_RETRIEVE_DRIVER_DATA_AS(buffer, VulkanBuffer*); VulkanDevice* vulkan_device = VULKAN_RETRIEVE_DRIVER_DATA_AS(buffer->device, VulkanDevice*); CHECK_VK_RETVAL(buffer->device->backend, vmaMapMemory(vulkan_device->allocator, vulkan_buffer->allocation, data), PULSE_ERROR_MAP_FAILED, false); diff --git a/Sources/Backends/Vulkan/VulkanBuffer.h b/Sources/Backends/Vulkan/VulkanBuffer.h index e0cfca6..19c1a36 100644 --- a/Sources/Backends/Vulkan/VulkanBuffer.h +++ b/Sources/Backends/Vulkan/VulkanBuffer.h @@ -23,7 +23,7 @@ typedef struct VulkanBuffer } VulkanBuffer; PulseBuffer VulkanCreateBuffer(PulseDevice device, const PulseBufferCreateInfo* create_infos); -bool VulkanMapBuffer(PulseBuffer buffer, void** data); +bool VulkanMapBuffer(PulseBuffer buffer, PulseMapMode mode, void** data); void VulkanUnmapBuffer(PulseBuffer buffer); bool VulkanCopyBufferToBuffer(PulseCommandList cmd, const PulseBufferRegion* src, const PulseBufferRegion* dst); bool VulkanCopyBufferToImage(PulseCommandList cmd, const PulseBufferRegion* src, const PulseImageRegion* dst); diff --git a/Sources/Backends/Vulkan/VulkanCommandList.c b/Sources/Backends/Vulkan/VulkanCommandList.c index 4566f1f..46aa1b4 100644 --- a/Sources/Backends/Vulkan/VulkanCommandList.c +++ b/Sources/Backends/Vulkan/VulkanCommandList.c @@ -109,7 +109,7 @@ bool VulkanSubmitCommandList(PulseDevice device, PulseCommandList cmd, PulseFenc default: break; } - VkFence vulkan_fence; + VkFence vulkan_fence = VK_NULL_HANDLE; if(fence != PULSE_NULL_HANDLE) { vulkan_fence = VULKAN_RETRIEVE_DRIVER_DATA_AS(fence, VkFence); @@ -132,7 +132,10 @@ bool VulkanSubmitCommandList(PulseDevice device, PulseCommandList cmd, PulseFenc submit_info.commandBufferCount = 1; submit_info.pCommandBuffers = &vulkan_cmd->cmd; res = vulkan_device->vkQueueSubmit(vulkan_queue->queue, 1, &submit_info, vulkan_fence); - cmd->state = PULSE_COMMAND_LIST_STATE_SENT; + if(fence != PULSE_NULL_HANDLE) + cmd->state = PULSE_COMMAND_LIST_STATE_SENT; + else + cmd->state = PULSE_COMMAND_LIST_STATE_READY; switch(res) { case VK_SUCCESS: return true; diff --git a/Sources/Backends/Vulkan/VulkanFence.c b/Sources/Backends/Vulkan/VulkanFence.c index 1c83f1d..5d3f269 100644 --- a/Sources/Backends/Vulkan/VulkanFence.c +++ b/Sources/Backends/Vulkan/VulkanFence.c @@ -93,10 +93,7 @@ bool VulkanWaitForFences(PulseDevice device, const PulseFence* fences, uint32_t free(vulkan_fences); switch(result) { - case VK_SUCCESS: - for(uint32_t i = 0; i < fences_count; i++) - fences[i]->cmd->state = PULSE_COMMAND_LIST_STATE_READY; - break; + case VK_SUCCESS: break; case VK_TIMEOUT: break; case VK_ERROR_DEVICE_LOST: PulseSetInternalError(PULSE_ERROR_DEVICE_LOST); return false; diff --git a/Sources/Backends/WebGPU/WebGPU.c b/Sources/Backends/WebGPU/WebGPU.c index 510c655..4ff778b 100644 --- a/Sources/Backends/WebGPU/WebGPU.c +++ b/Sources/Backends/WebGPU/WebGPU.c @@ -24,6 +24,8 @@ PulseBackendFlags WebGPUCheckSupport(PulseBackendFlags candidates, PulseShaderFo bool WebGPULoadBackend(PulseBackend backend, PulseDebugLevel debug_level) { + PULSE_UNUSED(backend); + PULSE_UNUSED(debug_level); WebGPUDriverData* driver_data = (WebGPUDriverData*)calloc(1, sizeof(WebGPUDriverData)); PULSE_CHECK_ALLOCATION_RETVAL(driver_data, false); driver_data->instance = wgpuCreateInstance(PULSE_NULLPTR); diff --git a/Sources/Backends/WebGPU/WebGPU.h b/Sources/Backends/WebGPU/WebGPU.h index 578a9f8..35a1298 100644 --- a/Sources/Backends/WebGPU/WebGPU.h +++ b/Sources/Backends/WebGPU/WebGPU.h @@ -23,4 +23,3 @@ PulseBackendFlags WebGPUCheckSupport(PulseBackendFlags candidates, PulseShaderFo #endif // PULSE_WEBGPU_H_ #endif // PULSE_ENABLE_WEBGPU_BACKEND - diff --git a/Sources/Backends/WebGPU/WebGPUBuffer.c b/Sources/Backends/WebGPU/WebGPUBuffer.c index ab7be42..4de3f41 100644 --- a/Sources/Backends/WebGPU/WebGPUBuffer.c +++ b/Sources/Backends/WebGPU/WebGPUBuffer.c @@ -2,19 +2,149 @@ // This file is part of "Pulse" // For conditions of distribution and use, see copyright notice in LICENSE +#include +#include + #include +#include "../../PulseInternal.h" #include "WebGPU.h" +#include "WebGPUDevice.h" +#include "webgpu.h" +#include "WebGPUBuffer.h" PulseBuffer WebGPUCreateBuffer(PulseDevice device, const PulseBufferCreateInfo* create_infos) { + WebGPUDevice* webgpu_device = WEBGPU_RETRIEVE_DRIVER_DATA_AS(device, WebGPUDevice*); + + PulseBufferHandler* buffer = (PulseBufferHandler*)calloc(1, sizeof(PulseBufferHandler)); + PULSE_CHECK_ALLOCATION_RETVAL(buffer, PULSE_NULL_HANDLE); + + WebGPUBuffer* webgpu_buffer = (WebGPUBuffer*)calloc(1, sizeof(WebGPUBuffer)); + PULSE_CHECK_ALLOCATION_RETVAL(webgpu_buffer, PULSE_NULL_HANDLE); + + buffer->device = device; + buffer->driver_data = webgpu_buffer; + buffer->size = create_infos->size; + buffer->usage = create_infos->usage; + + bool is_storage = false; + + WGPUBufferDescriptor descriptor = { 0 }; + descriptor.mappedAtCreation = false; + descriptor.size = buffer->size; + if(buffer->usage & PULSE_BUFFER_USAGE_STORAGE_READ || buffer->usage & PULSE_BUFFER_USAGE_STORAGE_WRITE) + { + descriptor.usage |= WGPUBufferUsage_Storage | WGPUBufferUsage_CopyDst | WGPUBufferUsage_CopySrc; + is_storage = true; + } + if(buffer->usage & PULSE_BUFFER_USAGE_TRANSFER_DOWNLOAD) + { + descriptor.usage |= WGPUBufferUsage_CopyDst; + if(!is_storage) + descriptor.usage |= WGPUBufferUsage_MapRead; + } + if(buffer->usage & PULSE_BUFFER_USAGE_TRANSFER_UPLOAD) + descriptor.usage |= WGPUBufferUsage_CopyDst | WGPUBufferUsage_CopySrc; + if(buffer->usage & PULSE_INTERNAL_BUFFER_USAGE_UNIFORM_ACCESS) + descriptor.usage |= WGPUBufferUsage_Uniform | WGPUBufferUsage_CopyDst | WGPUBufferUsage_CopySrc; + + webgpu_buffer->buffer = wgpuDeviceCreateBuffer(webgpu_device->device, &descriptor); + if(webgpu_buffer->buffer == PULSE_NULLPTR) + { + free(webgpu_buffer); + free(buffer); + return PULSE_NULL_HANDLE; + } + return buffer; } -bool WebGPUMapBuffer(PulseBuffer buffer, void** data) +#include + +static void WebGPUMapBufferCallback(WGPUMapAsyncStatus status, WGPUStringView message, void* userdata1, void* userdata2) { + atomic_int* mapping_finished = (atomic_int*)userdata1; + PulseBuffer buffer = (PulseBuffer)userdata2; + puts("test"); + if(status == WGPUMapAsyncStatus_Success) + atomic_store(mapping_finished, 1); + else + { + const char* reasons[] = { + "nvm it was successfull", + "instance has been dropped", + "an error occured", + "mapping was aborted", + }; + if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(buffer->device->backend)) + PulseLogErrorFmt(buffer->device->backend, "(WebGPU) buffer mapping failed because %s. %.*s", reasons[status], message.length, message.data); + PulseSetInternalError(PULSE_ERROR_MAP_FAILED); + atomic_store(mapping_finished, 2); + } +} + +bool WebGPUMapBuffer(PulseBuffer buffer, PulseMapMode mode, void** data) +{ + WebGPUBuffer* webgpu_buffer = WEBGPU_RETRIEVE_DRIVER_DATA_AS(buffer, WebGPUBuffer*); + + // If we only upload we can just use wgpuQueueWriteBuffer + // https://toji.dev/webgpu-best-practices/buffer-uploads.html + if(mode == PULSE_MAP_WRITE) + { + if(webgpu_buffer->map == PULSE_NULLPTR) + webgpu_buffer->map = malloc(buffer->size); + else + webgpu_buffer->map = realloc(webgpu_buffer->map, buffer->size); + PULSE_CHECK_ALLOCATION_RETVAL(webgpu_buffer->map, false); + } + else + { + atomic_int mapping_finished; + atomic_store(&mapping_finished, 0); + + const uint32_t timeout = 5000; + clock_t start = clock(); + + webgpu_buffer->map = PULSE_NULLPTR; + + WGPUBufferMapCallbackInfo callback_info = { 0 }; + callback_info.mode = WGPUCallbackMode_AllowSpontaneous; + callback_info.callback = WebGPUMapBufferCallback; + callback_info.userdata1 = &mapping_finished; + callback_info.userdata2 = buffer; + wgpuBufferMapAsync(webgpu_buffer->buffer, WGPUMapMode_Read, 0, buffer->size, callback_info); + + while(atomic_load(&mapping_finished) == 0) + { + clock_t elapsed = clock() - start; + if(elapsed > timeout) + { + if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(buffer->device->backend)) + PulseLogError(buffer->device->backend, "(WebGPU) buffer mapping failed (timeout)"); + PulseSetInternalError(PULSE_ERROR_MAP_FAILED); + return false; + } + PulseSleep(1); // 1ms + } + + if(atomic_load(&mapping_finished) == 1) + webgpu_buffer->map = (void*)wgpuBufferGetConstMappedRange(webgpu_buffer->buffer, 0, WGPU_WHOLE_MAP_SIZE); + } + if(webgpu_buffer->map == PULSE_NULLPTR) + return false; + webgpu_buffer->current_map_mode = mode; + *data = webgpu_buffer->map; + return true; } void WebGPUUnmapBuffer(PulseBuffer buffer) { + WebGPUDevice* webgpu_device = WEBGPU_RETRIEVE_DRIVER_DATA_AS(buffer->device, WebGPUDevice*); + WebGPUBuffer* webgpu_buffer = WEBGPU_RETRIEVE_DRIVER_DATA_AS(buffer, WebGPUBuffer*); + if(webgpu_buffer->current_map_mode == PULSE_MAP_WRITE) + wgpuQueueWriteBuffer(webgpu_device->queue, webgpu_buffer->buffer, 0, webgpu_buffer->map, buffer->size); + else + wgpuBufferUnmap(webgpu_buffer->buffer); + webgpu_buffer->map = PULSE_NULLPTR; } bool WebGPUCopyBufferToBuffer(PulseCommandList cmd, const PulseBufferRegion* src, const PulseBufferRegion* dst) @@ -27,4 +157,9 @@ bool WebGPUCopyBufferToImage(PulseCommandList cmd, const PulseBufferRegion* src, void WebGPUDestroyBuffer(PulseDevice device, PulseBuffer buffer) { + PULSE_UNUSED(device); + WebGPUBuffer* webgpu_buffer = WEBGPU_RETRIEVE_DRIVER_DATA_AS(buffer, WebGPUBuffer*); + wgpuBufferRelease(webgpu_buffer->buffer); + free(webgpu_buffer); + free(buffer); } diff --git a/Sources/Backends/WebGPU/WebGPUBuffer.h b/Sources/Backends/WebGPU/WebGPUBuffer.h index e6fc560..af7846b 100644 --- a/Sources/Backends/WebGPU/WebGPUBuffer.h +++ b/Sources/Backends/WebGPU/WebGPUBuffer.h @@ -10,14 +10,16 @@ #include #include -#include "../../PulseInternal.h" typedef struct WebGPUBuffer { + WGPUBuffer buffer; + void* map; + PulseMapMode current_map_mode; } WebGPUBuffer; PulseBuffer WebGPUCreateBuffer(PulseDevice device, const PulseBufferCreateInfo* create_infos); -bool WebGPUMapBuffer(PulseBuffer buffer, void** data); +bool WebGPUMapBuffer(PulseBuffer buffer, PulseMapMode mode, void** data); void WebGPUUnmapBuffer(PulseBuffer buffer); bool WebGPUCopyBufferToBuffer(PulseCommandList cmd, const PulseBufferRegion* src, const PulseBufferRegion* dst); bool WebGPUCopyBufferToImage(PulseCommandList cmd, const PulseBufferRegion* src, const PulseImageRegion* dst); diff --git a/Sources/Backends/WebGPU/WebGPUCommandList.c b/Sources/Backends/WebGPU/WebGPUCommandList.c index 3aef373..5a2b057 100644 --- a/Sources/Backends/WebGPU/WebGPUCommandList.c +++ b/Sources/Backends/WebGPU/WebGPUCommandList.c @@ -41,15 +41,17 @@ PulseCommandList WebGPURequestCommandList(PulseDevice device, PulseCommandListUs return cmd; } -#include - static void WebGPUFenceCallback(WGPUQueueWorkDoneStatus status, void* userdata1, void* userdata2) { PULSE_UNUSED(userdata2); WebGPUFence* webgpu_fence = (WebGPUFence*)userdata1; + PulseCommandList cmd = (PulseCommandList)userdata2; if(status == WGPUQueueWorkDoneStatus_Success) - atomic_store(&webgpu_fence->signal, true); - puts("test"); + { + if(webgpu_fence != PULSE_NULLPTR) + atomic_store(&webgpu_fence->signal, true); + cmd->state = PULSE_COMMAND_LIST_STATE_READY; + } } bool WebGPUSubmitCommandList(PulseDevice device, PulseCommandList cmd, PulseFence fence) @@ -60,17 +62,23 @@ bool WebGPUSubmitCommandList(PulseDevice device, PulseCommandList cmd, PulseFenc WGPUCommandBufferDescriptor command_buffer_descriptor = { 0 }; WGPUCommandBuffer command_buffer = wgpuCommandEncoderFinish(webgpu_cmd->encoder, &command_buffer_descriptor); - wgpuQueueSubmit(webgpu_device->queue, 1, &command_buffer); - - WebGPUFence* webgpu_fence = WEBGPU_RETRIEVE_DRIVER_DATA_AS(fence, WebGPUFence*); - atomic_store(&webgpu_fence->signal, false); - WGPUQueueWorkDoneCallbackInfo callback = { 0 }; callback.mode = WGPUCallbackMode_AllowSpontaneous; callback.callback = WebGPUFenceCallback; - callback.userdata1 = webgpu_fence; + callback.userdata1 = PULSE_NULLPTR; + callback.userdata2 = cmd; + if(fence != PULSE_NULL_HANDLE) + { + WebGPUFence* webgpu_fence = WEBGPU_RETRIEVE_DRIVER_DATA_AS(fence, WebGPUFence*); + callback.userdata1 = webgpu_fence; + atomic_store(&webgpu_fence->signal, false); + fence->cmd = cmd; + } wgpuQueueOnSubmittedWorkDone(webgpu_device->queue, callback); + cmd->state = PULSE_COMMAND_LIST_STATE_SENT; + wgpuQueueSubmit(webgpu_device->queue, 1, &command_buffer); + wgpuCommandBufferRelease(command_buffer); return true; } diff --git a/Sources/Backends/WebGPU/WebGPUComputePass.c b/Sources/Backends/WebGPU/WebGPUComputePass.c index c06309b..c60bd6a 100644 --- a/Sources/Backends/WebGPU/WebGPUComputePass.c +++ b/Sources/Backends/WebGPU/WebGPUComputePass.c @@ -3,6 +3,7 @@ // For conditions of distribution and use, see copyright notice in LICENSE #include +#include "../../PulseInternal.h" #include "WebGPU.h" #include "WebGPUDevice.h" #include "WebGPUComputePass.h" diff --git a/Sources/Backends/WebGPU/WebGPUDevice.c b/Sources/Backends/WebGPU/WebGPUDevice.c index 2f89e80..9335741 100644 --- a/Sources/Backends/WebGPU/WebGPUDevice.c +++ b/Sources/Backends/WebGPU/WebGPUDevice.c @@ -35,7 +35,7 @@ device->adapter = adapter; } #else - static uint64_t WebGPUScoreAdapter(WGPUInstance instance, WGPUAdapter adapter) + static uint64_t WebGPUScoreAdapter(WGPUAdapter adapter) { uint64_t score = 0; WGPUAdapterInfo infos; @@ -172,7 +172,7 @@ PulseDevice WebGPUCreateDevice(PulseBackend backend, PulseDevice* forbiden_devic { if(WebGPUIsDeviceForbidden(adapters[i], forbiden_devices, forbiden_devices_count)) continue; - uint64_t current_device_score = WebGPUScoreAdapter(instance, adapters[i]); + uint64_t current_device_score = WebGPUScoreAdapter(adapters[i]); if(current_device_score > best_device_score) { best_device_score = current_device_score; @@ -196,7 +196,7 @@ PulseDevice WebGPUCreateDevice(PulseBackend backend, PulseDevice* forbiden_devic WGPUDeviceLostCallbackInfo lost_callback = { 0 }; lost_callback.callback = WebGPUDeviceLostCallback; - lost_callback.mode = WGPUCallbackMode_AllowProcessEvents; + lost_callback.mode = WGPUCallbackMode_AllowSpontaneous; lost_callback.userdata1 = device; lost_callback.userdata2 = backend; WGPUUncapturedErrorCallbackInfo uncaptured_callback = { 0 }; @@ -209,7 +209,7 @@ PulseDevice WebGPUCreateDevice(PulseBackend backend, PulseDevice* forbiden_devic descriptor.uncapturedErrorCallbackInfo = uncaptured_callback; WGPURequestDeviceCallbackInfo device_callback = { 0 }; device_callback.callback = WebGPURequestDeviceCallback; - device_callback.mode = WGPUCallbackMode_AllowProcessEvents; + device_callback.mode = WGPUCallbackMode_AllowSpontaneous; device_callback.userdata1 = device; device_callback.userdata2 = backend; wgpuAdapterRequestDevice(device->adapter, &descriptor, device_callback); diff --git a/Sources/Backends/WebGPU/WebGPUFence.c b/Sources/Backends/WebGPU/WebGPUFence.c index 102deab..dc33005 100644 --- a/Sources/Backends/WebGPU/WebGPUFence.c +++ b/Sources/Backends/WebGPU/WebGPUFence.c @@ -35,11 +35,9 @@ bool WebGPUIsFenceReady(PulseDevice device, PulseFence fence) { PULSE_UNUSED(device); WebGPUFence* webgpu_fence = WEBGPU_RETRIEVE_DRIVER_DATA_AS(fence, WebGPUFence*); - return atomic_load(&webgpu_fence->signal) == true; + return atomic_load(&webgpu_fence->signal); } -#include - bool WebGPUWaitForFences(PulseDevice device, const PulseFence* fences, uint32_t fences_count, bool wait_for_all) { PULSE_UNUSED(device); diff --git a/Sources/PulseBuffer.c b/Sources/PulseBuffer.c index c28b34d..8c8fa47 100644 --- a/Sources/PulseBuffer.c +++ b/Sources/PulseBuffer.c @@ -19,6 +19,7 @@ PULSE_API PulseBuffer PulseCreateBuffer(PulseDevice device, const PulseBufferCre return PULSE_NULL_HANDLE; } } + PulseBuffer buffer = device->PFN_CreateBuffer(device, create_infos); if(buffer == PULSE_NULL_HANDLE) return PULSE_NULL_HANDLE; @@ -28,11 +29,53 @@ PULSE_API PulseBuffer PulseCreateBuffer(PulseDevice device, const PulseBufferCre return buffer; } -PULSE_API bool PulseMapBuffer(PulseBuffer buffer, void** data) +PULSE_API bool PulseMapBuffer(PulseBuffer buffer, PulseMapMode mode, void** data) { PULSE_CHECK_HANDLE_RETVAL(buffer, false); PULSE_CHECK_PTR_RETVAL(data, false); - bool res = buffer->device->PFN_MapBuffer(buffer, data); + + if(buffer->is_mapped) + { + if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(buffer->device->backend)) + PulseLogError(buffer->device->backend, "buffer is already mapped"); + PulseSetInternalError(PULSE_ERROR_MAP_FAILED); + return false; + } + + PulseFlags storage_flags = PULSE_BUFFER_USAGE_STORAGE_READ | PULSE_BUFFER_USAGE_STORAGE_WRITE; + if((buffer->usage & storage_flags) != 0) + { + if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(buffer->device->backend)) + PulseLogError(buffer->device->backend, "cannot map a buffer that has been created with storage flags"); + PulseSetInternalError(PULSE_ERROR_MAP_FAILED); + return false; + } + + PulseFlags transfer_flags = PULSE_BUFFER_USAGE_TRANSFER_UPLOAD | PULSE_BUFFER_USAGE_TRANSFER_DOWNLOAD; + if((buffer->usage & transfer_flags) == 0) + { + if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(buffer->device->backend)) + PulseLogError(buffer->device->backend, "cannot map a buffer that has not been created with upload or download flags"); + PulseSetInternalError(PULSE_ERROR_MAP_FAILED); + return false; + } + + if((buffer->usage & PULSE_BUFFER_USAGE_TRANSFER_UPLOAD) == 0 && mode == PULSE_MAP_WRITE) + { + if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(buffer->device->backend)) + PulseLogError(buffer->device->backend, "cannot map a buffer that has not been created with upload flags for writting"); + PulseSetInternalError(PULSE_ERROR_MAP_FAILED); + return false; + } + if((buffer->usage & PULSE_BUFFER_USAGE_TRANSFER_DOWNLOAD) == 0 && mode == PULSE_MAP_READ) + { + if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(buffer->device->backend)) + PulseLogError(buffer->device->backend, "cannot map a buffer that has not been created with download flags for reading"); + PulseSetInternalError(PULSE_ERROR_MAP_FAILED); + return false; + } + + bool res = buffer->device->PFN_MapBuffer(buffer, mode, data); if(res) buffer->is_mapped = true; return res; @@ -41,6 +84,14 @@ PULSE_API bool PulseMapBuffer(PulseBuffer buffer, void** data) PULSE_API void PulseUnmapBuffer(PulseBuffer buffer) { PULSE_CHECK_HANDLE(buffer); + + if(!buffer->is_mapped) + { + if(PULSE_IS_BACKEND_HIGH_LEVEL_DEBUG(buffer->device->backend)) + PulseLogError(buffer->device->backend, "buffer is not mapped"); + return; + } + buffer->device->PFN_UnmapBuffer(buffer); buffer->is_mapped = false; } diff --git a/Sources/PulseDefs.h b/Sources/PulseDefs.h index 3c610d2..520a45d 100644 --- a/Sources/PulseDefs.h +++ b/Sources/PulseDefs.h @@ -61,35 +61,37 @@ array[defrag_i] = array[defrag_i + 1]; \ #define PULSE_CHECK_COMMAND_LIST_STATE_RETVAL(cmd, retval) \ - if(cmd->state != PULSE_COMMAND_LIST_STATE_RECORDING) \ - { \ - switch(cmd->state) \ + do { \ + if(cmd->state != PULSE_COMMAND_LIST_STATE_RECORDING) \ { \ - case PULSE_COMMAND_LIST_STATE_INVALID: \ - if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(cmd->device->backend)) \ - PulseLogError(cmd->device->backend, "command list is in invalid state"); \ - return retval; \ - case PULSE_COMMAND_LIST_STATE_READY: \ - if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(cmd->device->backend)) \ - PulseLogError(cmd->device->backend, "command list is not recording"); \ - return retval; \ - case PULSE_COMMAND_LIST_STATE_SENT: \ - if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(cmd->device->backend)) \ - PulseLogWarning(cmd->device->backend, "command list has already been submitted"); \ - return retval; \ - default: break; \ + switch(cmd->state) \ + { \ + case PULSE_COMMAND_LIST_STATE_INVALID: \ + if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(cmd->device->backend)) \ + PulseLogError(cmd->device->backend, "command list is in invalid state"); \ + return retval; \ + case PULSE_COMMAND_LIST_STATE_READY: \ + if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(cmd->device->backend)) \ + PulseLogError(cmd->device->backend, "command list is not recording"); \ + return retval; \ + case PULSE_COMMAND_LIST_STATE_SENT: \ + if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(cmd->device->backend)) \ + PulseLogWarning(cmd->device->backend, "command list has already been submitted"); \ + return retval; \ + default: break; \ + } \ } \ - } \ + } while(0); \ #define PULSE_CHECK_COMMAND_LIST_STATE(cmd) PULSE_CHECK_COMMAND_LIST_STATE_RETVAL(cmd, ) #ifndef PULSE_STATIC_ASSERT #ifdef __cplusplus #if __cplusplus >= 201103L - #define PULSE_STATIC_ASSERT(name, x) static_assert(x, #x) + #define PULSE_STATIC_ASSERT(name, x) static_assert(x, #x) #endif #elif PULSE_C_VERSION >= 2023 - #define PULSE_STATIC_ASSERT(name, x) static_assert(x, #x) + #define PULSE_STATIC_ASSERT(name, x) static_assert(x, #x) #elif PULSE_C_VERSION >= 2011 #define PULSE_STATIC_ASSERT(name, x) _Static_assert(x, #x) #else diff --git a/Sources/PulseFence.c b/Sources/PulseFence.c index 18fd439..0bf16d3 100644 --- a/Sources/PulseFence.c +++ b/Sources/PulseFence.c @@ -28,5 +28,11 @@ PULSE_API bool PulseWaitForFences(PulseDevice device, const PulseFence* fences, { PULSE_CHECK_HANDLE_RETVAL(device, false); PULSE_CHECK_PTR_RETVAL(fences, false); - return device->PFN_WaitForFences(device, fences, fences_count, wait_for_all); + bool res = device->PFN_WaitForFences(device, fences, fences_count, wait_for_all); + if(res) + { + for(uint32_t i = 0; i < fences_count; i++) + fences[i]->cmd->state = PULSE_COMMAND_LIST_STATE_READY; + } + return res; } diff --git a/Sources/PulsePFNs.h b/Sources/PulsePFNs.h index fabdda1..fcbad22 100644 --- a/Sources/PulsePFNs.h +++ b/Sources/PulsePFNs.h @@ -25,7 +25,7 @@ typedef PulseCommandList (*PulseRequestCommandListPFN)(PulseDevice, PulseCommand typedef bool (*PulseSubmitCommandListPFN)(PulseDevice, PulseCommandList, PulseFence); typedef void (*PulseReleaseCommandListPFN)(PulseDevice, PulseCommandList); typedef PulseBuffer (*PulseCreateBufferPFN)(PulseDevice, const PulseBufferCreateInfo*); -typedef bool (*PulseMapBufferPFN)(PulseBuffer, void**); +typedef bool (*PulseMapBufferPFN)(PulseBuffer, PulseMapMode, void**); typedef void (*PulseUnmapBufferPFN)(PulseBuffer); typedef void (*PulseDestroyBufferPFN)(PulseDevice, PulseBuffer); typedef PulseImage (*PulseCreateImagePFN)(PulseDevice, const PulseImageCreateInfo*); diff --git a/Tests/Vulkan/Buffer.c b/Tests/Vulkan/Buffer.c index ff16bf3..712f09c 100644 --- a/Tests/Vulkan/Buffer.c +++ b/Tests/Vulkan/Buffer.c @@ -78,20 +78,20 @@ void TestBufferMapping() PulseBufferCreateInfo buffer_create_info = { 0 }; buffer_create_info.size = 8; - buffer_create_info.usage = PULSE_BUFFER_USAGE_TRANSFER_UPLOAD; + buffer_create_info.usage = PULSE_BUFFER_USAGE_TRANSFER_UPLOAD | PULSE_BUFFER_USAGE_TRANSFER_DOWNLOAD; PulseBuffer buffer = PulseCreateBuffer(device, &buffer_create_info); TEST_ASSERT_NOT_EQUAL_MESSAGE(buffer, PULSE_NULL_HANDLE, PulseVerbaliseErrorType(PulseGetLastErrorType())); { void* ptr; - TEST_ASSERT_NOT_EQUAL_MESSAGE(PulseMapBuffer(buffer, &ptr), false, PulseVerbaliseErrorType(PulseGetLastErrorType())); + TEST_ASSERT_NOT_EQUAL_MESSAGE(PulseMapBuffer(buffer, PULSE_MAP_WRITE, &ptr), false, PulseVerbaliseErrorType(PulseGetLastErrorType())); TEST_ASSERT_NOT_NULL(ptr); memcpy(ptr, data, 8); PulseUnmapBuffer(buffer); } { void* ptr; - TEST_ASSERT_NOT_EQUAL_MESSAGE(PulseMapBuffer(buffer, &ptr), false, PulseVerbaliseErrorType(PulseGetLastErrorType())); + TEST_ASSERT_NOT_EQUAL_MESSAGE(PulseMapBuffer(buffer, PULSE_MAP_READ, &ptr), false, PulseVerbaliseErrorType(PulseGetLastErrorType())); TEST_ASSERT_NOT_NULL(ptr); TEST_ASSERT_EQUAL(memcmp(ptr, data, 8), 0); PulseUnmapBuffer(buffer); @@ -99,7 +99,7 @@ void TestBufferMapping() DISABLE_ERRORS; void* ptr; - TEST_ASSERT_NOT_EQUAL_MESSAGE(PulseMapBuffer(buffer, &ptr), false, PulseVerbaliseErrorType(PulseGetLastErrorType())); + TEST_ASSERT_NOT_EQUAL_MESSAGE(PulseMapBuffer(buffer, PULSE_MAP_READ, &ptr), false, PulseVerbaliseErrorType(PulseGetLastErrorType())); PulseDestroyBuffer(device, buffer); ENABLE_ERRORS; @@ -124,7 +124,7 @@ void TestBufferCopy() { void* ptr; - TEST_ASSERT_NOT_EQUAL_MESSAGE(PulseMapBuffer(src_buffer, &ptr), false, PulseVerbaliseErrorType(PulseGetLastErrorType())); + TEST_ASSERT_NOT_EQUAL_MESSAGE(PulseMapBuffer(src_buffer, PULSE_MAP_WRITE, &ptr), false, PulseVerbaliseErrorType(PulseGetLastErrorType())); TEST_ASSERT_NOT_NULL(ptr); memcpy(ptr, data, 8); PulseUnmapBuffer(src_buffer); @@ -158,7 +158,7 @@ void TestBufferCopy() { void* ptr; - TEST_ASSERT_NOT_EQUAL_MESSAGE(PulseMapBuffer(dst_buffer, &ptr), false, PulseVerbaliseErrorType(PulseGetLastErrorType())); + TEST_ASSERT_NOT_EQUAL_MESSAGE(PulseMapBuffer(dst_buffer, PULSE_MAP_READ, &ptr), false, PulseVerbaliseErrorType(PulseGetLastErrorType())); TEST_ASSERT_NOT_NULL(ptr); TEST_ASSERT_EQUAL(memcmp(ptr, data, 8), 0); PulseUnmapBuffer(dst_buffer); @@ -192,7 +192,7 @@ void TestBufferCopy() DISABLE_ERRORS; void* ptr; - TEST_ASSERT_NOT_EQUAL_MESSAGE(PulseMapBuffer(src_buffer, &ptr), false, PulseVerbaliseErrorType(PulseGetLastErrorType())); + TEST_ASSERT_NOT_EQUAL_MESSAGE(PulseMapBuffer(src_buffer, PULSE_MAP_READ, &ptr), false, PulseVerbaliseErrorType(PulseGetLastErrorType())); PulseDestroyBuffer(device, src_buffer); ENABLE_ERRORS; @@ -219,7 +219,7 @@ void TestBufferCopyImage() { void* ptr; - TEST_ASSERT_NOT_EQUAL_MESSAGE(PulseMapBuffer(buffer, &ptr), false, PulseVerbaliseErrorType(PulseGetLastErrorType())); + TEST_ASSERT_NOT_EQUAL_MESSAGE(PulseMapBuffer(buffer, PULSE_MAP_WRITE, &ptr), false, PulseVerbaliseErrorType(PulseGetLastErrorType())); TEST_ASSERT_NOT_NULL(ptr); memcpy(ptr, data, 8); PulseUnmapBuffer(buffer); @@ -310,7 +310,7 @@ void TestBufferComputeWrite() void* ptr; uint32_t data[256]; memset(data, 0xFF, 256 * sizeof(uint32_t)); - TEST_ASSERT_NOT_EQUAL_MESSAGE(PulseMapBuffer(buffer, &ptr), false, PulseVerbaliseErrorType(PulseGetLastErrorType())); + TEST_ASSERT_NOT_EQUAL_MESSAGE(PulseMapBuffer(buffer, PULSE_MAP_READ, &ptr), false, PulseVerbaliseErrorType(PulseGetLastErrorType())); TEST_ASSERT_NOT_NULL(ptr); TEST_ASSERT_EQUAL(memcmp(ptr, data, 256 * sizeof(uint32_t)), 0); PulseUnmapBuffer(buffer); @@ -347,7 +347,7 @@ void TestBufferComputeCopy() { void* ptr; - TEST_ASSERT_NOT_EQUAL_MESSAGE(PulseMapBuffer(read_buffer, &ptr), false, PulseVerbaliseErrorType(PulseGetLastErrorType())); + TEST_ASSERT_NOT_EQUAL_MESSAGE(PulseMapBuffer(read_buffer, PULSE_MAP_WRITE, &ptr), false, PulseVerbaliseErrorType(PulseGetLastErrorType())); TEST_ASSERT_NOT_NULL(ptr); memcpy(ptr, data, 256 * sizeof(uint32_t)); PulseUnmapBuffer(read_buffer); @@ -378,7 +378,7 @@ void TestBufferComputeCopy() { void* ptr; - TEST_ASSERT_NOT_EQUAL_MESSAGE(PulseMapBuffer(write_buffer, &ptr), false, PulseVerbaliseErrorType(PulseGetLastErrorType())); + TEST_ASSERT_NOT_EQUAL_MESSAGE(PulseMapBuffer(write_buffer, PULSE_MAP_READ, &ptr), false, PulseVerbaliseErrorType(PulseGetLastErrorType())); TEST_ASSERT_NOT_NULL(ptr); TEST_ASSERT_EQUAL(memcmp(ptr, data, 256 * sizeof(uint32_t)), 0); PulseUnmapBuffer(write_buffer);