yes

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

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

Sources/Backends/Vulkan/VulkanLoader.c

@@ -28,12 +28,7 @@
 	#endif
 #endif
 
-#ifdef PULSE_PLAT_WINDOWS
+#ifndef 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
 	#include <dlfcn.h>
 	typedef void* LibModule;
 #endif

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

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 <webgpu/webgpu.h>
 
 #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

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)

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

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)

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

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

Tests/Backend.c

@@ -0,0 +1,72 @@
+#include "Common.h"
+
+#include <unity/unity.h>
+#include <Pulse.h>
+
+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);
+}

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);
 
+	#if defined(VULKAN_ENABLED)
 		const uint8_t shader_bytecode[] = {
-		#include "Shaders/SimpleBufferWrite.spv.h"
+			#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);
 
+	#if defined(VULKAN_ENABLED)
 		const uint8_t shader_bytecode[] = {
-		#include "Shaders/BufferCopy.spv.h"
+			#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()));

Tests/Common.c

@@ -1,5 +1,7 @@
 #include "Common.h"
 #include <unity/unity.h>
+#include <stdio.h>
+#include <string.h>
 
 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)
 {
+	#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 };
+	#if defined(WEBGPU_ENABLED)
+		info.code_size = strlen(code);
+	#else
 		info.code_size = code_size;
+	#endif
 	info.code = code;
 	info.entrypoint = "main";
+	#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;

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

Tests/Pipeline.c

@@ -10,9 +10,14 @@ void TestPipelineSetup()
 	PulseDevice device;
 	SetupDevice(backend, &device);
 
+	#if defined(VULKAN_ENABLED)
 		const uint8_t shader_bytecode[] = {
-		#include "Shaders/Simple.spv.h"
+			#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);
 
+	#if defined(VULKAN_ENABLED)
 		const uint8_t shader_bytecode[] = {
-		#include "Shaders/ReadOnlyBindings.spv.h"
+			#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);
 
+	#if defined(VULKAN_ENABLED)
 		const uint8_t shader_bytecode[] = {
-		#include "Shaders/ReadOnlyBindings.spv.h"
+			#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);
 
+	#if defined(VULKAN_ENABLED)
 		const uint8_t shader_bytecode[] = {
-		#include "Shaders/ReadOnlyBindings.spv.h"
+			#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()));

Tests/Shaders/.gitignore

@@ -0,0 +1,2 @@
+*.spv.h
+*.wgsl.h

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

Tests/Shaders/Vulkan/ReadOnlyBindings.nzsl

@@ -19,7 +19,7 @@ external
 }
 
 [entry(compute)]
-[workgroup(32, 32, 1)]
+[workgroup(16, 16, 1)]
 fn main(input: Input)
 {
 }

Tests/Shaders/Vulkan/ReadWriteBindings.nzsl

@@ -21,7 +21,7 @@ external
 }
 
 [entry(compute)]
-[workgroup(32, 32, 1)]
+[workgroup(16, 16, 1)]
 fn main(input: Input)
 {
 }

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

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)
 {
 }

Tests/Shaders/WebGPU/BufferCopy.wgsl

@@ -0,0 +1,8 @@
+@group(0) @binding(0) var<storage, read_only> read_ssbo: array<u32>;
+@group(1) @binding(0) var<storage, read_write> write_ssbo: array<u32>;
+
+@compute @workgroup_size(16, 16, 1)
+fn main(@builtin(global_invocation_id) grid: vec3<u32>)
+{
+    write_ssbo[grid.x * grid.y] = read_ssbo[grid.x * grid.y];
+}

Tests/Shaders/WebGPU/ReadOnlyBindings.wgsl

@@ -0,0 +1,7 @@
+@group(0) @binding(0) var<storage, read_only> read_ssbo: array<u32>;
+@group(0) @binding(1) var read_texture: texture_storage_2d<rgba8unorm, read>;
+
+@compute @workgroup_size(16, 16, 1)
+fn main(@builtin(global_invocation_id) grid: vec3<u32>)
+{
+}

Tests/Shaders/WebGPU/ReadWriteBindings.wgsl

@@ -0,0 +1,9 @@
+@group(0) @binding(0) var<storage, read_only> read_ssbo: array<u32>;
+@group(0) @binding(1) var read_texture: texture_storage_2d<rgba8unorm, read>;
+@group(1) @binding(0) var<storage, read_write> write_ssbo: array<u32>;
+@group(1) @binding(1) var write_texture: texture_storage_2d<rgba8unorm, read_write>;
+
+@compute @workgroup_size(16, 16, 1)
+fn main(@builtin(global_invocation_id) grid: vec3<u32>)
+{
+}

Tests/Shaders/WebGPU/Simple.wgsl

@@ -0,0 +1,4 @@
+@compute @workgroup_size(16, 16, 1)
+fn main(@builtin(global_invocation_id) grid: vec3<u32>)
+{
+}

Tests/Shaders/WebGPU/SimpleBufferWrite.wgsl

@@ -0,0 +1,7 @@
+@group(1) @binding(0) var<storage, read_write> write_ssbo: array<u32>;
+
+@compute @workgroup_size(16, 16, 1)
+fn main(@builtin(global_invocation_id) grid: vec3<u32>)
+{
+    write_ssbo[grid.x * grid.y] = u32(0xFFFFFFFF);
+}

Tests/Shaders/WebGPU/WriteOnlyBindings.wgsl

@@ -0,0 +1,7 @@
+@group(1) @binding(0) var<storage, read_write> write_ssbo: array<u32>;
+@group(1) @binding(1) var write_texture: texture_storage_2d<rgba8unorm, write>;
+
+@compute @workgroup_size(16, 16, 1)
+fn main(@builtin(global_invocation_id) grid: vec3<u32>)
+{
+}

Tests/Vulkan/Backend.c

@@ -1,52 +0,0 @@
-#include "Common.h"
-
-#include <unity/unity.h>
-#include <Pulse.h>
-
-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);
-}

Tests/Vulkan/Shaders/.gitignore

@@ -1 +0,0 @@
-*.spv.h

Tests/Vulkan/xmake.lua

@@ -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

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

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