yes

.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

Examples/Vulkan/main.c

@@ -3,13 +3,6 @@
 #include <stdio.h>
 #include <stdlib.h>
 
-#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,14 +16,21 @@ 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);
 
+	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);
+
+	// GPU computations
+	{
 		const uint8_t shader_bytecode[] = {
 			#include "shader.spv.h"
 		};
@@ -41,44 +41,60 @@ int main(void)
 		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.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);
 
 		PulseComputePass pass = PulseBeginComputePass(cmd);
-	CHECK_PULSE_HANDLE_RETVAL(pass, 1);
 			PulseBindStorageBuffers(pass, &buffer, 1);
 			PulseBindComputePipeline(pass, pipeline);
 			PulseDispatchComputations(pass, 32, 32, 1);
 		PulseEndComputePass(pass);
 
-	if(!PulseSubmitCommandList(device, cmd, fence))
+		PulseSubmitCommandList(device, cmd, fence);
-		fprintf(stderr, "Could not submit command list, %s\n", PulseVerbaliseErrorType(PulseGetLastErrorType()));
+		PulseWaitForFences(device, &fence, 1, true);
-	if(!PulseWaitForFences(device, &fence, 1, true))
-		fprintf(stderr, "Could not wait for fences, %s\n", PulseVerbaliseErrorType(PulseGetLastErrorType()));
-
-	void* ptr;
-	PulseMapBuffer(buffer, &ptr);
-	for(uint32_t i = 0; i < 256; i++)
-		printf("%d, ", ((int32_t*)ptr)[i]);
-	puts("");
-	PulseUnmapBuffer(buffer);
 
 		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);
 

Examples/WebGPU/main.c

@@ -5,12 +5,6 @@
 #include <string.h>
 
 #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);
 
+	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);
+
+	// 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);
-	CHECK_PULSE_HANDLE_RETVAL(pipeline, 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);
 
 		PulseComputePass pass = PulseBeginComputePass(cmd);
-	CHECK_PULSE_HANDLE_RETVAL(pass, 1);
+		//	PulseBindStorageBuffers(pass, &buffer, 1);
 			PulseBindComputePipeline(pass, pipeline);
 			PulseDispatchComputations(pass, 32, 32, 1);
 		PulseEndComputePass(pass);
 
-	if(!PulseSubmitCommandList(device, cmd, fence))
+		PulseSubmitCommandList(device, cmd, fence);
-		fprintf(stderr, "Could not submit command list, %s\n", PulseVerbaliseErrorType(PulseGetLastErrorType()));
+		PulseWaitForFences(device, &fence, 1, true);
-	if(!PulseWaitForFences(device, &fence, 1, true))
-		fprintf(stderr, "Could not wait for fences, %s\n", PulseVerbaliseErrorType(PulseGetLastErrorType()));
 
 		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);

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
+	if has_config("webgpu") then
 		includes("WebGPU/xmake.lua")
 	end
+end

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);

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);

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);

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);
+	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;

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:
+		case VK_SUCCESS: break;
-			for(uint32_t i = 0; i < fences_count; i++)
-				fences[i]->cmd->state = PULSE_COMMAND_LIST_STATE_READY;
-		break;
 		case VK_TIMEOUT: break;
 
 		case VK_ERROR_DEVICE_LOST: PulseSetInternalError(PULSE_ERROR_DEVICE_LOST); return false;

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);

Sources/Backends/WebGPU/WebGPU.h

@@ -23,4 +23,3 @@ PulseBackendFlags WebGPUCheckSupport(PulseBackendFlags candidates, PulseShaderFo
 #endif // PULSE_WEBGPU_H_
 
 #endif // PULSE_ENABLE_WEBGPU_BACKEND
-

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 <stdatomic.h>
+#include <time.h>
+
 #include <Pulse.h>
+#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 <stdio.h>
+
+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);
 }

Sources/Backends/WebGPU/WebGPUBuffer.h

@@ -10,14 +10,16 @@
 #include <webgpu/webgpu.h>
 
 #include <Pulse.h>
-#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);

Sources/Backends/WebGPU/WebGPUCommandList.c

@@ -41,15 +41,17 @@ PulseCommandList WebGPURequestCommandList(PulseDevice device, PulseCommandListUs
 	return cmd;
 }
 
-#include <stdio.h>
-
 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)
+	{
+		if(webgpu_fence != PULSE_NULLPTR)
 			atomic_store(&webgpu_fence->signal, true);
-	puts("test");
+		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 = 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;
 }

Sources/Backends/WebGPU/WebGPUComputePass.c

@@ -3,6 +3,7 @@
 // For conditions of distribution and use, see copyright notice in LICENSE
 
 #include <Pulse.h>
+#include "../../PulseInternal.h"
 #include "WebGPU.h"
 #include "WebGPUDevice.h"
 #include "WebGPUComputePass.h"

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);

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 <stdio.h>
-
 bool WebGPUWaitForFences(PulseDevice device, const PulseFence* fences, uint32_t fences_count, bool wait_for_all)
 {
 	PULSE_UNUSED(device);

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;
 }

Sources/PulseDefs.h

@@ -61,6 +61,7 @@
 		array[defrag_i] = array[defrag_i + 1]; \
 
 #define PULSE_CHECK_COMMAND_LIST_STATE_RETVAL(cmd, retval) \
+	do { \
 		if(cmd->state != PULSE_COMMAND_LIST_STATE_RECORDING) \
 		{ \
 			switch(cmd->state) \
@@ -80,6 +81,7 @@
 				default: break; \
 			} \
 		} \
+	} while(0); \
 
 #define PULSE_CHECK_COMMAND_LIST_STATE(cmd) PULSE_CHECK_COMMAND_LIST_STATE_RETVAL(cmd, )
 

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;
 }

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*);

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);