fixing WebGPU unit tests

Examples/Software/xmake.lua

@@ -1,6 +1,4 @@
-add_repositories("nazara-engine-repo https://github.com/NazaraEngine/xmake-repo")
+add_requires("nzsl", { configs = { shared = false, nzslc = true } })
-
-add_requires("nzsl >=2023.12.31", { configs = { shared = false, nzslc = true } })
 
 if is_cross() then
 	add_requires("nzsl~host", { kind = "binary", host = true })

Examples/Vulkan/xmake.lua

@@ -1,6 +1,4 @@
-add_repositories("nazara-engine-repo https://github.com/NazaraEngine/xmake-repo")
+add_requires("nzsl", { configs = { shared = false, nzslc = true } })
-
-add_requires("nzsl >=2023.12.31", { configs = { shared = false, nzslc = true } })
 
 if is_cross() then
 	add_requires("nzsl~host", { kind = "binary", host = true })

Sources/Backends/WebGPU/WebGPUComputePipeline.c

@@ -93,9 +93,10 @@ static void WebGPUReflectWGSLBindings(WebGPUBindGroupLayoutEntryInfo* infos, con
 
 	char* token;
 	char* copy = calloc(strlen(wgsl) + 1, 1);
+	char* copy_save = copy;
 	PulseStrlcpy(copy, wgsl, strlen(wgsl));
 
-	while((token = PulseStrtokR(copy, "\n", &copy)))
+	while((token = PulseStrtokR(copy, ";", &copy)))
 	{
 		found = false;
 		char* storage_pos;
@@ -151,7 +152,7 @@ static void WebGPUReflectWGSLBindings(WebGPUBindGroupLayoutEntryInfo* infos, con
 			infos->read_dimensions[count_read_images] = webgpu_dimension;
 			count_read_images++;
 		}
-		else
+		else if(webgpu_access != WGPUStorageTextureAccess_Undefined)
 		{
 			infos->write_access[count_write_images] = webgpu_access;
 			infos->write_formats[count_write_images] = webgpu_format;
@@ -160,7 +161,7 @@ static void WebGPUReflectWGSLBindings(WebGPUBindGroupLayoutEntryInfo* infos, con
 		}
 	}
 
-	free(copy);
+	free(copy_save);
 }
 
 static WGPUBindGroupLayout WebGPUCreateBindGroupLayout(PulseDevice device, const WebGPUBindGroupLayoutEntryInfo* infos,
@@ -189,6 +190,7 @@ static WGPUBindGroupLayout WebGPUCreateBindGroupLayout(PulseDevice device, const
 		{
 			entries[count].binding = count;
 			entries[count].visibility = WGPUShaderStage_Compute;
+			entries[count].nextInChain = PULSE_NULLPTR;
 
 			entries[count].buffer.nextInChain = PULSE_NULLPTR;
 			entries[count].buffer.hasDynamicOffset = false;
@@ -213,6 +215,7 @@ static WGPUBindGroupLayout WebGPUCreateBindGroupLayout(PulseDevice device, const
 		{
 			entries[count].binding = count;
 			entries[count].visibility = WGPUShaderStage_Compute;
+			entries[count].nextInChain = PULSE_NULLPTR;
 
 			entries[count].buffer.nextInChain = PULSE_NULLPTR;
 			entries[count].buffer.hasDynamicOffset = false;
@@ -239,6 +242,7 @@ static WGPUBindGroupLayout WebGPUCreateBindGroupLayout(PulseDevice device, const
 		{
 			entries[count].binding = count;
 			entries[count].visibility = WGPUShaderStage_Compute;
+			entries[count].nextInChain = PULSE_NULLPTR;
 
 			entries[count].buffer.nextInChain = PULSE_NULLPTR;
 			entries[count].buffer.hasDynamicOffset = false;
@@ -263,6 +267,7 @@ static WGPUBindGroupLayout WebGPUCreateBindGroupLayout(PulseDevice device, const
 		{
 			entries[count].binding = count;
 			entries[count].visibility = WGPUShaderStage_Compute;
+			entries[count].nextInChain = PULSE_NULLPTR;
 
 			entries[count].buffer.nextInChain = PULSE_NULLPTR;
 			entries[count].buffer.hasDynamicOffset = false;
@@ -289,6 +294,7 @@ static WGPUBindGroupLayout WebGPUCreateBindGroupLayout(PulseDevice device, const
 		{
 			entries[count].binding = count;
 			entries[count].visibility = WGPUShaderStage_Compute;
+			entries[count].nextInChain = PULSE_NULLPTR;
 
 			entries[count].buffer.nextInChain = PULSE_NULLPTR;
 			entries[count].buffer.hasDynamicOffset = false;
@@ -311,6 +317,7 @@ static WGPUBindGroupLayout WebGPUCreateBindGroupLayout(PulseDevice device, const
 	}
 
 	WGPUBindGroupLayoutDescriptor descriptor = { 0 };
+	descriptor.nextInChain = PULSE_NULLPTR;
 	descriptor.entryCount = count;
 	descriptor.entries = entries;
 	return wgpuDeviceCreateBindGroupLayout(webgpu_device->device, &descriptor);

Sources/Backends/WebGPU/WebGPUDevice.c

@@ -43,6 +43,9 @@
 		WGPULimits limits = { 0 };
 		wgpuAdapterGetLimits(adapter, &limits);
 
+		if(!wgpuAdapterHasFeature(adapter, WGPUNativeFeature_TextureAdapterSpecificFormatFeatures))
+			return 0;
+
 		if(infos.adapterType == WGPUAdapterType_DiscreteGPU)
 			score += 10000;
 
@@ -199,14 +202,25 @@ PulseDevice WebGPUCreateDevice(PulseBackend backend, PulseDevice* forbiden_devic
 	lost_callback.mode = WGPUCallbackMode_AllowSpontaneous;
 	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;
+
+	#ifndef PULSE_PLAT_WASM
+		WGPUFeatureName features[] = {
+			WGPUNativeFeature_TextureAdapterSpecificFormatFeatures,
+		};
+		descriptor.requiredFeatures = features;
+		descriptor.requiredFeatureCount = 1;
+	#endif
+
 	WGPURequestDeviceCallbackInfo device_callback = { 0 };
 	device_callback.callback = WebGPURequestDeviceCallback;
 	device_callback.mode = WGPUCallbackMode_AllowSpontaneous;

Sources/PulseInternal.c

@@ -113,14 +113,50 @@ size_t PulseStrlcpy(char* dst, const char* src, size_t maxlen)
 
 char* PulseStrtokR(char* str, const char* delim, char** saveptr)
 {
+	if(!delim || !saveptr || (!str && !*saveptr))
+		return PULSE_NULLPTR;
+
 	if(str != PULSE_NULLPTR)
 		*saveptr = str;
-	else if(*saveptr == PULSE_NULLPTR)
+	else
+	{
+		if(*saveptr == PULSE_NULLPTR)
+			return PULSE_NULLPTR;
+		str = *saveptr;
+	}
+
+	const char* p = delim;
+	while(*p && *str)
+	{
+		if(*str == *p)
+		{
+			++str;
+			p = delim;
+			continue;
+		}
+		++p;
+	}
+	if(!*str)
+	{
+		*saveptr = str;
 		return PULSE_NULLPTR;
-	char* token = strtok(*saveptr, delim);
+	}
-	if(token != PULSE_NULLPTR)
+
-		*saveptr = PULSE_NULLPTR;
+	*saveptr = str;
-	return token;
+	while(**saveptr)
+	{
+		p = delim;
+		while(*p)
+		{
+			if(**saveptr == *p++)
+			{
+				*((*saveptr)++) = '\0';
+				return str;
+			}
+		}
+		++(*saveptr);
+	}
+	return str;
 }
 
 static int PulseIsSpace(int x)

Tests/Shaders/WebGPU/ReadWriteBindings.wgsl

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

Tests/xmake.lua

@@ -8,7 +8,6 @@ local nzsl_included = false
 
 function nzsl(backend)
 	if not nzsl_included then
-		add_repositories("nazara-engine-repo https://github.com/NazaraEngine/xmake-repo")
 		add_requires("nzsl", { configs = { shared = false, nzslc = true } })
 		if is_cross() then
 			add_requires("nzsl~host", { kind = "binary", host = true })