kbz_8/Pulse
1
0
Fork 0
mirror of https://github.com/Kbz-8/Pulse.git synced 2026-01-11 15:33:34 +00:00
Code Issues Packages Projects Releases Wiki Activity

adding bindsgroup reflection to WebGPU backend

Browse Source
This commit is contained in:
kbz_8
2025-04-09 12:03:13 +02:00
parent 5470779d4c
commit 63b336783b
9 changed files with 301 additions and 97 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
Sources/Backends/WebGPU/WebGPUBuffer.c
View File

@@ -211,7 +211,7 @@ bool WebGPUCopyBufferToImage(PulseCommandList cmd, const PulseBufferRegion* src,
WGPUTexelCopyTextureInfo texture_copy_info = { 0 }; WGPUTexelCopyTextureInfo texture_copy_info = { 0 };
texture_copy_info.texture = webgpu_dst_image->texture; 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.aspect = WGPUTextureAspect_All;
texture_copy_info.origin.x = dst->x; texture_copy_info.origin.x = dst->x;
texture_copy_info.origin.y = dst->y; texture_copy_info.origin.y = dst->y;

330
Sources/Backends/WebGPU/WebGPUComputePipeline.c
View File

@@ -2,6 +2,10 @@
// This file is part of "Pulse" // This file is part of "Pulse"
// For conditions of distribution and use, see copyright notice in LICENSE // For conditions of distribution and use, see copyright notice in LICENSE
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <Pulse.h> #include <Pulse.h>
#include "../../PulseInternal.h" #include "../../PulseInternal.h"
#include "WebGPU.h" #include "WebGPU.h"
@@ -9,7 +13,157 @@
#include "webgpu.h" #include "webgpu.h"
#include "WebGPUComputePipeline.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", &copy)))
{
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_images_count,
uint32_t read_storage_buffers_count, uint32_t read_storage_buffers_count,
uint32_t write_storage_images_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++) for(uint32_t i = 0; i < read_storage_images_count; i++, count++)
{ {
entries[i].binding = i; entries[count].binding = count;
entries[i].visibility = WGPUShaderStage_Compute; entries[count].visibility = WGPUShaderStage_Compute;
entries[i].buffer.nextInChain = PULSE_NULLPTR; entries[count].buffer.nextInChain = PULSE_NULLPTR;
entries[i].buffer.hasDynamicOffset = false; entries[count].buffer.hasDynamicOffset = false;
entries[i].buffer.type = WGPUBufferBindingType_BindingNotUsed; entries[count].buffer.type = WGPUBufferBindingType_BindingNotUsed;
entries[i].buffer.minBindingSize = 0; entries[count].buffer.minBindingSize = 0;
entries[i].sampler.nextInChain = PULSE_NULLPTR; entries[count].sampler.nextInChain = PULSE_NULLPTR;
entries[i].sampler.type = WGPUSamplerBindingType_BindingNotUsed; entries[count].sampler.type = WGPUSamplerBindingType_BindingNotUsed;
entries[i].storageTexture.nextInChain = PULSE_NULLPTR; entries[count].storageTexture.nextInChain = PULSE_NULLPTR;
entries[i].storageTexture.access = WGPUStorageTextureAccess_Undefined; entries[count].storageTexture.access = infos->read_access[i];
entries[i].storageTexture.format = WGPUTextureFormat_Undefined; entries[count].storageTexture.format = infos->read_formats[i];
entries[i].storageTexture.viewDimension = WGPUTextureViewDimension_Undefined; entries[count].storageTexture.viewDimension = infos->read_dimensions[i];
entries[i].texture.nextInChain = PULSE_NULLPTR; entries[count].texture.nextInChain = PULSE_NULLPTR;
entries[i].texture.multisampled = false; entries[count].texture.multisampled = false;
entries[i].texture.sampleType = WGPUTextureSampleType_BindingNotUsed; entries[count].texture.sampleType = WGPUTextureSampleType_BindingNotUsed;
entries[i].texture.viewDimension = WGPUTextureViewDimension_Undefined; 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[count].binding = count;
entries[i].visibility = WGPUShaderStage_Compute; entries[count].visibility = WGPUShaderStage_Compute;
entries[i].buffer.nextInChain = PULSE_NULLPTR; entries[count].buffer.nextInChain = PULSE_NULLPTR;
entries[i].buffer.hasDynamicOffset = false; entries[count].buffer.hasDynamicOffset = false;
entries[i].buffer.type = WGPUBufferBindingType_ReadOnlyStorage; entries[count].buffer.type = WGPUBufferBindingType_ReadOnlyStorage;
entries[i].buffer.minBindingSize = 0; entries[count].buffer.minBindingSize = 0;
entries[i].sampler.nextInChain = PULSE_NULLPTR; entries[count].sampler.nextInChain = PULSE_NULLPTR;
entries[i].sampler.type = WGPUSamplerBindingType_BindingNotUsed; entries[count].sampler.type = WGPUSamplerBindingType_BindingNotUsed;
entries[i].storageTexture.nextInChain = PULSE_NULLPTR; entries[count].storageTexture.nextInChain = PULSE_NULLPTR;
entries[i].storageTexture.access = WGPUStorageTextureAccess_BindingNotUsed; entries[count].storageTexture.access = WGPUStorageTextureAccess_BindingNotUsed;
entries[i].storageTexture.format = WGPUTextureFormat_Undefined; entries[count].storageTexture.format = WGPUTextureFormat_Undefined;
entries[i].storageTexture.viewDimension = WGPUTextureViewDimension_Undefined; entries[count].storageTexture.viewDimension = WGPUTextureViewDimension_Undefined;
entries[i].texture.nextInChain = PULSE_NULLPTR; entries[count].texture.nextInChain = PULSE_NULLPTR;
entries[i].texture.multisampled = false; entries[count].texture.multisampled = false;
entries[i].texture.sampleType = WGPUTextureSampleType_BindingNotUsed; entries[count].texture.sampleType = WGPUTextureSampleType_BindingNotUsed;
entries[i].texture.viewDimension = WGPUTextureViewDimension_Undefined; entries[count].texture.viewDimension = WGPUTextureViewDimension_Undefined;
} }
} }
else if(category == 2) else if(category == 2)
{ {
for(uint32_t i = 0; i < write_storage_images_count; i++, count++) for(uint32_t i = 0; i < write_storage_images_count; i++, count++)
{ {
entries[i].binding = i; entries[count].binding = count;
entries[i].visibility = WGPUShaderStage_Compute; entries[count].visibility = WGPUShaderStage_Compute;
entries[i].buffer.nextInChain = PULSE_NULLPTR; entries[count].buffer.nextInChain = PULSE_NULLPTR;
entries[i].buffer.hasDynamicOffset = false; entries[count].buffer.hasDynamicOffset = false;
entries[i].buffer.type = WGPUBufferBindingType_BindingNotUsed; entries[count].buffer.type = WGPUBufferBindingType_BindingNotUsed;
entries[i].buffer.minBindingSize = 0; entries[count].buffer.minBindingSize = 0;
entries[i].sampler.nextInChain = PULSE_NULLPTR; entries[count].sampler.nextInChain = PULSE_NULLPTR;
entries[i].sampler.type = WGPUSamplerBindingType_BindingNotUsed; entries[count].sampler.type = WGPUSamplerBindingType_BindingNotUsed;
entries[i].storageTexture.nextInChain = PULSE_NULLPTR; entries[count].storageTexture.nextInChain = PULSE_NULLPTR;
entries[i].storageTexture.access = WGPUStorageTextureAccess_BindingNotUsed; entries[count].storageTexture.access = infos->write_access[i];
entries[i].storageTexture.format = WGPUTextureFormat_Undefined; entries[count].storageTexture.format = infos->write_formats[i];
entries[i].storageTexture.viewDimension = WGPUTextureViewDimension_Undefined; entries[count].storageTexture.viewDimension = infos->write_dimensions[i];
entries[i].texture.nextInChain = PULSE_NULLPTR; entries[count].texture.nextInChain = PULSE_NULLPTR;
entries[i].texture.multisampled = false; entries[count].texture.multisampled = false;
entries[i].texture.sampleType = WGPUTextureSampleType_BindingNotUsed; entries[count].texture.sampleType = WGPUTextureSampleType_BindingNotUsed;
entries[i].texture.viewDimension = WGPUTextureViewDimension_Undefined; 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[count].binding = count;
entries[i].visibility = WGPUShaderStage_Compute; entries[count].visibility = WGPUShaderStage_Compute;
entries[i].buffer.nextInChain = PULSE_NULLPTR; entries[count].buffer.nextInChain = PULSE_NULLPTR;
entries[i].buffer.hasDynamicOffset = false; entries[count].buffer.hasDynamicOffset = false;
entries[i].buffer.type = WGPUBufferBindingType_Storage; entries[count].buffer.type = WGPUBufferBindingType_Storage;
entries[i].buffer.minBindingSize = 0; entries[count].buffer.minBindingSize = 0;
entries[i].sampler.nextInChain = PULSE_NULLPTR; entries[count].sampler.nextInChain = PULSE_NULLPTR;
entries[i].sampler.type = WGPUSamplerBindingType_BindingNotUsed; entries[count].sampler.type = WGPUSamplerBindingType_BindingNotUsed;
entries[i].storageTexture.nextInChain = PULSE_NULLPTR; entries[count].storageTexture.nextInChain = PULSE_NULLPTR;
entries[i].storageTexture.access = WGPUStorageTextureAccess_BindingNotUsed; entries[count].storageTexture.access = WGPUStorageTextureAccess_BindingNotUsed;
entries[i].storageTexture.format = WGPUTextureFormat_Undefined; entries[count].storageTexture.format = WGPUTextureFormat_Undefined;
entries[i].storageTexture.viewDimension = WGPUTextureViewDimension_Undefined; entries[count].storageTexture.viewDimension = WGPUTextureViewDimension_Undefined;
entries[i].texture.nextInChain = PULSE_NULLPTR; entries[count].texture.nextInChain = PULSE_NULLPTR;
entries[i].texture.multisampled = false; entries[count].texture.multisampled = false;
entries[i].texture.sampleType = WGPUTextureSampleType_BindingNotUsed; entries[count].texture.sampleType = WGPUTextureSampleType_BindingNotUsed;
entries[i].texture.viewDimension = WGPUTextureViewDimension_Undefined; entries[count].texture.viewDimension = WGPUTextureViewDimension_Undefined;
} }
} }
else if(category == 3) else if(category == 3)
{ {
for(uint32_t i = 0; i < uniform_buffers_count; i++, count++) for(uint32_t i = 0; i < uniform_buffers_count; i++, count++)
{ {
entries[i].binding = i; entries[count].binding = count;
entries[i].visibility = WGPUShaderStage_Compute; entries[count].visibility = WGPUShaderStage_Compute;
entries[i].buffer.nextInChain = PULSE_NULLPTR; entries[count].buffer.nextInChain = PULSE_NULLPTR;
entries[i].buffer.hasDynamicOffset = false; entries[count].buffer.hasDynamicOffset = false;
entries[i].buffer.type = WGPUBufferBindingType_Uniform; entries[count].buffer.type = WGPUBufferBindingType_Uniform;
entries[i].buffer.minBindingSize = 0; entries[count].buffer.minBindingSize = 0;
entries[i].sampler.nextInChain = PULSE_NULLPTR; entries[count].sampler.nextInChain = PULSE_NULLPTR;
entries[i].sampler.type = WGPUSamplerBindingType_BindingNotUsed; entries[count].sampler.type = WGPUSamplerBindingType_BindingNotUsed;
entries[i].storageTexture.nextInChain = PULSE_NULLPTR; entries[count].storageTexture.nextInChain = PULSE_NULLPTR;
entries[i].storageTexture.access = WGPUStorageTextureAccess_BindingNotUsed; entries[count].storageTexture.access = WGPUStorageTextureAccess_BindingNotUsed;
entries[i].storageTexture.format = WGPUTextureFormat_Undefined; entries[count].storageTexture.format = WGPUTextureFormat_Undefined;
entries[i].storageTexture.viewDimension = WGPUTextureViewDimension_Undefined; entries[count].storageTexture.viewDimension = WGPUTextureViewDimension_Undefined;
entries[i].texture.nextInChain = PULSE_NULLPTR; entries[count].texture.nextInChain = PULSE_NULLPTR;
entries[i].texture.multisampled = false; entries[count].texture.multisampled = false;
entries[i].texture.sampleType = WGPUTextureSampleType_BindingNotUsed; entries[count].texture.sampleType = WGPUTextureSampleType_BindingNotUsed;
entries[i].texture.viewDimension = WGPUTextureViewDimension_Undefined; entries[count].texture.viewDimension = WGPUTextureViewDimension_Undefined;
} }
} }
@@ -194,9 +348,13 @@ PulseComputePipeline WebGPUCreateComputePipeline(PulseDevice device, const Pulse
shader_descriptor.nextInChain = (const WGPUChainedStruct*)&source; shader_descriptor.nextInChain = (const WGPUChainedStruct*)&source;
webgpu_pipeline->shader = wgpuDeviceCreateShaderModule(webgpu_device->device, &shader_descriptor); 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); WebGPUBindGroupLayoutEntryInfo bindgroup_infos = { 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); 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] = { WGPUBindGroupLayout bind_group_layouts[3] = {
webgpu_pipeline->readonly_group, webgpu_pipeline->readonly_group,

6
Sources/Backends/WebGPU/WebGPUImage.c
View File

@@ -121,9 +121,9 @@ PulseImage WebGPUCreateImage(PulseDevice device, const PulseImageCreateInfo* cre
descriptor.usage = 0; descriptor.usage = 0;
if(create_infos->usage & PULSE_IMAGE_USAGE_STORAGE_READ) 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; 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) if(create_infos->usage & PULSE_IMAGE_USAGE_STORAGE_SIMULTANEOUS_READWRITE)
descriptor.usage |= WGPUTextureUsage_StorageBinding | WGPUTextureUsage_CopySrc | WGPUTextureUsage_CopyDst; descriptor.usage |= WGPUTextureUsage_StorageBinding | WGPUTextureUsage_CopySrc | WGPUTextureUsage_CopyDst;
@@ -224,7 +224,7 @@ bool WebGPUCopyImageToBuffer(PulseCommandList cmd, const PulseImageRegion* src,
WGPUTexelCopyTextureInfo texture_copy_info = { 0 }; WGPUTexelCopyTextureInfo texture_copy_info = { 0 };
texture_copy_info.texture = webgpu_src_image->texture; 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.aspect = WGPUTextureAspect_All;
texture_copy_info.origin.x = src->x; texture_copy_info.origin.x = src->x;
texture_copy_info.origin.y = src->y; texture_copy_info.origin.y = src->y;

42
Sources/PulseInternal.c
View File

@@ -98,3 +98,45 @@ uint32_t PulseHashCombine(uint32_t lhs, uint32_t rhs)
lhs ^= rhs + 0x9e3779b9 + (lhs << 6) + (lhs >> 2); lhs ^= rhs + 0x9e3779b9 + (lhs << 6) + (lhs >> 2);
return lhs; 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);
}

4
Sources/PulseInternal.h
View File

@@ -192,6 +192,10 @@ void PulseSetInternalError(PulseErrorType error);
uint32_t PulseHashString(const char* str); uint32_t PulseHashString(const char* str);
uint32_t PulseHashCombine(uint32_t lhs, uint32_t rhs); 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, ...); 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__) #define PulseLogError(backend, msg) PulseLogBackend(backend, PULSE_DEBUG_MESSAGE_SEVERITY_ERROR, msg, __FILE__, __FUNCTION__, __LINE__)

