working on WebGPU support

Examples/WebGPU/main.c

@@ -0,0 +1,38 @@
+#include <Pulse.h>
+
+#include <stdio.h>
+#include <stdlib.h>
+
+#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)
+{
+	if(severity == PULSE_DEBUG_MESSAGE_SEVERITY_ERROR)
+	{
+		fprintf(stderr, "Pulse Error: %s\n", message);
+		exit(1);
+	}
+	else if(severity == PULSE_DEBUG_MESSAGE_SEVERITY_WARNING)
+		fprintf(stderr, "Pulse Warning: %s\n", message);
+	else
+		printf("Pulse: %s\n", message);
+}
+
+int main(void)
+{
+	PulseBackend backend = PulseLoadBackend(PULSE_BACKEND_WEBGPU, 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);
+
+	PulseDestroyDevice(device);
+	PulseUnloadBackend(backend);
+	puts("Successfully executed Pulse example using WebGPU !");
+	return 0;
+}

Examples/WebGPU/shader.nzsl

@@ -0,0 +1,13 @@
+[nzsl_version("1.0")]
+module;
+
+struct Input
+{
+    [builtin(global_invocation_indices)] indices: vec3[u32]
+}
+
+[entry(compute)]
+[workgroup(32, 32, 1)]
+fn main(input: Input)
+{
+}

Examples/WebGPU/shader.spv.h

@@ -0,0 +1,16 @@
+3,2,35,7,0,0,1,0,39,0,0,0,16,0,0,0,0,0,0,0,17,0,2,0,1,0,0,0,14,0,
+3,0,0,0,0,0,1,0,0,0,15,0,6,0,5,0,0,0,12,0,0,0,109,97,105,110,0,0,0,0,
+6,0,0,0,16,0,6,0,12,0,0,0,17,0,0,0,32,0,0,0,32,0,0,0,1,0,0,0,3,0,
+3,0,9,0,0,0,100,0,0,0,5,0,4,0,10,0,0,0,73,110,112,117,116,0,0,0,6,0,5,0,
+10,0,0,0,0,0,0,0,105,110,100,105,99,101,115,0,5,0,9,0,6,0,0,0,103,108,111,98,97,108,
+95,105,110,118,111,99,97,116,105,111,110,95,105,110,100,105,99,101,115,0,0,0,5,0,4,0,12,0,0,0,
+109,97,105,110,0,0,0,0,71,0,4,0,6,0,0,0,11,0,0,0,28,0,0,0,72,0,5,0,10,0,
+0,0,0,0,0,0,35,0,0,0,0,0,0,0,19,0,2,0,1,0,0,0,33,0,3,0,2,0,0,0,
+1,0,0,0,21,0,4,0,3,0,0,0,32,0,0,0,0,0,0,0,23,0,4,0,4,0,0,0,3,0,
+0,0,3,0,0,0,32,0,4,0,5,0,0,0,1,0,0,0,4,0,0,0,21,0,4,0,7,0,0,0,
+32,0,0,0,1,0,0,0,43,0,4,0,7,0,0,0,8,0,0,0,0,0,0,0,32,0,4,0,9,0,
+0,0,7,0,0,0,4,0,0,0,30,0,3,0,10,0,0,0,4,0,0,0,32,0,4,0,11,0,0,0,
+7,0,0,0,10,0,0,0,59,0,4,0,5,0,0,0,6,0,0,0,1,0,0,0,54,0,5,0,1,0,
+0,0,12,0,0,0,0,0,0,0,2,0,0,0,248,0,2,0,13,0,0,0,59,0,4,0,11,0,0,0,
+14,0,0,0,7,0,0,0,65,0,5,0,9,0,0,0,15,0,0,0,14,0,0,0,8,0,0,0,63,0,
+3,0,15,0,0,0,6,0,0,0,253,0,1,0,56,0,1,0

Examples/WebGPU/xmake.lua

@@ -0,0 +1,99 @@
+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()
+
+target("WebGPUExample")
+	add_deps("pulse_gpu")
+	if is_plat("linux") then
+		set_extension(".x86_64")
+	end
+	add_rules("compile_shaders")
+	add_files("*.c")
+	add_files("*.nzsl")
+target_end()

Includes/Pulse.h

@@ -38,6 +38,7 @@ typedef enum PulseBackendBits
 	PULSE_BACKEND_ANY     = PULSE_BIT(2),
 	PULSE_BACKEND_VULKAN  = PULSE_BIT(3),
 	PULSE_BACKEND_METAL   = PULSE_BIT(4),
+	PULSE_BACKEND_WEBGPU  = PULSE_BIT(5),
 	// More to come
 } PulseBackendBits;
 typedef PulseFlags PulseBackendFlags;
@@ -67,9 +68,10 @@ typedef PulseFlags PulseImageUsageFlags;
 
 typedef enum PulseShaderFormatsBits
 {
-	PULSE_SHADER_FORMAT_SPIRV_BIT    = PULSE_BIT(1), // Can be used by Vulkan
+	PULSE_SHADER_FORMAT_SPIRV_BIT    = PULSE_BIT(1), // Can be used by Vulkan and WebGPU
 	PULSE_SHADER_FORMAT_MSL_BIT      = PULSE_BIT(2), // Can be used by Metal
 	PULSE_SHADER_FORMAT_METALLIB_BIT = PULSE_BIT(3), // Can be used by Metal
+	PULSE_SHADER_FORMAT_WGSL_BIT     = PULSE_BIT(4), // Can be used by WebGPU
 	// More to come
 } PulseShaderFormatsBits;
 typedef PulseFlags PulseShaderFormatsFlags;

Includes/PulseProfile.h

