Skip to content

Commit

Permalink
Add storage texel format tests for textureLoad (gpuweb#3849)
Browse files Browse the repository at this point in the history 
Contributes to gpuweb#2268
Contributes to gpuweb#2156
Contributes to gpuweb#1262

* Add textureLoad tests for storage texture formats
  • Loading branch information
@alan-baker
alan-baker authored Jul 9, 2024
1 parent 982138f commit 07d7c07
Show file tree
Hide file tree
Showing 2 changed files with 338 additions and 0 deletions.
1 change: 1 addition & 0 deletions src/webgpu/listing_meta.json
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1552,6 +1552,7 @@
"webgpu:shader,execution,expression,call,builtin,textureLoad:sampled_1d:*": { "subcaseMS": 83.312 },
"webgpu:shader,execution,expression,call,builtin,textureLoad:sampled_2d:*": { "subcaseMS": 96.737 },
"webgpu:shader,execution,expression,call,builtin,textureLoad:sampled_3d:*": { "subcaseMS": 158.534 },
"webgpu:shader,execution,expression,call,builtin,textureLoad:storage_texel_formats:*": { "subcaseMS": 471.569 },
"webgpu:shader,execution,expression,call,builtin,textureNumLayers:arrayed:*": { "subcaseMS": 8.102 },
"webgpu:shader,execution,expression,call,builtin,textureNumLayers:sampled:*": { "subcaseMS": 2.101 },
"webgpu:shader,execution,expression,call,builtin,textureNumLayers:storage:*": { "subcaseMS": 8.000 },
Expand Down
337 changes: 337 additions & 0 deletions src/webgpu/shader/execution/expression/call/builtin/textureLoad.spec.ts
View file
Original file line number Diff line number Diff line change
Expand Up @@ -17,7 +17,16 @@ If an out of bounds access occurs, the built-in function returns one of:
`;

import { makeTestGroup } from '../../../../../../common/framework/test_group.js';
import { unreachable, iterRange } from '../../../../../../common/util/util.js';
import { GPUTest } from '../../../../../gpu_test.js';
import {
kFloat32Format,
kFloat16Format,
numberToFloatBits,
pack4x8unorm,
pack4x8snorm,
} from '../../../../../util/conversion.js';
import { TexelFormats } from '../../../../types.js';

import { generateCoordBoundaries } from './utils.js';

Expand Down Expand Up @@ -183,3 +192,331 @@ Parameters:
.combine('level', [-1, 0, `numlevels-1`, `numlevels`] as const)
)
.unimplemented();

// Returns texel values to use as inputs for textureLoad.
// Values are kept simple to avoid rounding issues.
function shaderValues(format: string, type: string) {
switch (type) {
case 'f32': {
switch (format) {
case 'rbga8snorm':
// prettier-ignore
return [
{ r: 0.0, g: 0.0, b: 0.0, a: 0.0, },
{ r: 0.2, g: 0.4, b: 0.6, a: 0.8, },
{ r: -0.2, g: -0.4, b: -0.6, a: -0.8, },
{ r: 0.2, g: -0.4, b: 0.6, a: -0.8, },
{ r: -0.2, g: 0.4, b: -0.6, a: 0.8, },
{ r: 0.2, g: 0.2, b: 0.2, a: 0.2, },
{ r: -0.2, g: -0.2, b: -0.2, a: -0.2, },
{ r: 0.4, g: 0.4, b: 0.4, a: 0.4, },
{ r: -0.4, g: -0.4, b: -0.4, a: -0.4, },
{ r: 0.6, g: 0.6, b: 0.6, a: 0.6, },
{ r: -0.6, g: -0.6, b: -0.6, a: -0.6, },
{ r: 0.8, g: 0.8, b: 0.8, a: 0.8, },
{ r: -0.8, g: -0.8, b: -0.8, a: -0.8, },
];
case 'rgba8unorm':
case 'bgra8unorm':
// prettier-ignore
return [
{ r: 0.0, g: 0.0, b: 0.0, a: 0.0, },
{ r: 0.2, g: 0.4, b: 0.6, a: 0.8, },
{ r: 0.9, g: 0.4, b: 0.6, a: 0.8, },
{ r: 0.2, g: 0.9, b: 0.6, a: 0.8, },
{ r: 0.2, g: 0.4, b: 0.9, a: 0.8, },
{ r: 0.2, g: 0.4, b: 0.6, a: 0.9, },
{ r: 0.2, g: 0.2, b: 0.2, a: 0.2, },
{ r: 0.4, g: 0.4, b: 0.4, a: 0.4, },
{ r: 0.6, g: 0.6, b: 0.6, a: 0.6, },
{ r: 0.8, g: 0.8, b: 0.8, a: 0.8, },
];
default:
// Stick within 16-bit ranges.
// prettier-ignore
return [
{ r: 100, g: 128, b: 100, a: 128, },
{ r: 64, g: 32, b: 32, a: 64, },
{ r: 8, g: 0, b: 8, a: 0, },
{ r: 0, g: 0, b: 0, a: 0, },
{ r: -100, g: 128, b: 100, a: 128, },
{ r: -64, g: 32, b: 32, a: 64, },
{ r: -8, g: 0, b: 8, a: 0, },
{ r: 100, g: -128, b: 100, a: 128, },
{ r: 64, g: -32, b: 32, a: 64, },
{ r: 8, g: 0, b: 8, a: 0, },
{ r: 100, g: 128, b: -100, a: 128, },
{ r: 64, g: 32, b: -32, a: 64, },
{ r: 8, g: 0, b: -8, a: 0, },
{ r: 100, g: 128, b: 100, a: -128, },
{ r: 64, g: 32, b: 32, a: -64, },
{ r: 8, g: 0, b: 8, a: 0, },
];
}
break;
}
case 'u32':
// Keep all ranges within u8.
// prettier-ignore
return [
{ r: 0, g: 0, b: 0, a: 0, },
{ r: 0, g: 8, b: 16, a: 128, },
{ r: 8, g: 16, b: 32, a: 64, },
{ r: 16, g: 32, b: 64, a: 128, },
{ r: 255, g: 254, b: 253, a: 252, },
{ r: 255, g: 255, b: 255, a: 255, },
{ r: 128, g: 64, b: 32, a: 16, },
{ r: 64, g: 32, b: 16, a: 8, },
{ r: 32, g: 16, b: 8, a: 0, },
];
case 'i32':
// Keep all ranges i8
// prettier-ignore
return [
{ r: 0, g: 0, b: 0, a: 0, },
{ r: 0, g: -8, b: 16, a: 127, },
{ r: 8, g: 16, b: -32, a: 64, },
{ r: -16, g: 32, b: 64, a: -128, },
{ r: 127, g: 126, b: 125, a: 124, },
{ r: -128, g: -127, b: -126, a: -125, },
{ r: 127, g: 127, b: 127, a: 127, },
{ r: -128, g: -128, b: -128, a: -128, },
];
default:
unreachable(`unhandled shader type ${type}`);
break;
}
return [];
}

g.test('storage_texel_formats')
.desc('Test loading of texel formats')
.params(u => u.combineWithParams([...TexelFormats, { format: 'bgra8unorm', _shaderType: 'f32' }]))
.beforeAllSubcases(t => {
t.skipIf(!t.hasLanguageFeature('readonly_and_readwrite_storage_textures'));
if (t.params.format === 'bgra8unorm') {
t.selectDeviceOrSkipTestCase('bgra8unorm-storage');
} else {
t.skipIfTextureFormatNotUsableAsStorageTexture(t.params.format as GPUTextureFormat);
}
})
.fn(t => {
const { format, _shaderType } = t.params;
const values = shaderValues(format, _shaderType);

// To avoid rounding issues, unorm and snorm values are repacked in the shader.
let useType = _shaderType;
let assignValue = `v`;
if (format === 'bgra8unorm' || format === 'rgba8unorm') {
useType = 'u32';
assignValue = `vec4u(pack4x8unorm(v),0,0,0)`;
} else if (format === 'rgba8snorm') {
useType = 'u32';
assignValue = `vec4u(pack4x8snorm(v),0,0,0)`;
}
const wgsl = `
requires readonly_and_readwrite_storage_textures;
@group(0) @binding(0)
var tex : texture_storage_1d<${format}, read>;
@group(0) @binding(1)
var<storage, read_write> out : array<vec4<${useType}>>;
@compute @workgroup_size(${values.length})
fn main(@builtin(global_invocation_id) gid : vec3u) {
let v = textureLoad(tex, gid.x);
out[gid.x] = ${assignValue};
}`;

const bytesPerRow = 256;
let bytesPerTexel = 4;
switch (format) {
case 'rgba16uint':
case 'rgba16sint':
case 'rgba16float':
case 'rg32uint':
case 'rg32sint':
case 'rg32float':
bytesPerTexel = 8;
break;
case 'rgba32uint':
case 'rgba32sint':
case 'rgba32float':
bytesPerTexel = 16;
break;
default:
break;
}

const textureSize: GPUExtent3D = {
width: bytesPerRow / bytesPerTexel,
height: 1,
depthOrArrayLayers: 1,
};
const texture = t.createTextureTracked({
format: format as GPUTextureFormat,
dimension: '1d',
size: textureSize,
mipLevelCount: 1,
usage: GPUTextureUsage.STORAGE_BINDING | GPUTextureUsage.COPY_DST,
});
const outputBuffer = t.makeBufferWithContents(
new Uint32Array([...iterRange(values.length * 4, x => 0)]),
GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE
);
t.trackForCleanup(outputBuffer);

const transformed = values.flatMap(x => {
switch (format) {
case 'rgba8unorm':
return pack4x8unorm(x.r, x.g, x.b, x.a);
case 'bgra8unorm':
return pack4x8unorm(x.b, x.g, x.r, x.a);
case 'rgba8snorm':
return pack4x8snorm(x.r, x.g, x.b, x.a);
case 'r32uint':
case 'r32sint':
return x.r;
case 'rg32uint':
case 'rg32sint':
return [x.r, x.g];
case 'rgba32uint':
case 'rgba32sint':
return [x.r, x.g, x.b, x.a];
case 'rgba8uint':
case 'rgba8sint':
return (x.r & 0xff) | ((x.g & 0xff) << 8) | ((x.b & 0xff) << 16) | ((x.a & 0xff) << 24);
case 'rgba16uint':
case 'rgba16sint':
return [(x.r & 0xffff) | ((x.g & 0xffff) << 16), (x.b & 0xffff) | ((x.a & 0xffff) << 16)];
case 'r32float':
return numberToFloatBits(x.r, kFloat32Format);
case 'rg32float':
return [numberToFloatBits(x.r, kFloat32Format), numberToFloatBits(x.g, kFloat32Format)];
case 'rgba32float':
return [
numberToFloatBits(x.r, kFloat32Format),
numberToFloatBits(x.g, kFloat32Format),
numberToFloatBits(x.b, kFloat32Format),
numberToFloatBits(x.a, kFloat32Format),
];
case 'rgba16float':
return [
(numberToFloatBits(x.r, kFloat16Format) & 0xffff) |
((numberToFloatBits(x.g, kFloat16Format) & 0xffff) << 16),
(numberToFloatBits(x.b, kFloat16Format) & 0xffff) |
((numberToFloatBits(x.a, kFloat16Format) & 0xffff) << 16),
];
default:
unreachable(`unhandled format ${format}`);
break;
}
return 0;
});

const texelBuffer = t.makeBufferWithContents(
new Uint32Array([
...iterRange(bytesPerRow, x => {
if (x < transformed.length) {
return transformed[x];
} else {
return 0;
}
}),
]),
GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE
);
t.trackForCleanup(texelBuffer);

const pipeline = t.device.createComputePipeline({
layout: 'auto',
compute: {
module: t.device.createShaderModule({
code: wgsl,
}),
entryPoint: 'main',
},
});
const bg = t.device.createBindGroup({
layout: pipeline.getBindGroupLayout(0),
entries: [
{
binding: 0,
resource: texture.createView({
format: format as GPUTextureFormat,
dimension: '1d',
}),
},
{
binding: 1,
resource: {
buffer: outputBuffer,
},
},
],
});

const encoder = t.device.createCommandEncoder();
encoder.copyBufferToTexture(
{
buffer: texelBuffer,
offset: 0,
bytesPerRow,
rowsPerImage: 1,
},
{ texture },
textureSize
);

const pass = encoder.beginComputePass();
pass.setPipeline(pipeline);
pass.setBindGroup(0, bg);
pass.dispatchWorkgroups(1, 1, 1);
pass.end();
t.queue.submit([encoder.finish()]);

const expected = new Uint32Array(
values.flatMap(x => {
switch (format) {
case 'r32uint':
case 'r32sint':
return [x.r, 0, 0, 1];
case 'rg32uint':
case 'rg32sint':
return [x.r, x.g, 0, 1];
case 'r32float':
return [
numberToFloatBits(x.r, kFloat32Format),
0,
0,
numberToFloatBits(1, kFloat32Format),
];
case 'rg32float':
return [
numberToFloatBits(x.r, kFloat32Format),
numberToFloatBits(x.g, kFloat32Format),
0,
numberToFloatBits(1, kFloat32Format),
];
case 'rgba32float':
case 'rgba16float':
return [
numberToFloatBits(x.r, kFloat32Format),
numberToFloatBits(x.g, kFloat32Format),
numberToFloatBits(x.b, kFloat32Format),
numberToFloatBits(x.a, kFloat32Format),
];
case 'rgba8unorm':
case 'bgra8unorm':
return [pack4x8unorm(x.r, x.g, x.b, x.a), 0, 0, 0];
case 'rgba8snorm':
return [pack4x8snorm(x.r, x.g, x.b, x.a), 0, 0, 0];
default:
break;
}
return [x.r, x.g, x.b, x.a];
})
);
t.expectGPUBufferValuesEqual(outputBuffer, expected);
});

0 comments on commit 07d7c07

Please sign in to comment.