working on WebGPU command lists

Examples/WebGPU/main.c

@@ -25,7 +25,7 @@ void DebugCallBack(PulseDebugMessageSeverity severity, const char* message)
 
 int main(void)
 {
-	PulseBackend backend = PulseLoadBackend(PULSE_BACKEND_WEBGPU, PULSE_SHADER_FORMAT_SPIRV_BIT, PULSE_HIGH_DEBUG);
+	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);

Examples/WebGPU/shader.nzsl

@@ -1,13 +0,0 @@
-[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

@@ -1,16 +0,0 @@
-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/shader.wgsl

@@ -0,0 +1,4 @@
+@compute @workgroup_size(32, 32, 1)
+fn main(@builtin(global_invocation_id) grid: vec3u)
+{
+}

Examples/WebGPU/xmake.lua

@@ -1,99 +1,7 @@
-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

@@ -68,7 +68,7 @@ typedef PulseFlags PulseImageUsageFlags;
 
 typedef enum PulseShaderFormatsBits
 {
-	PULSE_SHADER_FORMAT_SPIRV_BIT    = PULSE_BIT(1), // Can be used by Vulkan and WebGPU
+	PULSE_SHADER_FORMAT_SPIRV_BIT    = PULSE_BIT(1), // Can be used by Vulkan
 	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

Sources/Backends/Vulkan/VulkanCommandList.c

@@ -2,7 +2,6 @@
 // This file is part of "Pulse"
 // For conditions of distribution and use, see copyright notice in LICENSE
 
-#include "Pulse.h"
 #include "Vulkan.h"
 #include "VulkanCommandList.h"
 #include "VulkanCommandPool.h"
@@ -98,8 +97,6 @@ PulseCommandList VulkanRequestCommandList(PulseDevice device, PulseCommandListUs
 
 bool VulkanSubmitCommandList(PulseDevice device, PulseCommandList cmd, PulseFence fence)
 {
-	PULSE_CHECK_HANDLE_RETVAL(cmd, false);
-
 	VulkanDevice* vulkan_device = VULKAN_RETRIEVE_DRIVER_DATA_AS(device, VulkanDevice*);
 	VulkanCommandList* vulkan_cmd = VULKAN_RETRIEVE_DRIVER_DATA_AS(cmd, VulkanCommandList*);
 

Sources/Backends/WebGPU/WebGPU.c

@@ -12,7 +12,7 @@ PulseBackendFlags WebGPUCheckSupport(PulseBackendFlags candidates, PulseShaderFo
 {
 	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)
+	if((shader_formats_used & PULSE_SHADER_FORMAT_WGSL_BIT) == 0)
 		return PULSE_BACKEND_INVALID;
 
 	WGPUInstance instance = wgpuCreateInstance(PULSE_NULLPTR);
@@ -48,6 +48,6 @@ PulseBackendHandler WebGPUDriver = {
 	.PFN_UnloadBackend = WebGPUUnloadBackend,
 	.PFN_CreateDevice = WebGPUCreateDevice,
 	.backend = PULSE_BACKEND_WEBGPU,
-	.supported_shader_formats = PULSE_SHADER_FORMAT_SPIRV_BIT | PULSE_SHADER_FORMAT_WGSL_BIT,
+	.supported_shader_formats = PULSE_SHADER_FORMAT_WGSL_BIT,
 	.driver_data = PULSE_NULLPTR
 };

Sources/Backends/WebGPU/WebGPUCommandList.c

@@ -4,15 +4,62 @@
 
 #include <Pulse.h>
 #include "WebGPU.h"
+#include "WebGPUDevice.h"
+#include "WebGPUCommandList.h"
+#include "WebGPUComputePass.h"
+#include "../../PulseInternal.h"
 
 PulseCommandList WebGPURequestCommandList(PulseDevice device, PulseCommandListUsage usage)
 {
+	PULSE_CHECK_HANDLE_RETVAL(device, PULSE_NULL_HANDLE);
+
+	PULSE_UNUSED(usage);
+
+	WebGPUDevice* webgpu_device = WEBGPU_RETRIEVE_DRIVER_DATA_AS(device, WebGPUDevice*);
+
+	PulseCommandList cmd = (PulseCommandList)calloc(1, sizeof(PulseCommandListHandler));
+	PULSE_CHECK_ALLOCATION_RETVAL(cmd, PULSE_NULL_HANDLE);
+
+	WebGPUCommandList* webgpu_cmd = (WebGPUCommandList*)calloc(1, sizeof(WebGPUCommandList));
+	PULSE_CHECK_ALLOCATION_RETVAL(webgpu_cmd, PULSE_NULL_HANDLE);
+
+	WGPUCommandEncoderDescriptor encoder_descriptor = { 0 };
+	webgpu_cmd->encoder = wgpuDeviceCreateCommandEncoder(webgpu_device->device, &encoder_descriptor);
+
+	cmd->usage = usage;
+	cmd->device = device;
+	cmd->driver_data = webgpu_cmd;
+	cmd->thread_id = PulseGetThreadID();
+
+	cmd->pass = WebGPUCreateComputePass(device, cmd);
+	cmd->state = PULSE_COMMAND_LIST_STATE_RECORDING;
+	cmd->is_available = false;
+
+	return cmd;
 }
 
 bool WebGPUSubmitCommandList(PulseDevice device, PulseCommandList cmd, PulseFence fence)
 {
+	WebGPUDevice* webgpu_device = WEBGPU_RETRIEVE_DRIVER_DATA_AS(device, WebGPUDevice*);
+	WebGPUCommandList* webgpu_cmd = WEBGPU_RETRIEVE_DRIVER_DATA_AS(cmd, WebGPUCommandList*);
+
+	WGPUCommandBufferDescriptor command_buffer_descriptor = { 0 };
+	WGPUCommandBuffer command_buffer = wgpuCommandEncoderFinish(webgpu_cmd->encoder, &command_buffer_descriptor);
+
+	wgpuQueueSubmit(webgpu_device->queue, 1, &command_buffer);
+
+	wgpuCommandBufferRelease(command_buffer);
+	return true;
 }
 
 void WebGPUReleaseCommandList(PulseDevice device, PulseCommandList cmd)
 {
+	PULSE_CHECK_HANDLE(device);
+
+	WebGPUCommandList* webgpu_cmd = WEBGPU_RETRIEVE_DRIVER_DATA_AS(cmd, WebGPUCommandList*);
+
+	wgpuCommandEncoderRelease(webgpu_cmd->encoder);
+
+	free(webgpu_cmd);
+	free(cmd);
 }

Sources/Backends/WebGPU/WebGPUCommandList.h

@@ -14,6 +14,7 @@
 
 typedef struct WebGPUCommandList
 {
+	WGPUCommandEncoder encoder;
 } WebGPUCommandList;
 
 PulseCommandList WebGPURequestCommandList(PulseDevice device, PulseCommandListUsage usage);

Sources/Backends/WebGPU/WebGPUDevice.c

@@ -13,6 +13,7 @@
 #include "WebGPUBuffer.h"
 #include "WebGPUImage.h"
 #include "WebGPUComputePass.h"
+#include "webgpu.h"
 
 #ifndef PULSE_PLAT_WASM
 	#include <wgpu.h>
@@ -26,12 +27,12 @@
 		if(status != WGPURequestAdapterStatus_Success)
 		{
 			if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(backend))
-				PulseLogErrorFmt(backend, "(WebGPU) Could not load adapter %.*s", message.length, message.data);
+				PulseLogErrorFmt(backend, "(WebGPU) could not load adapter, %.*s", message.length, message.data);
 			PulseSetInternalError(PULSE_ERROR_INITIALIZATION_FAILED);
-			pulse_device->has_error = true;
+			device->has_error = true;
 			return;
 		}
-		pulse_device->has_error = false;
+		device->has_error = false;
 		device->adapter = adapter;
 	}
 #else
@@ -66,20 +67,20 @@
 
 		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 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;
+}
+
 static void WebGPURequestDeviceCallback(WGPURequestDeviceStatus status, WGPUDevice device, WGPUStringView message, void* userdata1, void* userdata2)
 {
 	WebGPUDevice* pulse_device = (WebGPUDevice*)userdata1;
@@ -87,7 +88,7 @@ static void WebGPURequestDeviceCallback(WGPURequestDeviceStatus status, WGPUDevi
 	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);
+			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;
@@ -96,6 +97,35 @@ static void WebGPURequestDeviceCallback(WGPURequestDeviceStatus status, WGPUDevi
 	pulse_device->device = device;
 }
 
+static void WebGPUDeviceLostCallback(const WGPUDevice* _, WGPUDeviceLostReason reason, WGPUStringView message, void* userdata1, void* userdata2)
+{
+	WebGPUDevice* device = (WebGPUDevice*)userdata1;
+	PulseBackend backend = (PulseBackend)userdata2;
+	const char* reasons[] = {
+		"of unknown reason",
+		"device has been destroyed",
+		"instance have been dropped",
+		"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);
+}
+
+static void WebGPUDeviceUncapturedErrorCallback(const WGPUDevice* _, WGPUErrorType type, WGPUStringView message, void* userdata1, void* userdata2)
+{
+	WebGPUDevice* device = (WebGPUDevice*)userdata1;
+	PulseBackend backend = (PulseBackend)userdata2;
+	const char* types[] = {
+		"has recieved no error",
+		"has recieved a validation error",
+		"is out of memory",
+		"has recieved an internal error",
+		"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);
+}
+
 PulseDevice WebGPUCreateDevice(PulseBackend backend, PulseDevice* forbiden_devices, uint32_t forbiden_devices_count)
 {
 	PULSE_CHECK_HANDLE_RETVAL(backend, PULSE_NULLPTR);
@@ -112,7 +142,7 @@ PulseDevice WebGPUCreateDevice(PulseBackend backend, PulseDevice* forbiden_devic
 		WGPURequestAdapterOptions adapter_options = { 0 };
 		adapter_options.powerPreference = WGPUPowerPreference_HighPerformance;
 		adapter_options.compatibleSurface = PULSE_NULLPTR;
-		adapter_options.backendType = WGPUBackendType_WebGPU;
+		adapter_options.backendType = WGPUBackendType_Undefined;
 		WGPURequestAdapterCallbackInfo adapter_callback = { 0 };
 		adapter_callback.callback = WebGPURequestAdapterCallback;
 		adapter_callback.mode = WGPUCallbackMode_AllowProcessEvents;
@@ -122,6 +152,16 @@ PulseDevice WebGPUCreateDevice(PulseBackend backend, PulseDevice* forbiden_devic
 
 		while(device->adapter == PULSE_NULLPTR && !device->has_error) // Wait for adapter request
 			PulseSleep(100);
+		
+		if(device->adapter != PULSE_NULLPTR && WebGPUIsDeviceForbidden(device->adapter, forbiden_devices, forbiden_devices_count))
+		{
+			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;
+		}
 	#else
 		WGPUInstanceEnumerateAdapterOptions adapter_options = { 0 };
 		adapter_options.nextInChain = PULSE_NULLPTR;
@@ -156,8 +196,19 @@ PulseDevice WebGPUCreateDevice(PulseBackend backend, PulseDevice* forbiden_devic
 	wgpuAdapterGetLimits(device->adapter, &device->limits);
 	wgpuAdapterGetInfo(device->adapter, &device->infos);
 
+	WGPUDeviceLostCallbackInfo lost_callback = { 0 };
+	lost_callback.callback = WebGPUDeviceLostCallback;
+	lost_callback.mode = WGPUCallbackMode_AllowProcessEvents;
+	lost_callback.userdata1 = device;
+	lost_callback.userdata2 = backend;
+	WGPUUncapturedErrorCallbackInfo uncaptured_callback = { 0 };
+	uncaptured_callback.callback = WebGPUDeviceUncapturedErrorCallback;
+	uncaptured_callback.userdata1 = device;
+	uncaptured_callback.userdata2 = backend;
 	WGPUDeviceDescriptor descriptor = { 0 };
 	descriptor.requiredLimits = &device->limits;
+	descriptor.deviceLostCallbackInfo = lost_callback;
+	descriptor.uncapturedErrorCallbackInfo = uncaptured_callback;
 	WGPURequestDeviceCallbackInfo device_callback = { 0 };
 	device_callback.callback = WebGPURequestDeviceCallback;
 	device_callback.mode = WGPUCallbackMode_AllowProcessEvents;
@@ -175,6 +226,8 @@ PulseDevice WebGPUCreateDevice(PulseBackend backend, PulseDevice* forbiden_devic
 		return PULSE_NULL_HANDLE;
 	}
 
+	device->queue = wgpuDeviceGetQueue(device->device);
+
 	pulse_device->driver_data = device;
 	pulse_device->backend = backend;
 	PULSE_LOAD_DRIVER_DEVICE(WebGPU);
@@ -192,7 +245,7 @@ PulseDevice WebGPUCreateDevice(PulseBackend backend, PulseDevice* forbiden_devic
 			"OpenGL",
 			"OpenGLES",
 		};
-		PulseLogInfoFmt(backend, "(WebGPU) created device from %.*s using backend %s", device->infos.device.length, device->infos.device.data, backends[device->infos.backendType]);
+		PulseLogInfoFmt(backend, "(WebGPU) created device from %.*s using %s backend", device->infos.device.length, device->infos.device.data, backends[device->infos.backendType]);
 	}
 	return pulse_device;
 }
@@ -202,6 +255,12 @@ 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;
+	wgpuQueueRelease(webgpu_device->queue);
+	wgpuDeviceRelease(webgpu_device->device);
+	wgpuAdapterRelease(webgpu_device->adapter);
+	if(PULSE_IS_BACKEND_HIGH_LEVEL_DEBUG(device->backend))
+		PulseLogInfoFmt(device->backend, "(WebGPU) destroyed device created from %.*s", webgpu_device->infos.device.length, webgpu_device->infos.device.data);
+	wgpuAdapterInfoFreeMembers(webgpu_device->infos);
 	free(webgpu_device);
 	free(device);
 }

Sources/Backends/WebGPU/WebGPUDevice.h

@@ -17,6 +17,7 @@ typedef struct WebGPUDevice
 	WGPULimits limits;
 	WGPUAdapter adapter;
 	WGPUDevice device;
+	WGPUQueue queue;
 
 	bool has_error;
 } WebGPUDevice;