@@ -35,6 +35,8 @@ extern "C" {
 	#define PULSE_COMPILER_INTEL
 #elif defined(_MSC_VER)
 	#define PULSE_COMPILER_MSVC
+#elif __EMSCRIPTEN__
+	#define PULSE_COMPILER_EMSCRIPTEN
 #else
 	#define PULSE_COMPILER_UNKNOWN
 	#warning "This compiler is not fully supported"

README.md

@@ -5,7 +5,7 @@
 [![Msys2 build](https://github.com/ft-grmhd/Pulse/actions/workflows/msys2-build.yml/badge.svg)](https://github.com/ft-grmhd/Pulse/actions/workflows/msys2-build.yml)
 [![Windows build](https://github.com/ft-grmhd/Pulse/actions/workflows/windows-build.yml/badge.svg)](https://github.com/ft-grmhd/Pulse/actions/workflows/windows-build.yml)
 
-Pulse is a low level GPGPU library designed for highly intensive general GPU computations with high control over the hardware. It is built on top of Vulkan and a Metal support is in discussion.
+Pulse is a low level GPGPU library designed for highly intensive general GPU computations with high control over the hardware. It is built on top of Vulkan. A Metal and WebGPU backends are in development.
 
 ```cpp
 #include <Pulse.h>
@@ -36,7 +36,7 @@ int main(void)
 	PulseCommandList cmd = PulseRequestCommandList(device, PULSE_COMMAND_LIST_GENERAL);
 
 	PulseComputePass pass = PulseBeginComputePass(cmd);
-		PulseBindStorageBuffers(pass, 0, &buffer, 1);
+		PulseBindStorageBuffers(pass, &buffer, 1);
 		PulseBindComputePipeline(pass, pipeline);
 		PulseDispatchComputations(pass, 32, 32, 1);
 	PulseEndComputePass(pass);

Sources/Backends/Vulkan/Vulkan.c

@@ -29,6 +29,7 @@ PulseBackendFlags VulkanCheckSupport(PulseBackendFlags candidates, PulseShaderFo
 	VulkanInstance instance;
 	if(!VulkanInitInstance(PULSE_NULL_HANDLE, &instance, PULSE_NO_DEBUG)) // Instance creation will fail if Vulkan is not supported
 	{
+		VulkanLoaderShutdown();
 		PulseGetLastErrorType(); // Clearing out the errors set by the failed instance creation
 		return PULSE_BACKEND_INVALID;
 	}
@@ -44,7 +45,10 @@ bool VulkanLoadBackend(PulseBackend backend, PulseDebugLevel debug_level)
 	VulkanDriverData* driver_data = (VulkanDriverData*)calloc(1, sizeof(VulkanDriverData));
 	PULSE_CHECK_ALLOCATION_RETVAL(driver_data, false);
 	if(!VulkanInitInstance(backend, &driver_data->instance, debug_level))
+	{
+		free(driver_data);
 		return false;
+	}
 	VulkanDriver.driver_data = driver_data;
 	return true;
 }
@@ -53,6 +57,7 @@ void VulkanUnloadBackend(PulseBackend backend)
 {
 	VulkanDestroyInstance(&VULKAN_RETRIEVE_DRIVER_DATA_AS(backend, VulkanDriverData*)->instance);
 	VulkanLoaderShutdown();
+	free(backend->driver_data);
 }
 
 const char* VulkanVerbaliseResult(VkResult res)

Sources/Backends/Vulkan/VulkanDevice.c

@@ -46,6 +46,7 @@ static int32_t VulkanScorePhysicalDevice(VulkanInstance* instance, VkPhysicalDev
 			break;
 		}
 	}
+	free(props);
 	if(are_there_required_device_extensions == false)
 		return -1;
 
@@ -76,7 +77,6 @@ static bool VulkanIsDeviceForbidden(VkPhysicalDevice device, PulseDevice* forbid
 {
 	if(device == VK_NULL_HANDLE)
 		return true;
-	#pragma omp parallel for
 	for(uint32_t i = 0; i < forbiden_devices_count; i++)
 	{
 		if(device == ((VulkanDevice*)forbiden_devices[i]->driver_data)->physical)
@@ -97,7 +97,6 @@ static VkPhysicalDevice VulkanPickPhysicalDevice(VulkanInstance* instance, Pulse
 	PULSE_CHECK_ALLOCATION_RETVAL(devices, VK_NULL_HANDLE);
 	instance->vkEnumeratePhysicalDevices(instance->instance, &device_count, devices);
 
-	#pragma omp parallel for
 	for(uint32_t i = 0; i < device_count; i++)
 	{
 		if(VulkanIsDeviceForbidden(devices[i], forbiden_devices, forbiden_devices_count))
@@ -120,12 +119,12 @@ PulseDevice VulkanCreateDevice(PulseBackend backend, PulseDevice* forbiden_devic
 	PULSE_CHECK_ALLOCATION_RETVAL(pulse_device, PULSE_NULL_HANDLE);
 
 	VulkanDevice* device = (VulkanDevice*)calloc(1, sizeof(VulkanDevice));
-	PULSE_CHECK_ALLOCATION_RETVAL(device, PULSE_NULLPTR);
+	PULSE_CHECK_ALLOCATION_RETVAL(device, PULSE_NULL_HANDLE);
 
 	VulkanInstance* instance = &VULKAN_RETRIEVE_DRIVER_DATA_AS(backend, VulkanDriverData*)->instance;
 
 	device->physical = VulkanPickPhysicalDevice(instance, forbiden_devices, forbiden_devices_count);
-	PULSE_CHECK_HANDLE_RETVAL(device->physical, PULSE_NULLPTR);
+	PULSE_CHECK_HANDLE_RETVAL(device->physical, PULSE_NULL_HANDLE);
 
 	const float queue_priority = 1.0f;
 
@@ -139,6 +138,8 @@ PulseDevice VulkanCreateDevice(PulseBackend backend, PulseDevice* forbiden_devic
 		if(!VulkanPrepareDeviceQueue(instance, device, i))
 		{
 			PulseSetInternalError(PULSE_ERROR_INITIALIZATION_FAILED);
+			free(device);
+			free(pulse_device);
 			return PULSE_NULLPTR;
 		}
 	}
@@ -185,6 +186,8 @@ PulseDevice VulkanCreateDevice(PulseBackend backend, PulseDevice* forbiden_devic
 		if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(backend))
 			PulseLogInfoFmt(backend, "(Vulkan) Could not load device functions from %s", device->properties.deviceName);
 		PulseSetInternalError(PULSE_ERROR_INITIALIZATION_FAILED);
+		free(device);
+		free(pulse_device);
 		return PULSE_NULLPTR;
 	}
 
@@ -193,6 +196,8 @@ PulseDevice VulkanCreateDevice(PulseBackend backend, PulseDevice* forbiden_devic
 		if(!VulkanRetrieveDeviceQueue(device, (VulkanQueueType)i))
 		{
 			PulseSetInternalError(PULSE_ERROR_INITIALIZATION_FAILED);
+			free(device);
+			free(pulse_device);
 			return PULSE_NULLPTR;
 		}
 	}

Sources/Backends/Vulkan/VulkanImage.c

@@ -2,7 +2,7 @@
 // This file is part of "Pulse"
 // For conditions of distribution and use, see copyright notice in LICENSE
 
-#include "Pulse.h"
+#include <Pulse.h>
 #include "Vulkan.h"
 #include "VulkanImage.h"
 #include "VulkanBuffer.h"

Sources/Backends/WebGPU/WebGPU.c

@@ -0,0 +1,53 @@
+// Copyright (C) 2025 kanel
+// This file is part of "Pulse"
+// For conditions of distribution and use, see copyright notice in LICENSE
+
+#include <Pulse.h>
+#include "../../PulseInternal.h"
+
+#include "WebGPU.h"
+#include "WebGPUDevice.h"
+
+PulseBackendFlags WebGPUCheckSupport(PulseBackendFlags candidates, PulseShaderFormatsFlags shader_formats_used)
+{
+	if(candidates != PULSE_BACKEND_ANY && (candidates & PULSE_BACKEND_WEBGPU) == 0)
+		return PULSE_BACKEND_INVALID;
+	if((shader_formats_used & PULSE_SHADER_FORMAT_SPIRV_BIT) == 0 && (shader_formats_used & PULSE_SHADER_FORMAT_WGSL_BIT) == 0)
+		return PULSE_BACKEND_INVALID;
+
+	WGPUInstance instance = wgpuCreateInstance(PULSE_NULLPTR);
+	if(instance == PULSE_NULLPTR)
+		return PULSE_BACKEND_INVALID;
+	wgpuInstanceRelease(instance);
+	return PULSE_BACKEND_WEBGPU;
+}
+
+bool WebGPULoadBackend(PulseBackend backend, PulseDebugLevel debug_level)
+{
+	WebGPUDriverData* driver_data = (WebGPUDriverData*)calloc(1, sizeof(WebGPUDriverData));
+	PULSE_CHECK_ALLOCATION_RETVAL(driver_data, false);
+	driver_data->instance = wgpuCreateInstance(PULSE_NULLPTR);
+	if(driver_data->instance == PULSE_NULLPTR)
+	{
+		PulseSetInternalError(PULSE_ERROR_INITIALIZATION_FAILED);
+		free(driver_data);
+		return false;
+	}
+	WebGPUDriver.driver_data = driver_data;
+	return true;
+}
+
+void WebGPUUnloadBackend(PulseBackend backend)
+{
+	wgpuInstanceRelease(WEBGPU_RETRIEVE_DRIVER_DATA_AS(backend, WebGPUDriverData*)->instance);
+	free(backend->driver_data);
+}
+
+PulseBackendHandler WebGPUDriver = {
+	.PFN_LoadBackend = WebGPULoadBackend,
+	.PFN_UnloadBackend = WebGPUUnloadBackend,
+	.PFN_CreateDevice = WebGPUCreateDevice,
+	.backend = PULSE_BACKEND_WEBGPU,
+	.supported_shader_formats = PULSE_SHADER_FORMAT_SPIRV_BIT | PULSE_SHADER_FORMAT_WGSL_BIT,
+	.driver_data = PULSE_NULLPTR
+};

Sources/Backends/WebGPU/WebGPU.h

@@ -0,0 +1,26 @@
+// Copyright (C) 2025 kanel
+// This file is part of "Pulse"
+// For conditions of distribution and use, see copyright notice in LICENSE
+
+#include <Pulse.h>
+
+#ifdef PULSE_ENABLE_WEBGPU_BACKEND
+
+#ifndef PULSE_WEBGPU_H_
+#define PULSE_WEBGPU_H_
+
+#include <webgpu/webgpu.h>
+
+#define WEBGPU_RETRIEVE_DRIVER_DATA_AS(handle, cast) ((cast)handle->driver_data)
+
+typedef struct WebGPUDriverData
+{
+	WGPUInstance instance;
+} WebGPUDriverData;
+
+PulseBackendFlags WebGPUCheckSupport(PulseBackendFlags candidates, PulseShaderFormatsFlags shader_formats_used); // Return PULSE_BACKEND_WEBGPU in case of success and PULSE_BACKEND_INVALID otherwise
+
+#endif // PULSE_WEBGPU_H_
+
+#endif // PULSE_ENABLE_WEBGPU_BACKEND
+

Sources/Backends/WebGPU/WebGPUBuffer.c

@@ -0,0 +1,30 @@
+// Copyright (C) 2025 kanel
+// This file is part of "Pulse"
+// For conditions of distribution and use, see copyright notice in LICENSE
+
+#include <Pulse.h>
+#include "WebGPU.h"
+
+PulseBuffer WebGPUCreateBuffer(PulseDevice device, const PulseBufferCreateInfo* create_infos)
+{
+}
+
+bool WebGPUMapBuffer(PulseBuffer buffer, 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)
+{
+}
+
+void WebGPUDestroyBuffer(PulseDevice device, PulseBuffer buffer)
+{
+}

Sources/Backends/WebGPU/WebGPUBuffer.h

@@ -0,0 +1,28 @@
+// Copyright (C) 2025 kanel
+// This file is part of "Pulse"
+// For conditions of distribution and use, see copyright notice in LICENSE
+
+#ifdef PULSE_ENABLE_WEBGPU_BACKEND
+
+#ifndef PULSE_WEBGPU_BUFFER_H_
+#define PULSE_WEBGPU_BUFFER_H_
+
+#include <webgpu/webgpu.h>
+
+#include <Pulse.h>
+#include "../../PulseInternal.h"
+
+typedef struct WebGPUBuffer
+{
+} WebGPUBuffer;
+
+PulseBuffer WebGPUCreateBuffer(PulseDevice device, const PulseBufferCreateInfo* create_infos);
+bool WebGPUMapBuffer(PulseBuffer buffer, 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);
+void WebGPUDestroyBuffer(PulseDevice device, PulseBuffer buffer);
+
+#endif // PULSE_WEBGPU_BUFFER_H_
+
+#endif // PULSE_ENABLE_WEBGPU_BACKEND

Sources/Backends/WebGPU/WebGPUCommandList.c

@@ -0,0 +1,18 @@
+// Copyright (C) 2025 kanel
+// This file is part of "Pulse"
+// For conditions of distribution and use, see copyright notice in LICENSE
+
+#include <Pulse.h>
+#include "WebGPU.h"
+
+PulseCommandList WebGPURequestCommandList(PulseDevice device, PulseCommandListUsage usage)
+{
+}
+
+bool WebGPUSubmitCommandList(PulseDevice device, PulseCommandList cmd, PulseFence fence)
+{
+}
+
+void WebGPUReleaseCommandList(PulseDevice device, PulseCommandList cmd)
+{
+}

Sources/Backends/WebGPU/WebGPUCommandList.h

@@ -0,0 +1,25 @@
+// Copyright (C) 2025 kanel
+// This file is part of "Pulse"
+// For conditions of distribution and use, see copyright notice in LICENSE
+
+#ifdef PULSE_ENABLE_WEBGPU_BACKEND
+
+#ifndef PULSE_WEBGPU_COMMAND_LIST_H_
+#define PULSE_WEBGPU_COMMAND_LIST_H_
+
+#include <webgpu/webgpu.h>
+
+#include <Pulse.h>
+#include "WebGPUBuffer.h"
+
+typedef struct WebGPUCommandList
+{
+} WebGPUCommandList;
+
+PulseCommandList WebGPURequestCommandList(PulseDevice device, PulseCommandListUsage usage);
+bool WebGPUSubmitCommandList(PulseDevice device, PulseCommandList cmd, PulseFence fence);
+void WebGPUReleaseCommandList(PulseDevice device, PulseCommandList cmd);
+
+#endif // PULSE_WEBGPU_COMMAND_LIST_H_
+
+#endif // PULSE_ENABLE_WEBGPU_BACKEND

Sources/Backends/WebGPU/WebGPUComputePass.c

@@ -0,0 +1,42 @@
+// Copyright (C) 2025 kanel
+// This file is part of "Pulse"
+// For conditions of distribution and use, see copyright notice in LICENSE
+
+#include <Pulse.h>
+#include "WebGPU.h"
+
+PulseComputePass WebGPUCreateComputePass(PulseDevice device, PulseCommandList cmd)
+{
+}
+
+void WebGPUDestroyComputePass(PulseDevice device, PulseComputePass pass)
+{
+}
+
+PulseComputePass WebGPUBeginComputePass(PulseCommandList cmd)
+{
+}
+
+void WebGPUEndComputePass(PulseComputePass pass)
+{
+}
+
+void WebGPUBindStorageBuffers(PulseComputePass pass, const PulseBuffer* buffers, uint32_t num_buffers)
+{
+}
+
+void WebGPUBindUniformData(PulseComputePass pass, uint32_t slot, const void* data, uint32_t data_size)
+{
+}
+
+void WebGPUBindStorageImages(PulseComputePass pass, const PulseImage* images, uint32_t num_images)
+{
+}
+
+void WebGPUBindComputePipeline(PulseComputePass pass, PulseComputePipeline pipeline)
+{
+}
+
+void WebGPUDispatchComputations(PulseComputePass pass, uint32_t groupcount_x, uint32_t groupcount_y, uint32_t groupcount_z)
+{
+}

Sources/Backends/WebGPU/WebGPUComputePass.h

@@ -0,0 +1,33 @@
+// Copyright (C) 2025 kanel
+// This file is part of "Pulse"
+// For conditions of distribution and use, see copyright notice in LICENSE
+
+#ifdef PULSE_ENABLE_WEBGPU_BACKEND
+
+#ifndef PULSE_WEBGPU_COMPUTE_PASS_H_
+#define PULSE_WEBGPU_COMPUTE_PASS_H_
+
+#include <webgpu/webgpu.h>
+
+#include <Pulse.h>
+#include "WebGPUBuffer.h"
+#include "WebGPUCommandList.h"
+
+typedef struct WebGPUComputePass
+{
+} WebGPUComputePass;
+
+PulseComputePass WebGPUCreateComputePass(PulseDevice device, PulseCommandList cmd);
+void WebGPUDestroyComputePass(PulseDevice device, PulseComputePass pass);
+
+PulseComputePass WebGPUBeginComputePass(PulseCommandList cmd);
+void WebGPUEndComputePass(PulseComputePass pass);
+void WebGPUBindStorageBuffers(PulseComputePass pass, const PulseBuffer* buffers, uint32_t num_buffers);
+void WebGPUBindUniformData(PulseComputePass pass, uint32_t slot, const void* data, uint32_t data_size);
+void WebGPUBindStorageImages(PulseComputePass pass, const PulseImage* images, uint32_t num_images);
+void WebGPUBindComputePipeline(PulseComputePass pass, PulseComputePipeline pipeline);
+void WebGPUDispatchComputations(PulseComputePass pass, uint32_t groupcount_x, uint32_t groupcount_y, uint32_t groupcount_z);
+
+#endif // PULSE_WEBGPU_COMPUTE_PASS_H_
+
+#endif // PULSE_ENABLE_WEBGPU_BACKEND

Sources/Backends/WebGPU/WebGPUComputePipeline.c

@@ -0,0 +1,14 @@
+// Copyright (C) 2025 kanel
+// This file is part of "Pulse"
+// For conditions of distribution and use, see copyright notice in LICENSE
+
+#include <Pulse.h>
+#include "WebGPU.h"
+
+PulseComputePipeline WebGPUCreateComputePipeline(PulseDevice device, const PulseComputePipelineCreateInfo* info)
+{
+}
+
+void WebGPUDestroyComputePipeline(PulseDevice device, PulseComputePipeline pipeline)
+{
+}

Sources/Backends/WebGPU/WebGPUComputePipeline.h

@@ -0,0 +1,23 @@
+// Copyright (C) 2025 kanel
+// This file is part of "Pulse"
+// For conditions of distribution and use, see copyright notice in LICENSE
+
+#ifdef PULSE_ENABLE_WEBGPU_BACKEND
+
+#ifndef PULSE_WEBGPU_COMPUTE_PIPELINE_H_
+#define PULSE_WEBGPU_COMPUTE_PIPELINE_H_
+
+#include <webgpu/webgpu.h>
+
+#include <Pulse.h>
+
+typedef struct WebGPUComputePipeline
+{
+} WebGPUComputePipeline;
+
+PulseComputePipeline WebGPUCreateComputePipeline(PulseDevice device, const PulseComputePipelineCreateInfo* info);
+void WebGPUDestroyComputePipeline(PulseDevice device, PulseComputePipeline pipeline);
+
+#endif // PULSE_WEBGPU_COMPUTE_PIPELINE_H_
+
+#endif // PULSE_ENABLE_WEBGPU_BACKEND

Sources/Backends/WebGPU/WebGPUDevice.c

@@ -0,0 +1,207 @@
+// Copyright (C) 2025 kanel
+// This file is part of "Pulse"
+// For conditions of distribution and use, see copyright notice in LICENSE
+
+#include <Pulse.h>
+
+#include "../../PulseInternal.h"
+#include "WebGPU.h"
+#include "WebGPUComputePipeline.h"
+#include "WebGPUCommandList.h"
+#include "WebGPUDevice.h"
+#include "WebGPUFence.h"
+#include "WebGPUBuffer.h"
+#include "WebGPUImage.h"
+#include "WebGPUComputePass.h"
+
+#ifndef PULSE_PLAT_WASM
+	#include <wgpu.h>
+#endif
+
+#ifdef PULSE_PLAT_WASM
+	static void WebGPURequestAdapterCallback(WGPURequestAdapterStatus status, WGPUAdapter adapter, WGPUStringView message, void* userdata1, void* userdata2)
+	{
+		WebGPUDevice* device = (WebGPUDevice*)userdata1;
+		PulseBackend backend = (PulseBackend)userdata2;
+		if(status != WGPURequestAdapterStatus_Success)
+		{
+			if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(backend))
+				PulseLogErrorFmt(backend, "(WebGPU) Could not load adapter %.*s", message.length, message.data);
+			PulseSetInternalError(PULSE_ERROR_INITIALIZATION_FAILED);
+			pulse_device->has_error = true;
+			return;
+		}
+		pulse_device->has_error = false;
+		device->adapter = adapter;
+	}
+#else
+	static uint64_t WebGPUScoreAdapter(WGPUInstance instance, WGPUAdapter adapter)
+	{
+		uint64_t score = 0;
+		WGPUAdapterInfo infos;
+		wgpuAdapterGetInfo(adapter, &infos);
+		WGPULimits limits;
+		wgpuAdapterGetLimits(adapter, &limits);
+
+		if(infos.adapterType == WGPUAdapterType_DiscreteGPU)
+			score += 10000;
+
+		switch(infos.backendType)
+		{
+			case WGPUBackendType_Metal:
+			case WGPUBackendType_D3D12:
+			case WGPUBackendType_WebGPU:
+			case WGPUBackendType_Vulkan: score += 1000;
+
+			default: break;
+		}
+
+		score += limits.maxComputeWorkgroupSizeX;
+		score += limits.maxComputeWorkgroupSizeY;
+		score += limits.maxComputeWorkgroupSizeZ;
+		score += limits.maxComputeWorkgroupStorageSize;
+		score += limits.maxComputeInvocationsPerWorkgroup;
+		score += limits.maxComputeWorkgroupsPerDimension;
+		score += limits.maxTextureDimension2D;
+
+		return score;
+	}
+
+	static bool WebGPUIsDeviceForbidden(WGPUAdapter adapter, PulseDevice* forbiden_devices, uint32_t forbiden_devices_count)
+	{
+		if(adapter == PULSE_NULLPTR)
+			return true;
+		for(uint32_t i = 0; i < forbiden_devices_count; i++)
+		{
+			if(adapter == ((WebGPUDevice*)forbiden_devices[i]->driver_data)->adapter)
+				return true;
+		}
+		return false;
+	}
+#endif
+
+static void WebGPURequestDeviceCallback(WGPURequestDeviceStatus status, WGPUDevice device, WGPUStringView message, void* userdata1, void* userdata2)
+{
+	WebGPUDevice* pulse_device = (WebGPUDevice*)userdata1;
+	PulseBackend backend = (PulseBackend)userdata2;
+	if(status != WGPURequestDeviceStatus_Success)
+	{
+		if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(backend))
+			PulseLogErrorFmt(backend, "(WebGPU) Could not create device from %.*s, %.*s", pulse_device->infos.device.length, pulse_device->infos.device.data, message.length, message.data);
+		PulseSetInternalError(PULSE_ERROR_INITIALIZATION_FAILED);
+		pulse_device->has_error = true;
+		return;
+	}
+	pulse_device->has_error = false;
+	pulse_device->device = device;
+}
+
+PulseDevice WebGPUCreateDevice(PulseBackend backend, PulseDevice* forbiden_devices, uint32_t forbiden_devices_count)
+{
+	PULSE_CHECK_HANDLE_RETVAL(backend, PULSE_NULLPTR);
+
+	PulseDevice pulse_device = (PulseDeviceHandler*)calloc(1, sizeof(PulseDeviceHandler));
+	PULSE_CHECK_ALLOCATION_RETVAL(pulse_device, PULSE_NULL_HANDLE);
+
+	WebGPUDevice* device = (WebGPUDevice*)calloc(1, sizeof(WebGPUDevice));
+	PULSE_CHECK_ALLOCATION_RETVAL(device, PULSE_NULL_HANDLE);
+
+	WGPUInstance instance = WEBGPU_RETRIEVE_DRIVER_DATA_AS(backend, WebGPUDriverData*)->instance;
+
+	#ifdef PULSE_PLAT_WASM
+		WGPURequestAdapterOptions adapter_options = { 0 };
+		adapter_options.powerPreference = WGPUPowerPreference_HighPerformance;
+		adapter_options.compatibleSurface = PULSE_NULLPTR;
+		adapter_options.backendType = WGPUBackendType_WebGPU;
+		WGPURequestAdapterCallbackInfo adapter_callback = { 0 };
+		adapter_callback.callback = WebGPURequestAdapterCallback;
+		adapter_callback.mode = WGPUCallbackMode_AllowProcessEvents;
+		adapter_callback.userdata1 = device;
+		adapter_callback.userdata2 = backend;
+		wgpuInstanceRequestAdapter(instance, &adapter_options, adapter_callback);
+
+		while(device->adapter == PULSE_NULLPTR && !device->has_error) // Wait for adapter request
+			PulseSleep(100);
+	#else
+		WGPUInstanceEnumerateAdapterOptions adapter_options = { 0 };
+		adapter_options.nextInChain = PULSE_NULLPTR;
+		adapter_options.backends = WGPUInstanceBackend_All;
+		size_t adapter_count = wgpuInstanceEnumerateAdapters(instance, &adapter_options, PULSE_NULLPTR);
+		WGPUAdapter* adapters = (WGPUAdapter*)calloc(adapter_count, sizeof(WGPUAdapter));
+		wgpuInstanceEnumerateAdapters(instance, &adapter_options, adapters);
+		uint64_t best_device_score = 0;
+		for(uint32_t i = 0; i < adapter_count; i++)
+		{
+			if(WebGPUIsDeviceForbidden(adapters[i], forbiden_devices, forbiden_devices_count))
+				continue;
+			uint64_t current_device_score = WebGPUScoreAdapter(instance, adapters[i]);
+			if(current_device_score > best_device_score)
+			{
+				best_device_score = current_device_score;
+				device->adapter = adapters[i];
+			}
+		}
+	#endif
+
+	if(device->adapter == PULSE_NULLPTR)
+	{
+		if(PULSE_IS_BACKEND_HIGH_LEVEL_DEBUG(backend))
+			PulseLogError(backend, "(WebGPU) could not find suitable adapter");
+		PulseSetInternalError(PULSE_ERROR_INITIALIZATION_FAILED);
+		free(pulse_device);
+		free(device);
+		return PULSE_NULL_HANDLE;
+	}
+
+	wgpuAdapterGetLimits(device->adapter, &device->limits);
+	wgpuAdapterGetInfo(device->adapter, &device->infos);
+
+	WGPUDeviceDescriptor descriptor = { 0 };
+	descriptor.requiredLimits = &device->limits;
+	WGPURequestDeviceCallbackInfo device_callback = { 0 };
+	device_callback.callback = WebGPURequestDeviceCallback;
+	device_callback.mode = WGPUCallbackMode_AllowProcessEvents;
+	device_callback.userdata1 = device;
+	device_callback.userdata2 = backend;
+	wgpuAdapterRequestDevice(device->adapter, &descriptor, device_callback);
+
+	while(device->device == PULSE_NULLPTR && !device->has_error) // Wait for device request
+		PulseSleep(100);
+
+	if(device->has_error)
+	{
+		free(pulse_device);
+		free(device);
+		return PULSE_NULL_HANDLE;
+	}
+
+	pulse_device->driver_data = device;
+	pulse_device->backend = backend;
+	PULSE_LOAD_DRIVER_DEVICE(WebGPU);
+
+	if(PULSE_IS_BACKEND_HIGH_LEVEL_DEBUG(backend))
+	{
+		const char* backends[] = {
+			"Undefined",
+			"Null",
+			"WebGPU",
+			"D3D11",
+			"D3D12",
+			"Metal",
+			"Vulkan",
+			"OpenGL",
+			"OpenGLES",
+		};
+		PulseLogInfoFmt(backend, "(WebGPU) created device from %.*s using backend %s", device->infos.device.length, device->infos.device.data, backends[device->infos.backendType]);
+	}
+	return pulse_device;
+}
+
+void WebGPUDestroyDevice(PulseDevice device)
+{
+	WebGPUDevice* webgpu_device = WEBGPU_RETRIEVE_DRIVER_DATA_AS(device, WebGPUDevice*);
+	if(webgpu_device == PULSE_NULLPTR || webgpu_device->device == PULSE_NULLPTR)
+		return;
+	free(webgpu_device);
+	free(device);
+}

Sources/Backends/WebGPU/WebGPUDevice.h

@@ -0,0 +1,29 @@
+// Copyright (C) 2025 kanel
+// This file is part of "Pulse"
+// For conditions of distribution and use, see copyright notice in LICENSE
+
+#ifdef PULSE_ENABLE_WEBGPU_BACKEND
+
+#ifndef PULSE_WEBGPU_DEVICE_H_
+#define PULSE_WEBGPU_DEVICE_H_
+
+#include <webgpu/webgpu.h>
+
+#include <Pulse.h>
+
+typedef struct WebGPUDevice
+{
+	WGPUAdapterInfo infos;
+	WGPULimits limits;
+	WGPUAdapter adapter;
+	WGPUDevice device;
+
+	bool has_error;
+} WebGPUDevice;
+
+PulseDevice WebGPUCreateDevice(PulseBackend backend, PulseDevice* forbiden_devices, uint32_t forbiden_devices_count);
+void WebGPUDestroyDevice(PulseDevice device);
+
+#endif // PULSE_WEBGPU_DEVICE_H_
+
+#endif // PULSE_ENABLE_WEBGPU_BACKEND

Sources/Backends/WebGPU/WebGPUFence.c

@@ -0,0 +1,22 @@
+// Copyright (C) 2025 kanel
+// This file is part of "Pulse"
+// For conditions of distribution and use, see copyright notice in LICENSE
+
+#include <Pulse.h>
+#include "WebGPU.h"
+
+PulseFence WebGPUCreateFence(PulseDevice device)
+{
+}
+
+void WebGPUDestroyFence(PulseDevice device, PulseFence fence)
+{
+}
+
+bool WebGPUIsFenceReady(PulseDevice device, PulseFence fence)
+{
+}
+
+bool WebGPUWaitForFences(PulseDevice device, const PulseFence* fences, uint32_t fences_count, bool wait_for_all)
+{
+}

Sources/Backends/WebGPU/WebGPUFence.h

@@ -0,0 +1,21 @@
+// Copyright (C) 2025 kanel
+// This file is part of "Pulse"
+// For conditions of distribution and use, see copyright notice in LICENSE
+
+#ifdef PULSE_ENABLE_WEBGPU_BACKEND
+
+#ifndef PULSE_WEBGPU_FENCE_H_
+#define PULSE_WEBGPU_FENCE_H_
+
+#include <webgpu/webgpu.h>
+
+#include <Pulse.h>
+
+PulseFence WebGPUCreateFence(PulseDevice device);
+void WebGPUDestroyFence(PulseDevice device, PulseFence fence);
+bool WebGPUIsFenceReady(PulseDevice device, PulseFence fence);
+bool WebGPUWaitForFences(PulseDevice device, const PulseFence* fences, uint32_t fences_count, bool wait_for_all);
+
+#endif // PULSE_WEBGPU_FENCE_H_
+
+#endif // PULSE_ENABLE_WEBGPU_BACKEND

Sources/Backends/WebGPU/WebGPUImage.c

@@ -0,0 +1,26 @@
+// Copyright (C) 2025 kanel
+// This file is part of "Pulse"
+// For conditions of distribution and use, see copyright notice in LICENSE
+
+#include <Pulse.h>
+#include "WebGPU.h"
+
+PulseImage WebGPUCreateImage(PulseDevice device, const PulseImageCreateInfo* create_infos)
+{
+}
+
+bool WebGPUIsImageFormatValid(PulseDevice device, PulseImageFormat format, PulseImageType type, PulseImageUsageFlags usage)
+{
+}
+
+bool WebGPUCopyImageToBuffer(PulseCommandList cmd, const PulseImageRegion* src, const PulseBufferRegion* dst)
+{
+}
+
+bool WebGPUBlitImage(PulseCommandList cmd, const PulseImageRegion* src, const PulseImageRegion* dst)
+{
+}
+
+void WebGPUDestroyImage(PulseDevice device, PulseImage image)
+{
+}

Sources/Backends/WebGPU/WebGPUImage.h

@@ -0,0 +1,27 @@
+// Copyright (C) 2025 kanel
+// This file is part of "Pulse"
+// For conditions of distribution and use, see copyright notice in LICENSE
+
+#ifdef PULSE_ENABLE_WEBGPU_BACKEND
+
+#ifndef PULSE_WEBGPU_IMAGE_H_
+#define PULSE_WEBGPU_IMAGE_H_
+
+#include <webgpu/webgpu.h>
+
+#include <Pulse.h>
+#include "../../PulseInternal.h"
+
+typedef struct WebGPUImage
+{
+} WebGPUImage;
+
+PulseImage WebGPUCreateImage(PulseDevice device, const PulseImageCreateInfo* create_infos);
+bool WebGPUIsImageFormatValid(PulseDevice device, PulseImageFormat format, PulseImageType type, PulseImageUsageFlags usage);
+bool WebGPUCopyImageToBuffer(PulseCommandList cmd, const PulseImageRegion* src, const PulseBufferRegion* dst);
+bool WebGPUBlitImage(PulseCommandList cmd, const PulseImageRegion* src, const PulseImageRegion* dst);
+void WebGPUDestroyImage(PulseDevice device, PulseImage image);
+
+#endif // PULSE_WEBGPU_IMAGE_H_
+
+#endif // PULSE_ENABLE_WEBGPU_BACKEND

Sources/PulseBackend.c

@@ -14,6 +14,9 @@
 #ifdef PULSE_ENABLE_METAL_BACKEND
 	#include "Backends/Metal/Metal.h"
 #endif
+#ifdef PULSE_ENABLE_WEBGPU_BACKEND
+	#include "Backends/WebGPU/WebGPU.h"
+#endif
 
 // Ordered by default preference
 static const PulseCheckBackendSupportPFN backends_supports[] = {
@@ -23,6 +26,9 @@ static const PulseCheckBackendSupportPFN backends_supports[] = {
 	#ifdef PULSE_ENABLE_METAL_BACKEND
 		MetalCheckSupport,
 	#endif
+	#ifdef PULSE_ENABLE_WEBGPU_BACKEND
+		WebGPUCheckSupport,
+	#endif
 	PULSE_NULLPTR
 };
 
@@ -90,6 +96,9 @@ static PulseBackend PulseGetBackendFromFlag(PulseBackendBits flag)
 		#ifdef PULSE_ENABLE_METAL_BACKEND
 			case PULSE_BACKEND_METAL: return &MetalDriver;
 		#endif
+		#ifdef PULSE_ENABLE_WEBGPU_BACKEND
+			case PULSE_BACKEND_WEBGPU: return &WebGPUDriver;
+		#endif
 
 		default: break;
 	}

Sources/PulseInternal.c

@@ -0,0 +1,19 @@
+// Copyright (C) 2025 kanel
+// This file is part of "Pulse"
+// For conditions of distribution and use, see copyright notice in LICENSE
+
+#include "PulseInternal.h"
+
+#include <threads.h>
+
+PulseThreadID PulseGetThreadID()
+{
+	return (PulseThreadID)thrd_current();
+}
+
+void PulseSleep(int32_t ms)
+{
+	if(ms <= 0)
+		return;
+	thrd_sleep(&(struct timespec){ .tv_sec = ms / 1000, .tv_nsec = (ms % 1000) * 1000000 }, PULSE_NULLPTR);
+}

Sources/PulseInternal.h

@@ -155,6 +155,7 @@ typedef struct PulseComputePassHandler
 } PulseComputePassHandler;
 
 PulseThreadID PulseGetThreadID();
+void PulseSleep(int32_t ms);
 
 void PulseSetInternalError(PulseErrorType error);
 
@@ -174,5 +175,8 @@ void PulseLogBackend(PulseBackend backend, PulseDebugMessageSeverity type, const
 #ifdef PULSE_ENABLE_METAL_BACKEND
 	extern PulseBackendHandler MetalDriver;
 #endif // PULSE_ENABLE_METAL_BACKEND
+#ifdef PULSE_ENABLE_WEBGPU_BACKEND
+	extern PulseBackendHandler WebGPUDriver;
+#endif // PULSE_ENABLE_WEBGPU_BACKEND
 
 #endif // PULSE_INTERNAL_H_

Sources/PulsePosix.c

@@ -1,15 +0,0 @@
-// Copyright (C) 2025 kanel
-// This file is part of "Pulse"
-// For conditions of distribution and use, see copyright notice in LICENSE
-
-#include "PulseInternal.h"
-#include "PulsePosix.h"
-
-#ifdef PULSE_PLAT_POSIX
-
-PulseThreadID PulseGetThreadID()
-{
-	return (PulseThreadID)pthread_self();
-}
-
-#endif // PULSE_PLAT_POSIX

Sources/PulsePosix.h

@@ -1,17 +0,0 @@
-// Copyright (C) 2025 kanel
-// This file is part of "Pulse"
-// For conditions of distribution and use, see copyright notice in LICENSE
-
-#ifndef PULSE_POSIX_H_
-#define PULSE_POSIX_H_
-
-#include <PulseProfile.h>
-#include "PulseInternal.h"
-
-#ifdef PULSE_PLAT_POSIX
-
-#include <pthread.h>
-
-#endif // PULSE_PLAT_WINDOWS
-
-#endif // PULSE_POSIX_H_

Sources/PulseWindows.c

@@ -1,15 +0,0 @@
-// Copyright (C) 2025 kanel
-// This file is part of "Pulse"
-// For conditions of distribution and use, see copyright notice in LICENSE
-
-#include "PulseInternal.h"
-#include "PulseWindows.h"
-
-#ifdef PULSE_PLAT_WINDOWS
-
-PulseThreadID PulseGetThreadID()
-{
-	return (PulseThreadID)GetCurrentThreadId();
-}
-
-#endif // PULSE_PLAT_WINDOWS

Sources/PulseWindows.h

@@ -1,18 +0,0 @@
-// Copyright (C) 2025 kanel
-// This file is part of "Pulse"
-// For conditions of distribution and use, see copyright notice in LICENSE
-
-#ifndef PULSE_WINDOWS_H_
-#define PULSE_WINDOWS_H_
-
-#include <PulseProfile.h>
-#include "PulseInternal.h"
-
-#ifdef PULSE_PLAT_WINDOWS
-
-#define NOMINMAX
-#include <windows.h>
-
-#endif // PULSE_PLAT_WINDOWS
-
-#endif // PULSE_WINDOWS_H_

xmake.lua

@@ -5,7 +5,7 @@
 local backends = {
 	Vulkan = {
 		option = "vulkan",
-		default = true,
+		default = not is_plat("wasm"),
 		packages = { "vulkan-headers", "vulkan-memory-allocator" },
 		custom = function()
 			add_defines("VK_NO_PROTOTYPES")
@@ -18,6 +18,16 @@ local backends = {
 		custom = function()
 			add_files("Sources/Backends/Metal/**.m")
 		end
+	},
+	WebGPU = {
+		option = "webgpu",
+		packages = { "wgpu-native" },
+		default = is_plat("wasm"),
+		custom = function()
+			if is_plat("wasm") then
+				add_defines("PULSE_PLAT_WASM")
+			end
+		end
 	}
 }
 
@@ -58,6 +68,10 @@ includes("Xmake/**.lua")
 
 option("unitybuild", { description = "Build the library using unity build", default = false })
 
+if is_plat("wasm") then
+	backends.Vulkan = nil
+end
+
 for name, module in pairs(backends) do
 	if has_config(module.option) then
 		if module.packages then