8
Tests/Buffer.c
View File

@@ -239,11 +239,11 @@ void TestBufferCopyImage()
PulseImageRegion dst_region = { 0 }; PulseImageRegion dst_region = { 0 };
dst_region.image = image; dst_region.image = image;
dst_region.width = 8; dst_region.width = 8;
dst_region.height = 1; dst_region.height = 0;
dst_region.depth = 1; dst_region.depth = 0;
dst_region.x = 1; dst_region.x = 0;
dst_region.y = 1; dst_region.y = 1;
dst_region.z = 1; dst_region.z = 0;
dst_region.layer = 1; dst_region.layer = 1;
TEST_ASSERT_TRUE_MESSAGE(PulseCopyBufferToImage(cmd, &src_region, &dst_region), PulseVerbaliseErrorType(PulseGetLastErrorType())); TEST_ASSERT_TRUE_MESSAGE(PulseCopyBufferToImage(cmd, &src_region, &dst_region), PulseVerbaliseErrorType(PulseGetLastErrorType()));

2
Tests/Shaders/WebGPU/BufferCopy.wgsl
View File

@@ -1,4 +1,4 @@
@group(0) @binding(0) var<storage, read_only> read_ssbo: array<u32>; @group(0) @binding(0) var<storage, read> read_ssbo: array<u32>;
@group(1) @binding(0) var<storage, read_write> write_ssbo: array<u32>; @group(1) @binding(0) var<storage, read_write> write_ssbo: array<u32>;
@compute @workgroup_size(16, 16, 1) @compute @workgroup_size(16, 16, 1)

2
Tests/Shaders/WebGPU/ReadWriteBindings.wgsl
View File

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

2
Tests/Shaders/WebGPU/WriteOnlyBindings.wgsl
View File

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