diff --git a/Examples/Software/xmake.lua b/Examples/Software/xmake.lua index ef43b42..b59a362 100644 --- a/Examples/Software/xmake.lua +++ b/Examples/Software/xmake.lua @@ -1,6 +1,4 @@ -add_repositories("nazara-engine-repo https://github.com/NazaraEngine/xmake-repo") - -add_requires("nzsl >=2023.12.31", { configs = { shared = false, nzslc = true } }) +add_requires("nzsl", { configs = { shared = false, nzslc = true } }) if is_cross() then add_requires("nzsl~host", { kind = "binary", host = true }) diff --git a/Examples/Vulkan/xmake.lua b/Examples/Vulkan/xmake.lua index 1a24af4..dfae680 100644 --- a/Examples/Vulkan/xmake.lua +++ b/Examples/Vulkan/xmake.lua @@ -1,6 +1,4 @@ -add_repositories("nazara-engine-repo https://github.com/NazaraEngine/xmake-repo") - -add_requires("nzsl >=2023.12.31", { configs = { shared = false, nzslc = true } }) +add_requires("nzsl", { configs = { shared = false, nzslc = true } }) if is_cross() then add_requires("nzsl~host", { kind = "binary", host = true }) diff --git a/Sources/Backends/WebGPU/WebGPUComputePipeline.c b/Sources/Backends/WebGPU/WebGPUComputePipeline.c index 11df01e..5c8f845 100644 --- a/Sources/Backends/WebGPU/WebGPUComputePipeline.c +++ b/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", ©))) + while((token = PulseStrtokR(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); diff --git a/Sources/Backends/WebGPU/WebGPUDevice.c b/Sources/Backends/WebGPU/WebGPUDevice.c index f425351..aa01333 100644 --- a/Sources/Backends/WebGPU/WebGPUDevice.c +++ b/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; diff --git a/Sources/PulseInternal.c b/Sources/PulseInternal.c index 9dfb960..5a3b12e 100644 --- a/Sources/PulseInternal.c +++ b/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; - return token; + } + + *saveptr = str; + while(**saveptr) + { + p = delim; + while(*p) + { + if(**saveptr == *p++) + { + *((*saveptr)++) = '\0'; + return str; + } + } + ++(*saveptr); + } + return str; } static int PulseIsSpace(int x) diff --git a/Tests/Shaders/WebGPU/ReadWriteBindings.wgsl b/Tests/Shaders/WebGPU/ReadWriteBindings.wgsl index f132dc3..9e06fda 100644 --- a/Tests/Shaders/WebGPU/ReadWriteBindings.wgsl +++ b/Tests/Shaders/WebGPU/ReadWriteBindings.wgsl @@ -1,7 +1,7 @@ -@group(0) @binding(0) var read_ssbo: array; -@group(0) @binding(1) var read_texture: texture_storage_2d; -@group(1) @binding(0) var write_ssbo: array; -@group(1) @binding(1) var write_texture: texture_storage_2d; +@group(0) @binding(0) var read_texture: texture_storage_2d; +@group(0) @binding(1) var read_ssbo: array; +@group(1) @binding(0) var write_texture: texture_storage_2d; +@group(1) @binding(1) var write_ssbo: array; @compute @workgroup_size(16, 16, 1) fn main(@builtin(global_invocation_id) grid: vec3) diff --git a/Tests/xmake.lua b/Tests/xmake.lua index 6b10fa7..e9822b6 100644 --- a/Tests/xmake.lua +++ b/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 })