diff --git a/Sources/Backends/WebGPU/WebGPUBuffer.c b/Sources/Backends/WebGPU/WebGPUBuffer.c index 05d70d4..8e867f1 100644 --- a/Sources/Backends/WebGPU/WebGPUBuffer.c +++ b/Sources/Backends/WebGPU/WebGPUBuffer.c @@ -211,7 +211,7 @@ bool WebGPUCopyBufferToImage(PulseCommandList cmd, const PulseBufferRegion* src, WGPUTexelCopyTextureInfo texture_copy_info = { 0 }; texture_copy_info.texture = webgpu_dst_image->texture; - texture_copy_info.mipLevel = 1; + texture_copy_info.mipLevel = 0; texture_copy_info.aspect = WGPUTextureAspect_All; texture_copy_info.origin.x = dst->x; texture_copy_info.origin.y = dst->y; diff --git a/Sources/Backends/WebGPU/WebGPUComputePipeline.c b/Sources/Backends/WebGPU/WebGPUComputePipeline.c index c0cf706..11df01e 100644 --- a/Sources/Backends/WebGPU/WebGPUComputePipeline.c +++ b/Sources/Backends/WebGPU/WebGPUComputePipeline.c @@ -2,6 +2,10 @@ // This file is part of "Pulse" // For conditions of distribution and use, see copyright notice in LICENSE +#include +#include +#include + #include #include "../../PulseInternal.h" #include "WebGPU.h" @@ -9,7 +13,157 @@ #include "webgpu.h" #include "WebGPUComputePipeline.h" -static WGPUBindGroupLayout WebGPUCreateBindGroupLayout(PulseDevice device, +typedef struct WebGPUBindGroupLayoutEntryInfo +{ + WGPUStorageTextureAccess read_access[PULSE_MAX_READ_TEXTURES_BOUND]; + WGPUTextureViewDimension read_dimensions[PULSE_MAX_READ_TEXTURES_BOUND]; + WGPUTextureFormat read_formats[PULSE_MAX_READ_TEXTURES_BOUND]; + + WGPUStorageTextureAccess write_access[PULSE_MAX_WRITE_TEXTURES_BOUND]; + WGPUTextureViewDimension write_dimensions[PULSE_MAX_WRITE_TEXTURES_BOUND]; + WGPUTextureFormat write_formats[PULSE_MAX_WRITE_TEXTURES_BOUND]; +} WebGPUBindGroupLayoutEntryInfo; + +static WGPUTextureFormat WebGPUMapFormat(const char* format) +{ + if(strcmp(format, "r32float") == 0) + return WGPUTextureFormat_R32Float; + if(strcmp(format, "rg32float") == 0) + return WGPUTextureFormat_RG32Float; + if(strcmp(format, "rgba32float") == 0) + return WGPUTextureFormat_RGBA32Float; + if(strcmp(format, "r32uint") == 0) + return WGPUTextureFormat_R32Uint; + if(strcmp(format, "r32sint") == 0) + return WGPUTextureFormat_R32Sint; + if(strcmp(format, "rg32uint") == 0) + return WGPUTextureFormat_RG32Uint; + if(strcmp(format, "rg32sint") == 0) + return WGPUTextureFormat_RG32Sint; + if(strcmp(format, "rgba32uint") == 0) + return WGPUTextureFormat_RGBA32Uint; + if(strcmp(format, "rgba32sint") == 0) + return WGPUTextureFormat_RGBA32Sint; + if(strcmp(format, "rgba8unorm") == 0) + return WGPUTextureFormat_RGBA8Unorm; + if(strcmp(format, "rgba8snorm") == 0) + return WGPUTextureFormat_RGBA8Snorm; + if(strcmp(format, "rgba8uint") == 0) + return WGPUTextureFormat_RGBA8Uint; + if(strcmp(format, "rgba8sint") == 0) + return WGPUTextureFormat_RGBA8Sint; + if(strcmp(format, "rg8unorm") == 0) + return WGPUTextureFormat_RG8Unorm; + if(strcmp(format, "rg8snorm") == 0) + return WGPUTextureFormat_RG8Snorm; + if(strcmp(format, "rg8uint") == 0) + return WGPUTextureFormat_RG8Uint; + if(strcmp(format, "rg8sint") == 0) + return WGPUTextureFormat_RG8Sint; + if(strcmp(format, "r8unorm") == 0) + return WGPUTextureFormat_R8Unorm; + if(strcmp(format, "r8snorm") == 0) + return WGPUTextureFormat_R8Snorm; + if(strcmp(format, "r8uint") == 0) + return WGPUTextureFormat_R8Uint; + if(strcmp(format, "r8sint") == 0) + return WGPUTextureFormat_R8Sint; + if(strcmp(format, "bgra8unorm") == 0) + return WGPUTextureFormat_BGRA8Unorm; + return WGPUTextureFormat_Undefined; +} + +static WGPUStorageTextureAccess WebGPUMapAccess(const char* access) +{ + if(strcmp(access, "read") == 0) + return WGPUStorageTextureAccess_ReadOnly; + if(strcmp(access, "write") == 0) + return WGPUStorageTextureAccess_WriteOnly; + if(strcmp(access, "read_write") == 0) + return WGPUStorageTextureAccess_ReadWrite; + return WGPUStorageTextureAccess_Undefined; +} + +static void WebGPUReflectWGSLBindings(WebGPUBindGroupLayoutEntryInfo* infos, const char* wgsl) +{ + bool found; + + uint32_t count_read_images = 0; + uint32_t count_write_images = 0; + + char* token; + char* copy = calloc(strlen(wgsl) + 1, 1); + PulseStrlcpy(copy, wgsl, strlen(wgsl)); + + while((token = PulseStrtokR(copy, "\n", ©))) + { + found = false; + char* storage_pos; + + WGPUStorageTextureAccess webgpu_access = WGPUStorageTextureAccess_Undefined; + WGPUTextureFormat webgpu_format = WGPUTextureFormat_Undefined; + WGPUTextureViewDimension webgpu_dimension = WGPUTextureViewDimension_Undefined; + + if((storage_pos = strstr(token, "texture_storage_2d<"))) + { + char format[32]; + char access[32]; + if(sscanf(storage_pos, "texture_storage_2d<%[^,],%[^>]>", format, access) == 2) + { + PulseTrimString(format); + PulseTrimString(access); + webgpu_access = WebGPUMapAccess(access); + webgpu_format = WebGPUMapFormat(format); + webgpu_dimension = WGPUTextureViewDimension_2D; + } + } + else if((storage_pos = strstr(token, "texture_storage_2d_array<"))) + { + char format[32]; + char access[32]; + if(sscanf(storage_pos, "texture_storage_2d_array<%[^,],%[^>]>", format, access) == 2) + { + PulseTrimString(format); + PulseTrimString(access); + webgpu_access = WebGPUMapAccess(access); + webgpu_format = WebGPUMapFormat(format); + webgpu_dimension = WGPUTextureViewDimension_2DArray; + } + } + else if((storage_pos = strstr(token, "texture_storage_3d<"))) + { + char format[32]; + char access[32]; + if(sscanf(storage_pos, "texture_storage_3d<%[^,],%[^>]>", format, access) == 2) + { + PulseTrimString(format); + PulseTrimString(access); + webgpu_access = WebGPUMapAccess(access); + webgpu_format = WebGPUMapFormat(format); + webgpu_dimension = WGPUTextureViewDimension_3D; + } + } + + if(webgpu_access == WGPUStorageTextureAccess_ReadOnly) + { + infos->read_access[count_read_images] = webgpu_access; + infos->read_formats[count_read_images] = webgpu_format; + infos->read_dimensions[count_read_images] = webgpu_dimension; + count_read_images++; + } + else + { + infos->write_access[count_write_images] = webgpu_access; + infos->write_formats[count_write_images] = webgpu_format; + infos->write_dimensions[count_write_images] = webgpu_dimension; + count_write_images++; + } + } + + free(copy); +} + +static WGPUBindGroupLayout WebGPUCreateBindGroupLayout(PulseDevice device, const WebGPUBindGroupLayoutEntryInfo* infos, uint32_t read_storage_images_count, uint32_t read_storage_buffers_count, uint32_t write_storage_images_count, @@ -33,126 +187,126 @@ static WGPUBindGroupLayout WebGPUCreateBindGroupLayout(PulseDevice device, { for(uint32_t i = 0; i < read_storage_images_count; i++, count++) { - entries[i].binding = i; - entries[i].visibility = WGPUShaderStage_Compute; + entries[count].binding = count; + entries[count].visibility = WGPUShaderStage_Compute; - entries[i].buffer.nextInChain = PULSE_NULLPTR; - entries[i].buffer.hasDynamicOffset = false; - entries[i].buffer.type = WGPUBufferBindingType_BindingNotUsed; - entries[i].buffer.minBindingSize = 0; + entries[count].buffer.nextInChain = PULSE_NULLPTR; + entries[count].buffer.hasDynamicOffset = false; + entries[count].buffer.type = WGPUBufferBindingType_BindingNotUsed; + entries[count].buffer.minBindingSize = 0; - entries[i].sampler.nextInChain = PULSE_NULLPTR; - entries[i].sampler.type = WGPUSamplerBindingType_BindingNotUsed; + entries[count].sampler.nextInChain = PULSE_NULLPTR; + entries[count].sampler.type = WGPUSamplerBindingType_BindingNotUsed; - entries[i].storageTexture.nextInChain = PULSE_NULLPTR; - entries[i].storageTexture.access = WGPUStorageTextureAccess_Undefined; - entries[i].storageTexture.format = WGPUTextureFormat_Undefined; - entries[i].storageTexture.viewDimension = WGPUTextureViewDimension_Undefined; + entries[count].storageTexture.nextInChain = PULSE_NULLPTR; + entries[count].storageTexture.access = infos->read_access[i]; + entries[count].storageTexture.format = infos->read_formats[i]; + entries[count].storageTexture.viewDimension = infos->read_dimensions[i]; - entries[i].texture.nextInChain = PULSE_NULLPTR; - entries[i].texture.multisampled = false; - entries[i].texture.sampleType = WGPUTextureSampleType_BindingNotUsed; - entries[i].texture.viewDimension = WGPUTextureViewDimension_Undefined; + entries[count].texture.nextInChain = PULSE_NULLPTR; + entries[count].texture.multisampled = false; + entries[count].texture.sampleType = WGPUTextureSampleType_BindingNotUsed; + entries[count].texture.viewDimension = WGPUTextureViewDimension_Undefined; } - for(uint32_t i = read_storage_images_count; i < read_storage_images_count + read_storage_buffers_count; i++, count++) + for(uint32_t i = 0; i < read_storage_buffers_count; i++, count++) { - entries[i].binding = i; - entries[i].visibility = WGPUShaderStage_Compute; + entries[count].binding = count; + entries[count].visibility = WGPUShaderStage_Compute; - entries[i].buffer.nextInChain = PULSE_NULLPTR; - entries[i].buffer.hasDynamicOffset = false; - entries[i].buffer.type = WGPUBufferBindingType_ReadOnlyStorage; - entries[i].buffer.minBindingSize = 0; + entries[count].buffer.nextInChain = PULSE_NULLPTR; + entries[count].buffer.hasDynamicOffset = false; + entries[count].buffer.type = WGPUBufferBindingType_ReadOnlyStorage; + entries[count].buffer.minBindingSize = 0; - entries[i].sampler.nextInChain = PULSE_NULLPTR; - entries[i].sampler.type = WGPUSamplerBindingType_BindingNotUsed; + entries[count].sampler.nextInChain = PULSE_NULLPTR; + entries[count].sampler.type = WGPUSamplerBindingType_BindingNotUsed; - entries[i].storageTexture.nextInChain = PULSE_NULLPTR; - entries[i].storageTexture.access = WGPUStorageTextureAccess_BindingNotUsed; - entries[i].storageTexture.format = WGPUTextureFormat_Undefined; - entries[i].storageTexture.viewDimension = WGPUTextureViewDimension_Undefined; + entries[count].storageTexture.nextInChain = PULSE_NULLPTR; + entries[count].storageTexture.access = WGPUStorageTextureAccess_BindingNotUsed; + entries[count].storageTexture.format = WGPUTextureFormat_Undefined; + entries[count].storageTexture.viewDimension = WGPUTextureViewDimension_Undefined; - entries[i].texture.nextInChain = PULSE_NULLPTR; - entries[i].texture.multisampled = false; - entries[i].texture.sampleType = WGPUTextureSampleType_BindingNotUsed; - entries[i].texture.viewDimension = WGPUTextureViewDimension_Undefined; + entries[count].texture.nextInChain = PULSE_NULLPTR; + entries[count].texture.multisampled = false; + entries[count].texture.sampleType = WGPUTextureSampleType_BindingNotUsed; + entries[count].texture.viewDimension = WGPUTextureViewDimension_Undefined; } } else if(category == 2) { for(uint32_t i = 0; i < write_storage_images_count; i++, count++) { - entries[i].binding = i; - entries[i].visibility = WGPUShaderStage_Compute; + entries[count].binding = count; + entries[count].visibility = WGPUShaderStage_Compute; - entries[i].buffer.nextInChain = PULSE_NULLPTR; - entries[i].buffer.hasDynamicOffset = false; - entries[i].buffer.type = WGPUBufferBindingType_BindingNotUsed; - entries[i].buffer.minBindingSize = 0; + entries[count].buffer.nextInChain = PULSE_NULLPTR; + entries[count].buffer.hasDynamicOffset = false; + entries[count].buffer.type = WGPUBufferBindingType_BindingNotUsed; + entries[count].buffer.minBindingSize = 0; - entries[i].sampler.nextInChain = PULSE_NULLPTR; - entries[i].sampler.type = WGPUSamplerBindingType_BindingNotUsed; + entries[count].sampler.nextInChain = PULSE_NULLPTR; + entries[count].sampler.type = WGPUSamplerBindingType_BindingNotUsed; - entries[i].storageTexture.nextInChain = PULSE_NULLPTR; - entries[i].storageTexture.access = WGPUStorageTextureAccess_BindingNotUsed; - entries[i].storageTexture.format = WGPUTextureFormat_Undefined; - entries[i].storageTexture.viewDimension = WGPUTextureViewDimension_Undefined; + entries[count].storageTexture.nextInChain = PULSE_NULLPTR; + entries[count].storageTexture.access = infos->write_access[i]; + entries[count].storageTexture.format = infos->write_formats[i]; + entries[count].storageTexture.viewDimension = infos->write_dimensions[i]; - entries[i].texture.nextInChain = PULSE_NULLPTR; - entries[i].texture.multisampled = false; - entries[i].texture.sampleType = WGPUTextureSampleType_BindingNotUsed; - entries[i].texture.viewDimension = WGPUTextureViewDimension_Undefined; + entries[count].texture.nextInChain = PULSE_NULLPTR; + entries[count].texture.multisampled = false; + entries[count].texture.sampleType = WGPUTextureSampleType_BindingNotUsed; + entries[count].texture.viewDimension = WGPUTextureViewDimension_Undefined; } - for(uint32_t i = write_storage_images_count; i < write_storage_images_count + write_storage_buffers_count; i++, count++) + for(uint32_t i = 0; i < write_storage_buffers_count; i++, count++) { - entries[i].binding = i; - entries[i].visibility = WGPUShaderStage_Compute; + entries[count].binding = count; + entries[count].visibility = WGPUShaderStage_Compute; - entries[i].buffer.nextInChain = PULSE_NULLPTR; - entries[i].buffer.hasDynamicOffset = false; - entries[i].buffer.type = WGPUBufferBindingType_Storage; - entries[i].buffer.minBindingSize = 0; + entries[count].buffer.nextInChain = PULSE_NULLPTR; + entries[count].buffer.hasDynamicOffset = false; + entries[count].buffer.type = WGPUBufferBindingType_Storage; + entries[count].buffer.minBindingSize = 0; - entries[i].sampler.nextInChain = PULSE_NULLPTR; - entries[i].sampler.type = WGPUSamplerBindingType_BindingNotUsed; + entries[count].sampler.nextInChain = PULSE_NULLPTR; + entries[count].sampler.type = WGPUSamplerBindingType_BindingNotUsed; - entries[i].storageTexture.nextInChain = PULSE_NULLPTR; - entries[i].storageTexture.access = WGPUStorageTextureAccess_BindingNotUsed; - entries[i].storageTexture.format = WGPUTextureFormat_Undefined; - entries[i].storageTexture.viewDimension = WGPUTextureViewDimension_Undefined; + entries[count].storageTexture.nextInChain = PULSE_NULLPTR; + entries[count].storageTexture.access = WGPUStorageTextureAccess_BindingNotUsed; + entries[count].storageTexture.format = WGPUTextureFormat_Undefined; + entries[count].storageTexture.viewDimension = WGPUTextureViewDimension_Undefined; - entries[i].texture.nextInChain = PULSE_NULLPTR; - entries[i].texture.multisampled = false; - entries[i].texture.sampleType = WGPUTextureSampleType_BindingNotUsed; - entries[i].texture.viewDimension = WGPUTextureViewDimension_Undefined; + entries[count].texture.nextInChain = PULSE_NULLPTR; + entries[count].texture.multisampled = false; + entries[count].texture.sampleType = WGPUTextureSampleType_BindingNotUsed; + entries[count].texture.viewDimension = WGPUTextureViewDimension_Undefined; } } else if(category == 3) { for(uint32_t i = 0; i < uniform_buffers_count; i++, count++) { - entries[i].binding = i; - entries[i].visibility = WGPUShaderStage_Compute; + entries[count].binding = count; + entries[count].visibility = WGPUShaderStage_Compute; - entries[i].buffer.nextInChain = PULSE_NULLPTR; - entries[i].buffer.hasDynamicOffset = false; - entries[i].buffer.type = WGPUBufferBindingType_Uniform; - entries[i].buffer.minBindingSize = 0; + entries[count].buffer.nextInChain = PULSE_NULLPTR; + entries[count].buffer.hasDynamicOffset = false; + entries[count].buffer.type = WGPUBufferBindingType_Uniform; + entries[count].buffer.minBindingSize = 0; - entries[i].sampler.nextInChain = PULSE_NULLPTR; - entries[i].sampler.type = WGPUSamplerBindingType_BindingNotUsed; + entries[count].sampler.nextInChain = PULSE_NULLPTR; + entries[count].sampler.type = WGPUSamplerBindingType_BindingNotUsed; - entries[i].storageTexture.nextInChain = PULSE_NULLPTR; - entries[i].storageTexture.access = WGPUStorageTextureAccess_BindingNotUsed; - entries[i].storageTexture.format = WGPUTextureFormat_Undefined; - entries[i].storageTexture.viewDimension = WGPUTextureViewDimension_Undefined; + entries[count].storageTexture.nextInChain = PULSE_NULLPTR; + entries[count].storageTexture.access = WGPUStorageTextureAccess_BindingNotUsed; + entries[count].storageTexture.format = WGPUTextureFormat_Undefined; + entries[count].storageTexture.viewDimension = WGPUTextureViewDimension_Undefined; - entries[i].texture.nextInChain = PULSE_NULLPTR; - entries[i].texture.multisampled = false; - entries[i].texture.sampleType = WGPUTextureSampleType_BindingNotUsed; - entries[i].texture.viewDimension = WGPUTextureViewDimension_Undefined; + entries[count].texture.nextInChain = PULSE_NULLPTR; + entries[count].texture.multisampled = false; + entries[count].texture.sampleType = WGPUTextureSampleType_BindingNotUsed; + entries[count].texture.viewDimension = WGPUTextureViewDimension_Undefined; } } @@ -194,9 +348,13 @@ PulseComputePipeline WebGPUCreateComputePipeline(PulseDevice device, const Pulse shader_descriptor.nextInChain = (const WGPUChainedStruct*)&source; webgpu_pipeline->shader = wgpuDeviceCreateShaderModule(webgpu_device->device, &shader_descriptor); - webgpu_pipeline->readonly_group = WebGPUCreateBindGroupLayout(device, info->num_readonly_storage_images, info->num_readonly_storage_buffers, 0, 0, 0); - webgpu_pipeline->readwrite_group = WebGPUCreateBindGroupLayout(device, 0, 0, info->num_readwrite_storage_images, info->num_readwrite_storage_buffers, 0); - webgpu_pipeline->uniform_group = WebGPUCreateBindGroupLayout(device, 0, 0, 0, 0, info->num_uniform_buffers); + WebGPUBindGroupLayoutEntryInfo bindgroup_infos = { 0 }; + + WebGPUReflectWGSLBindings(&bindgroup_infos, (const char*)info->code); + + webgpu_pipeline->readonly_group = WebGPUCreateBindGroupLayout(device, &bindgroup_infos, info->num_readonly_storage_images, info->num_readonly_storage_buffers, 0, 0, 0); + webgpu_pipeline->readwrite_group = WebGPUCreateBindGroupLayout(device, &bindgroup_infos, 0, 0, info->num_readwrite_storage_images, info->num_readwrite_storage_buffers, 0); + webgpu_pipeline->uniform_group = WebGPUCreateBindGroupLayout(device, &bindgroup_infos, 0, 0, 0, 0, info->num_uniform_buffers); WGPUBindGroupLayout bind_group_layouts[3] = { webgpu_pipeline->readonly_group, diff --git a/Sources/Backends/WebGPU/WebGPUImage.c b/Sources/Backends/WebGPU/WebGPUImage.c index 8a7663d..3274284 100644 --- a/Sources/Backends/WebGPU/WebGPUImage.c +++ b/Sources/Backends/WebGPU/WebGPUImage.c @@ -121,9 +121,9 @@ PulseImage WebGPUCreateImage(PulseDevice device, const PulseImageCreateInfo* cre descriptor.usage = 0; if(create_infos->usage & PULSE_IMAGE_USAGE_STORAGE_READ) - descriptor.usage |= WGPUTextureUsage_StorageBinding | WGPUTextureUsage_CopyDst; - if(create_infos->usage & PULSE_IMAGE_USAGE_STORAGE_WRITE) descriptor.usage |= WGPUTextureUsage_StorageBinding | WGPUTextureUsage_CopySrc; + if(create_infos->usage & PULSE_IMAGE_USAGE_STORAGE_WRITE) + descriptor.usage |= WGPUTextureUsage_StorageBinding | WGPUTextureUsage_CopyDst; if(create_infos->usage & PULSE_IMAGE_USAGE_STORAGE_SIMULTANEOUS_READWRITE) descriptor.usage |= WGPUTextureUsage_StorageBinding | WGPUTextureUsage_CopySrc | WGPUTextureUsage_CopyDst; @@ -224,7 +224,7 @@ bool WebGPUCopyImageToBuffer(PulseCommandList cmd, const PulseImageRegion* src, WGPUTexelCopyTextureInfo texture_copy_info = { 0 }; texture_copy_info.texture = webgpu_src_image->texture; - texture_copy_info.mipLevel = 1; + texture_copy_info.mipLevel = 0; texture_copy_info.aspect = WGPUTextureAspect_All; texture_copy_info.origin.x = src->x; texture_copy_info.origin.y = src->y; diff --git a/Sources/PulseInternal.c b/Sources/PulseInternal.c index a0227b8..8686088 100644 --- a/Sources/PulseInternal.c +++ b/Sources/PulseInternal.c @@ -98,3 +98,45 @@ uint32_t PulseHashCombine(uint32_t lhs, uint32_t rhs) lhs ^= rhs + 0x9e3779b9 + (lhs << 6) + (lhs >> 2); return lhs; } + +size_t PulseStrlcpy(char* dst, const char* src, size_t maxlen) +{ + size_t srclen = strlen(src); + if(maxlen > 0) + { + size_t len = (srclen < maxlen - 1 ? srclen : maxlen - 1); + memcpy(dst, src, len); + dst[len] = '\0'; + } + return srclen; +} + +char* PulseStrtokR(char* str, const char* delim, char** saveptr) +{ + if(str != PULSE_NULLPTR) + *saveptr = str; + else if(*saveptr == PULSE_NULLPTR) + return PULSE_NULLPTR; + char* token = strtok(*saveptr, delim); + if(token != PULSE_NULLPTR) + *saveptr = PULSE_NULLPTR; + return token; +} + +static int isspace(int x) +{ + return ((x) == ' ') || ((x) == '\t') || ((x) == '\r') || ((x) == '\n') || ((x) == '\f') || ((x) == '\v'); +} + +void PulseTrimString(char* str) +{ + char* start = str; + while(*start && isspace((unsigned char)*start)) + start++; + char* end = start + strlen(start) - 1; + while(end > start && isspace((unsigned char)*end)) + end--; + *(end + 1) = '\0'; + if(start != str) + memmove(str, start, end - start + 2); +} diff --git a/Sources/PulseInternal.h b/Sources/PulseInternal.h index 9ca4fd6..a3b197f 100644 --- a/Sources/PulseInternal.h +++ b/Sources/PulseInternal.h @@ -192,6 +192,10 @@ void PulseSetInternalError(PulseErrorType error); uint32_t PulseHashString(const char* str); uint32_t PulseHashCombine(uint32_t lhs, uint32_t rhs); +size_t PulseStrlcpy(char* dst, const char* src, size_t maxlen); +char* PulseStrtokR(char* str, const char* delim, char** saveptr); +void PulseTrimString(char* str); + void PulseLogBackend(PulseBackend backend, PulseDebugMessageSeverity type, const char* message, const char* file, const char* function, int line, ...); #define PulseLogError(backend, msg) PulseLogBackend(backend, PULSE_DEBUG_MESSAGE_SEVERITY_ERROR, msg, __FILE__, __FUNCTION__, __LINE__) diff --git a/Tests/Buffer.c b/Tests/Buffer.c index a1012aa..5905c75 100644 --- a/Tests/Buffer.c +++ b/Tests/Buffer.c @@ -239,11 +239,11 @@ void TestBufferCopyImage() PulseImageRegion dst_region = { 0 }; dst_region.image = image; dst_region.width = 8; - dst_region.height = 1; - dst_region.depth = 1; - dst_region.x = 1; + dst_region.height = 0; + dst_region.depth = 0; + dst_region.x = 0; dst_region.y = 1; - dst_region.z = 1; + dst_region.z = 0; dst_region.layer = 1; TEST_ASSERT_TRUE_MESSAGE(PulseCopyBufferToImage(cmd, &src_region, &dst_region), PulseVerbaliseErrorType(PulseGetLastErrorType())); diff --git a/Tests/Shaders/WebGPU/BufferCopy.wgsl b/Tests/Shaders/WebGPU/BufferCopy.wgsl index a273145..e002736 100644 --- a/Tests/Shaders/WebGPU/BufferCopy.wgsl +++ b/Tests/Shaders/WebGPU/BufferCopy.wgsl @@ -1,4 +1,4 @@ -@group(0) @binding(0) var read_ssbo: array; +@group(0) @binding(0) var read_ssbo: array; @group(1) @binding(0) var write_ssbo: array; @compute @workgroup_size(16, 16, 1) diff --git a/Tests/Shaders/WebGPU/ReadWriteBindings.wgsl b/Tests/Shaders/WebGPU/ReadWriteBindings.wgsl index 44a312b..f132dc3 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(1) @binding(1) var write_texture: texture_storage_2d; @compute @workgroup_size(16, 16, 1) fn main(@builtin(global_invocation_id) grid: vec3) diff --git a/Tests/Shaders/WebGPU/WriteOnlyBindings.wgsl b/Tests/Shaders/WebGPU/WriteOnlyBindings.wgsl index c20f6c3..96decc5 100644 --- a/Tests/Shaders/WebGPU/WriteOnlyBindings.wgsl +++ b/Tests/Shaders/WebGPU/WriteOnlyBindings.wgsl @@ -1,5 +1,5 @@ @group(1) @binding(0) var write_ssbo: array; -@group(1) @binding(1) var write_texture: texture_storage_2d; +@group(1) @binding(1) var write_texture: texture_storage_2d; @compute @workgroup_size(16, 16, 1) fn main(@builtin(global_invocation_id) grid: vec3)