import { keysOf } from '../../../../../../common/util/data_tables.js';
import { assert, range, unreachable, hasFeature } from '../../../../../../common/util/util.js';
import { Float16Array } from '../../../../../../external/petamoriken/float16/float16.js';
import {
ColorTextureFormat,
EncodableTextureFormat,
getBlockInfoForColorTextureFormat,
getBlockInfoForTextureFormat,
getTextureFormatType,
isColorTextureFormat,
isCompressedFloatTextureFormat,
isCompressedTextureFormat,
isDepthOrStencilTextureFormat,
isDepthTextureFormat,
isEncodableTextureFormat,
isTextureFormatFilterable,
isSintOrUintFormat,
isStencilTextureFormat,
kEncodableTextureFormats,
} from '../../../../../format_info.js';
import { GPUTest } from '../../../../../gpu_test.js';
import {
align,
clamp,
dotProduct,
hashU32,
lcm,
lerp,
quantizeToF32,
} from '../../../../../util/math.js';
import {
effectiveViewDimensionForDimension,
physicalMipSize,
physicalMipSizeFromTexture,
reifyTextureDescriptor,
SampleCoord,
virtualMipSize,
} from '../../../../../util/texture/base.js';
import {
kTexelRepresentationInfo,
NumericRange,
PerComponentNumericRange,
PerTexelComponent,
TexelComponent,
TexelRepresentationInfo,
} from '../../../../../util/texture/texel_data.js';
import { PerPixelAtLevel, TexelView } from '../../../../../util/texture/texel_view.js';
import {
copyTexelViewsToTexture,
createTextureFromTexelViews,
} from '../../../../../util/texture.js';
import { reifyExtent3D, reifyOrigin3D } from '../../../../../util/unions.js';
import { ShaderStage } from '../../../../validation/decl/util.js';
export const kShortShaderStageToShaderStage = {
c: 'compute' as ShaderStage,
f: 'fragment' as ShaderStage,
v: 'vertex' as ShaderStage,
} as const;
export const kShortShaderStages = keysOf(kShortShaderStageToShaderStage);
export type ShortShaderStage = (typeof kShortShaderStages)[number];
export const kShortAddressModeToAddressMode: Record<string, GPUAddressMode> = {
c: 'clamp-to-edge',
r: 'repeat',
m: 'mirror-repeat',
};
export const kShortAddressModes = keysOf(kShortAddressModeToAddressMode);
export const kSampleTypeInfo = {
f32: {
format: 'rgba8unorm',
},
i32: {
format: 'rgba8sint',
},
u32: {
format: 'rgba8uint',
},
} as const;
* Return the texture type for a given view dimension
*/
export function getTextureTypeForTextureViewDimension(viewDimension: GPUTextureViewDimension) {
switch (viewDimension) {
case '1d':
return 'texture_1d<f32>';
case '2d':
return 'texture_2d<f32>';
case '2d-array':
return 'texture_2d_array<f32>';
case '3d':
return 'texture_3d<f32>';
case 'cube':
return 'texture_cube<f32>';
case 'cube-array':
return 'texture_cube_array<f32>';
default:
unreachable();
}
}
const isUnencodableDepthFormat = (format: GPUTextureFormat) =>
format === 'depth24plus' ||
format === 'depth24plus-stencil8' ||
format === 'depth32float-stencil8';
* Skips a subcase if the filter === 'linear' and the format is type
* 'unfilterable-float' and we cannot enable filtering.
*/
export function skipIfTextureFormatNotSupportedOrNeedsFilteringAndIsUnfilterable(
t: GPUTest,
filter: GPUFilterMode,
format: GPUTextureFormat
) {
t.skipIfTextureFormatNotSupported(format);
if (filter === 'linear') {
t.skipIfTextureFormatNotFilterable(format);
}
}
* Returns if a texture format can be filled with random data.
*/
export function isFillable(format: GPUTextureFormat) {
return !isCompressedTextureFormat(format) || !format.endsWith('float');
}
* Returns if a texture format can potentially be filtered and can be filled with random data.
*/
export function isPotentiallyFilterableAndFillable(format: GPUTextureFormat) {
const type = getTextureFormatType(format);
const canPotentiallyFilter =
type === 'float' || type === 'unfilterable-float' || type === 'depth';
const result = canPotentiallyFilter && isFillable(format);
return result;
}
const builtinNeedsMipLevelWeights = (builtin: TextureBuiltin) =>
builtin !== 'textureLoad' &&
builtin !== 'textureGather' &&
builtin !== 'textureGatherCompare' &&
builtin !== 'textureSampleBaseClampToEdge';
* Splits in array into multiple arrays where every Nth value goes to a different array
*/
function unzip<T>(array: T[], num: number, srcStride?: number) {
srcStride = srcStride === undefined ? num : srcStride;
const arrays: T[][] = range(num, () => []);
const numEntries = Math.ceil(array.length / srcStride);
for (let i = 0; i < numEntries; ++i) {
for (let j = 0; j < num; ++j) {
arrays[j].push(array[i * srcStride + j]);
}
}
return arrays;
}
type MipWeights = {
sampleLevelWeights?: number[];
softwareMixToGPUMixGradWeights?: number[];
};
type MipWeightType = keyof MipWeights;
function makeGraph(width: number, height: number) {
const data = new Uint8Array(width * height);
return {
plot(norm: number, x: number, c: number) {
const y = clamp(Math.floor(norm * height), { min: 0, max: height - 1 });
const offset = (height - y - 1) * width + x;
data[offset] = c;
},
plotValues(values: Iterable<number>, c: number) {
let i = 0;
for (const v of values) {
this.plot(v, i, c);
++i;
}
},
toString(conversion = ['.', 'e', 'A']) {
const lines = [];
for (let y = 0; y < height; ++y) {
const offset = y * width;
lines.push([...data.subarray(offset, offset + width)].map(v => conversion[v]).join(''));
}
return lines.join('\n');
},
};
}
function* linear0to1OverN(n: number) {
for (let i = 0; i <= n; ++i) {
yield i / n;
}
}
* Generates an ascii graph of weights
*/
export function graphWeights(height: number, weights: number[]) {
const graph = makeGraph(weights.length, height);
graph.plotValues(linear0to1OverN(weights.length - 1), 1);
graph.plotValues(weights, 2);
return graph.toString();
}
* Validates the weights go from 0 to 1 in increasing order.
*/
function validateWeights(t: GPUTest, stage: string, weights: number[]) {
const showWeights = t.rec.debugging
? () => `
${weights.map((v, i) => `${i.toString().padStart(2)}: ${v}`).join('\n')}
e = expected
A = actual
${graphWeights(32, weights)}
`
: () => ``;
assert(
weights[0] === 0,
`stage: ${stage}, weight 0 expected 0 but was ${weights[0]}\n${showWeights()}`
);
assert(
weights[kMipLevelWeightSteps] === 1,
`stage: ${stage}, top weight expected 1 but was ${
weights[kMipLevelWeightSteps]
}\n${showWeights()}`
);
const kMinPercentUniqueWeights = 25;
assert(
new Set(weights).size >= ((weights.length * kMinPercentUniqueWeights * 0.01) | 0),
`stage: ${stage}, expected at least ~${kMinPercentUniqueWeights}% unique weights\n${showWeights()}`
);
}
* In an attempt to pass on more devices without lowering the tolerances
* so low they are meaningless, we ask the hardware to tell us, for a given
* gradient, level, what mix weights are being used.
*
* This is done by drawing instanced quads and using instance_index to
* write out results into an array. We sample a 2x2 pixel texture with
* 2 mip levels and set the 2nd mip level to white. This means the value
* we get back represents the weight used to mix the 2 mip levels.
*
* Just as a record of some differences across GPUs
*
* level weights: mapping from the mip level
* parameter of `textureSampleLevel` to
* the mix weight used by the GPU
*
* +--------+--------+--------+--------+
* | | | intel | amd |
* | | m1 | gen-9 | rna-1 |
* | level | mac | mac | mac |
* +--------+--------+--------+--------+
* | 0.0000 | 0.0000 | 0.0000 | 0.0000 |
* | 0.0313 | 0.0314 | 0.0313 | 0.0000 |
* | 0.0625 | 0.0625 | 0.0625 | 0.0000 |
* | 0.0938 | 0.0939 | 0.0938 | 0.0000 |
* | 0.1250 | 0.1250 | 0.1250 | 0.0313 |
* | 0.1563 | 0.1564 | 0.1563 | 0.0703 |
* | 0.1875 | 0.1875 | 0.1875 | 0.1094 |
* | 0.2188 | 0.2189 | 0.2188 | 0.1484 |
* | 0.2500 | 0.2500 | 0.2500 | 0.1875 |
* | 0.2813 | 0.2814 | 0.2813 | 0.2266 |
* | 0.3125 | 0.3125 | 0.3125 | 0.2656 |
* | 0.3438 | 0.3439 | 0.3438 | 0.3047 |
* | 0.3750 | 0.3750 | 0.3750 | 0.3438 |
* | 0.4063 | 0.4064 | 0.4063 | 0.3828 |
* | 0.4375 | 0.4375 | 0.4375 | 0.4219 |
* | 0.4688 | 0.4689 | 0.4688 | 0.4609 |
* | 0.5000 | 0.5000 | 0.5000 | 0.5000 |
* | 0.5313 | 0.5314 | 0.5313 | 0.5391 |
* | 0.5625 | 0.5625 | 0.5625 | 0.5781 |
* | 0.5938 | 0.5939 | 0.5938 | 0.6172 |
* | 0.6250 | 0.6250 | 0.6250 | 0.6563 |
* | 0.6563 | 0.6564 | 0.6563 | 0.6953 |
* | 0.6875 | 0.6875 | 0.6875 | 0.7344 |
* | 0.7188 | 0.7189 | 0.7188 | 0.7734 |
* | 0.7500 | 0.7500 | 0.7500 | 0.8125 |
* | 0.7813 | 0.7814 | 0.7813 | 0.8516 |
* | 0.8125 | 0.8125 | 0.8125 | 0.8906 |
* | 0.8438 | 0.8439 | 0.8438 | 0.9297 |
* | 0.8750 | 0.8750 | 0.8750 | 0.9688 |
* | 0.9063 | 0.9064 | 0.9063 | 1.0000 |
* | 0.9375 | 0.9375 | 0.9375 | 1.0000 |
* | 0.9688 | 0.9689 | 0.9688 | 1.0000 |
* | 1.0000 | 1.0000 | 1.0000 | 1.0000 |
* +--------+--------+--------+--------+
*
* grad weights: mapping from ddx value
* passed into `textureSampleGrad` to
* the mix weight used by the GPU
*
* +--------+--------+--------+--------+
* | | | intel | amd |
* | | m1 | gen-9 | rna-1 |
* | ddx | mac | mac | mac |
* +--------+--------+--------+--------+
* | 0.5000 | 0.0000 | 0.0000 | 0.0000 |
* | 0.5109 | 0.0390 | 0.0430 | 0.0000 |
* | 0.5221 | 0.0821 | 0.0859 | 0.0000 |
* | 0.5336 | 0.1211 | 0.1289 | 0.0352 |
* | 0.5453 | 0.1600 | 0.1719 | 0.0898 |
* | 0.5572 | 0.2032 | 0.2109 | 0.1328 |
* | 0.5694 | 0.2422 | 0.2461 | 0.1797 |
* | 0.5819 | 0.2814 | 0.2852 | 0.2305 |
* | 0.5946 | 0.3203 | 0.3203 | 0.2773 |
* | 0.6076 | 0.3554 | 0.3594 | 0.3164 |
* | 0.6209 | 0.3868 | 0.3906 | 0.3633 |
* | 0.6345 | 0.4218 | 0.4258 | 0.4063 |
* | 0.6484 | 0.4532 | 0.4609 | 0.4492 |
* | 0.6626 | 0.4882 | 0.4922 | 0.4883 |
* | 0.6771 | 0.5196 | 0.5234 | 0.5273 |
* | 0.6920 | 0.5507 | 0.5547 | 0.5664 |
* | 0.7071 | 0.5860 | 0.5859 | 0.6055 |
* | 0.7226 | 0.6132 | 0.6133 | 0.6406 |
* | 0.7384 | 0.6407 | 0.6445 | 0.6797 |
* | 0.7546 | 0.6679 | 0.6719 | 0.7148 |
* | 0.7711 | 0.6953 | 0.6992 | 0.7461 |
* | 0.7880 | 0.7225 | 0.7266 | 0.7813 |
* | 0.8052 | 0.7500 | 0.7539 | 0.8164 |
* | 0.8229 | 0.7814 | 0.7813 | 0.8516 |
* | 0.8409 | 0.8086 | 0.8086 | 0.8828 |
* | 0.8593 | 0.8321 | 0.8320 | 0.9141 |
* | 0.8781 | 0.8554 | 0.8594 | 0.9492 |
* | 0.8974 | 0.8789 | 0.8828 | 0.9766 |
* | 0.9170 | 0.9025 | 0.9063 | 1.0000 |
* | 0.9371 | 0.9297 | 0.9297 | 1.0000 |
* | 0.9576 | 0.9532 | 0.9531 | 1.0000 |
* | 0.9786 | 0.9765 | 0.9766 | 1.0000 |
* | 1.0000 | 1.0000 | 1.0000 | 1.0000 |
* +--------+--------+--------+--------+
*/
export async function queryMipLevelMixWeightsForDevice(t: GPUTest, stage: ShaderStage) {
const { device } = t;
const kNumWeightTypes = 2;
assert(kNumWeightTypes <= 4);
const module = device.createShaderModule({
code: `
@group(0) @binding(0) var tex: texture_2d<f32>;
@group(0) @binding(1) var smp: sampler;
@group(0) @binding(2) var<storage, read_write> result: array<vec4f>;
struct VSOutput {
@builtin(position) pos: vec4f,
@location(0) @interpolate(flat, either) ndx: u32,
@location(1) @interpolate(flat, either) result: vec4f,
};
fn getMixLevels(wNdx: u32) -> vec4f {
let mipLevel = f32(wNdx) / ${kMipLevelWeightSteps};
let size = textureDimensions(tex);
let g = mix(1.0, 2.0, mipLevel) / f32(size.x);
let ddx = vec2f(g, 0);
return vec4f(
textureSampleLevel(tex, smp, vec2f(0.5), mipLevel).r,
textureSampleGrad(tex, smp, vec2f(0.5), ddx, vec2f(0)).r,
0,
0);
}
fn getPosition(vNdx: u32) -> vec4f {
let pos = array(
vec2f(-1, 3),
vec2f( 3, -1),
vec2f(-1, -1),
);
let p = pos[vNdx];
return vec4f(p, 0, 1);
}
// -- for getting fragment stage weights --
@vertex fn vs(@builtin(vertex_index) vNdx: u32, @builtin(instance_index) iNdx: u32) -> VSOutput {
return VSOutput(getPosition(vNdx), iNdx, vec4f(0));
}
@fragment fn fsRecord(v: VSOutput) -> @location(0) vec4u {
return bitcast<vec4u>(getMixLevels(v.ndx));
}
// -- for getting compute stage weights --
@compute @workgroup_size(1) fn csRecord(@builtin(global_invocation_id) id: vec3u) {
result[id.x] = getMixLevels(id.x);
}
// -- for getting vertex stage weights --
@vertex fn vsRecord(@builtin(vertex_index) vNdx: u32, @builtin(instance_index) iNdx: u32) -> VSOutput {
return VSOutput(getPosition(vNdx), iNdx, getMixLevels(iNdx));
}
@fragment fn fsSaveVs(v: VSOutput) -> @location(0) vec4u {
return bitcast<vec4u>(v.result);
}
`,
});
const texture = t.createTextureTracked({
size: [2, 2, 1],
format: 'r8unorm',
usage: GPUTextureUsage.TEXTURE_BINDING | GPUTextureUsage.COPY_DST,
mipLevelCount: 2,
});
device.queue.writeTexture(
{ texture, mipLevel: 1 },
new Uint8Array([255]),
{ bytesPerRow: 1 },
[1, 1]
);
const sampler = device.createSampler({
minFilter: 'linear',
magFilter: 'linear',
mipmapFilter: 'linear',
});
const target = t.createTextureTracked({
size: [kMipLevelWeightSteps + 1, 1],
format: 'rgba32uint',
usage: GPUTextureUsage.RENDER_ATTACHMENT | GPUTextureUsage.COPY_SRC,
});
const storageBuffer = t.createBufferTracked({
label: 'queryMipLevelMixWeightsForDevice:storageBuffer',
size: 4 * 4 * (kMipLevelWeightSteps + 1),
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
});
const resultBuffer = t.createBufferTracked({
label: 'queryMipLevelMixWeightsForDevice:resultBuffer',
size: align(storageBuffer.size, 256),
usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ,
});
const createBindGroup = (pipeline: GPUComputePipeline | GPURenderPipeline) =>
device.createBindGroup({
layout: pipeline.getBindGroupLayout(0),
entries: [
{ binding: 0, resource: texture.createView() },
{ binding: 1, resource: sampler },
...(stage === 'compute' ? [{ binding: 2, resource: { buffer: storageBuffer } }] : []),
],
});
const encoder = device.createCommandEncoder({ label: 'queryMipLevelMixWeightsForDevice' });
switch (stage) {
case 'compute': {
const pipeline = device.createComputePipeline({
layout: 'auto',
compute: { module },
});
const pass = encoder.beginComputePass();
pass.setPipeline(pipeline);
pass.setBindGroup(0, createBindGroup(pipeline));
pass.dispatchWorkgroups(kMipLevelWeightSteps + 1);
pass.end();
encoder.copyBufferToBuffer(storageBuffer, 0, resultBuffer, 0, storageBuffer.size);
break;
}
case 'fragment': {
const pipeline = device.createRenderPipeline({
layout: 'auto',
vertex: { module, entryPoint: 'vs' },
fragment: { module, entryPoint: 'fsRecord', targets: [{ format: 'rgba32uint' }] },
});
const pass = encoder.beginRenderPass({
colorAttachments: [
{
view: target.createView(),
loadOp: 'clear',
storeOp: 'store',
},
],
});
pass.setPipeline(pipeline);
pass.setBindGroup(0, createBindGroup(pipeline));
for (let x = 0; x <= kMipLevelWeightSteps; ++x) {
pass.setViewport(x, 0, 1, 1, 0, 1);
pass.draw(3, 1, 0, x);
}
pass.end();
encoder.copyTextureToBuffer({ texture: target }, { buffer: resultBuffer }, [target.width]);
break;
}
case 'vertex': {
const pipeline = device.createRenderPipeline({
layout: 'auto',
vertex: { module, entryPoint: 'vsRecord' },
fragment: { module, entryPoint: 'fsSaveVs', targets: [{ format: 'rgba32uint' }] },
});
const pass = encoder.beginRenderPass({
colorAttachments: [
{
view: target.createView(),
loadOp: 'clear',
storeOp: 'store',
},
],
});
pass.setPipeline(pipeline);
pass.setBindGroup(0, createBindGroup(pipeline));
for (let x = 0; x <= kMipLevelWeightSteps; ++x) {
pass.setViewport(x, 0, 1, 1, 0, 1);
pass.draw(3, 1, 0, x);
}
pass.end();
encoder.copyTextureToBuffer({ texture: target }, { buffer: resultBuffer }, [target.width]);
break;
}
}
device.queue.submit([encoder.finish()]);
await resultBuffer.mapAsync(GPUMapMode.READ);
const result = Array.from(
new Float32Array(resultBuffer.getMappedRange(0, (kMipLevelWeightSteps + 1) * 16))
);
resultBuffer.unmap();
resultBuffer.destroy();
const [sampleLevelWeights, gradWeights] = unzip(result, kNumWeightTypes, 4);
validateWeights(t, stage, sampleLevelWeights);
validateWeights(t, stage, gradWeights);
texture.destroy();
storageBuffer.destroy();
return {
sampleLevelWeights,
softwareMixToGPUMixGradWeights: generateSoftwareMixToGPUMixGradWeights(gradWeights, texture),
};
}
function getIndexAndWeight(values: readonly number[], v: number) {
assert(v >= values[0] && v <= values[values.length - 1]);
let lo = 0;
let hi = values.length - 1;
for (;;) {
const i = (lo + (hi - lo) / 2) | 0;
const w0 = values[i];
const w1 = values[i + 1];
if (lo === hi || (v >= w0 && v <= w1)) {
const weight = (v - w0) / (w1 - w0);
return [i, weight];
}
if (v < w0) {
hi = i;
} else {
lo = i + 1;
}
}
}
* Given a fractional number between 0 and values.length returns the value between
* 2 values. Effectively lerp(values[ndx], values[ndx + 1], weight)
*/
function bilinearFilter(values: readonly number[], ndx: number, weight: number) {
const v0 = values[ndx];
const v1 = values[ndx + 1] ?? 0;
assert(ndx < values.length - 1 || (ndx === values.length - 1 && weight === 0));
return lerp(v0, v1, weight);
}
* Generates an array of values that maps between the software renderer's gradient computed
* mip level and the GPUs gradient computed mip level for mip level 0 to 1.
*/
function generateSoftwareMixToGPUMixGradWeights(gpuWeights: number[], texture: GPUTexture) {
const numSteps = gpuWeights.length - 1;
const size = [texture.width, texture.height, texture.depthOrArrayLayers];
const softwareWeights = range(numSteps + 1, i => {
const u = i / numSteps;
const g = lerp(1, 2, u) / texture.width;
const mipLevel = computeMipLevelFromGradients([g], [0], size);
assert(mipLevel >= 0 && mipLevel <= 1);
return mipLevel;
});
const softwareMixToGPUMixMap = range(numSteps + 1, i => {
const mix = i / numSteps;
const [ndx, weight] = getIndexAndWeight(softwareWeights, mix);
return bilinearFilter(gpuWeights, ndx, weight);
});
return softwareMixToGPUMixMap;
}
function mapSoftwareMipLevelToGPUMipLevel(t: GPUTest, stage: ShaderStage, mipLevel: number) {
const baseLevel = Math.floor(mipLevel);
const softwareMix = mipLevel - baseLevel;
const gpuMix = getMixWeightByTypeForMipLevel(
t,
stage,
'softwareMixToGPUMixGradWeights',
softwareMix
);
return baseLevel + gpuMix;
}
const euclideanModulo = (n: number, m: number) => ((n % m) + m) % m;
* Gets the mip gradient values for the current device.
* The issue is, different GPUs have different ways of mixing between mip levels.
* For most GPUs it's linear but for AMD GPUs on Mac in particular, it's something
* else (which AFAICT is against all the specs).
*
* We seemingly have 3 options:
*
* 1. Increase the tolerances of tests so they pass on AMD.
* 2. Mark AMD as failing
* 3. Try to figure out how the GPU converts mip levels into weights
*
* We're doing 3.
*
* There's an assumption that the gradient will be the same for all formats
* and usages.
*
* Note: The code below has 2 maps. One device->Promise, the other device->weights
* device->weights is meant to be used synchronously by other code so we don't
* want to leave initMipGradientValuesForDevice until the weights have been read.
* But, multiple subcases will run because this function is async. So, subcase 1
* runs, hits this init code, this code waits for the weights. Then, subcase 2
* runs and hits this init code. The weights will not be in the device->weights map
* yet which is why we have the device->Promise map. This is so subcase 2 waits
* for subcase 1's "query the weights" step. Otherwise, all subcases would do the
* "get the weights" step separately.
*/
const kMipLevelWeightSteps = 64;
const s_deviceToMipLevelWeightsPromise = new WeakMap<
GPUDevice,
Record<ShaderStage, Promise<MipWeights>>
>();
const s_deviceToMipLevelWeights = new WeakMap<GPUDevice, Record<ShaderStage, MipWeights>>();
async function initMipLevelWeightsForDevice(t: GPUTest, stage: ShaderStage) {
const { device } = t;
const stageWeightsP =
s_deviceToMipLevelWeightsPromise.get(device) ??
({} as Record<ShaderStage, Promise<MipWeights>>);
s_deviceToMipLevelWeightsPromise.set(device, stageWeightsP);
let weightsP = stageWeightsP[stage];
if (!weightsP) {
weightsP = queryMipLevelMixWeightsForDevice(t, stage);
weightsP
.then(weights => {
const stageWeights =
s_deviceToMipLevelWeights.get(device) ?? ({} as Record<ShaderStage, MipWeights>);
s_deviceToMipLevelWeights.set(device, stageWeights);
stageWeights[stage] = weights;
})
.catch(e => {
throw e;
});
stageWeightsP[stage] = weightsP;
}
return await weightsP;
}
function getMixWeightByTypeForMipLevel(
t: GPUTest,
stage: ShaderStage,
weightType: MipWeightType | 'identity',
mipLevel: number
) {
if (weightType === 'identity') {
return euclideanModulo(mipLevel, 1);
}
const weights = s_deviceToMipLevelWeights.get(t.device)![stage][weightType];
assert(
!!weights,
'you must use WGSLTextureSampleTest or call initializeDeviceMipWeights before calling this function'
);
const steps = weights.length - 1;
const w = euclideanModulo(mipLevel, 1) * steps;
const lowerNdx = Math.floor(w);
const upperNdx = Math.ceil(w);
const mix = w % 1;
return lerp(weights[lowerNdx], weights[upperNdx], mix);
}
function getWeightForMipLevel(
t: GPUTest,
stage: ShaderStage,
weightType: MipWeightType | 'identity',
mipLevelCount: number,
mipLevel: number
) {
if (mipLevel < 0 || mipLevel >= mipLevelCount) {
return 1;
}
return getMixWeightByTypeForMipLevel(t, stage, weightType, mipLevel);
}
* Skip a test if the specific stage doesn't support storage textures.
*/
export function skipIfNoStorageTexturesInStage(t: GPUTest, stage: ShaderStage) {
if (t.isCompatibility) {
t.skipIf(
stage === 'fragment' && !(t.device.limits.maxStorageTexturesInFragmentStage! > 0),
'device does not support storage textures in fragment shaders'
);
t.skipIf(
stage === 'vertex' && !(t.device.limits.maxStorageTexturesInVertexStage! > 0),
'device does not support storage textures in vertex shaders'
);
}
}
* Runs a texture query like textureDimensions, textureNumLevels and expects
* a particular result.
*/
export function executeTextureQueryAndExpectResult(
t: GPUTest,
stage: ShaderStage,
code: string,
texture: GPUTexture | GPUExternalTexture,
viewDescriptor: GPUTextureViewDescriptor | undefined,
expected: number[]
) {
const { device } = t;
const returnType = `vec4<u32>`;
const castWGSL = `${returnType}(getValue()${range(4 - expected.length, () => ', 0').join('')})`;
const stageWGSL =
stage === 'vertex'
? `
// --------------------------- vertex stage shaders --------------------------------
@vertex fn vsVertex(
@builtin(vertex_index) vertex_index : u32,
@builtin(instance_index) instance_index : u32) -> VOut {
let positions = array(vec2f(-1, 3), vec2f(3, -1), vec2f(-1, -1));
return VOut(vec4f(positions[vertex_index], 0, 1),
instance_index,
${castWGSL});
}
@fragment fn fsVertex(v: VOut) -> @location(0) vec4u {
return bitcast<vec4u>(v.result);
}
`
: stage === 'fragment'
? `
// --------------------------- fragment stage shaders --------------------------------
@vertex fn vsFragment(
@builtin(vertex_index) vertex_index : u32,
@builtin(instance_index) instance_index : u32) -> VOut {
let positions = array(vec2f(-1, 3), vec2f(3, -1), vec2f(-1, -1));
return VOut(vec4f(positions[vertex_index], 0, 1), instance_index, ${returnType}(0));
}
@fragment fn fsFragment(v: VOut) -> @location(0) vec4u {
return bitcast<vec4u>(${castWGSL});
}
`
: `
// --------------------------- compute stage shaders --------------------------------
@group(1) @binding(0) var<storage, read_write> results: array<${returnType}>;
@compute @workgroup_size(1) fn csCompute(@builtin(global_invocation_id) id: vec3u) {
results[id.x] = ${castWGSL};
}
`;
const wgsl = `
${code}
struct VOut {
@builtin(position) pos: vec4f,
@location(0) @interpolate(flat, either) ndx: u32,
@location(1) @interpolate(flat, either) result: ${returnType},
};
${stageWGSL}
`;
const module = device.createShaderModule({ code: wgsl });
const visibility =
stage === 'compute'
? GPUShaderStage.COMPUTE
: stage === 'fragment'
? GPUShaderStage.FRAGMENT
: GPUShaderStage.VERTEX;
const entries: GPUBindGroupLayoutEntry[] = [];
if (code.includes('texture_external')) {
entries.push({
binding: 0,
visibility,
externalTexture: {},
});
} else if (code.includes('texture_storage')) {
assert(texture instanceof GPUTexture);
entries.push({
binding: 0,
visibility,
storageTexture: {
access: code.includes(', read>')
? 'read-only'
: code.includes(', write>')
? 'write-only'
: 'read-write',
viewDimension: viewDescriptor?.dimension ?? '2d',
format: texture.format,
},
});
} else {
assert(texture instanceof GPUTexture);
const sampleType =
viewDescriptor?.aspect === 'stencil-only'
? 'uint'
: code.includes('texture_depth')
? 'depth'
: isDepthTextureFormat(texture.format)
? 'unfilterable-float'
: isStencilTextureFormat(texture.format)
? 'uint'
: texture.sampleCount > 1 && getTextureFormatType(texture.format) === 'float'
? 'unfilterable-float'
: getTextureFormatType(texture.format) ?? 'unfilterable-float';
entries.push({
binding: 0,
visibility,
texture: {
sampleType,
viewDimension: viewDescriptor?.dimension ?? '2d',
multisampled: texture.sampleCount > 1,
},
});
}
const bindGroupLayouts: GPUBindGroupLayout[] = [device.createBindGroupLayout({ entries })];
if (stage === 'compute') {
bindGroupLayouts.push(
device.createBindGroupLayout({
entries: [
{
binding: 0,
visibility: GPUShaderStage.COMPUTE,
buffer: {
type: 'storage',
hasDynamicOffset: false,
minBindingSize: 16,
},
},
],
})
);
}
const layout = device.createPipelineLayout({
bindGroupLayouts,
});
let pipeline: GPUComputePipeline | GPURenderPipeline;
switch (stage) {
case 'compute':
pipeline = device.createComputePipeline({
layout,
compute: { module },
});
break;
case 'fragment':
case 'vertex':
pipeline = device.createRenderPipeline({
layout,
vertex: { module },
fragment: {
module,
targets: [{ format: 'rgba32uint' }],
},
});
break;
}
const bindGroup0 = device.createBindGroup({
layout: pipeline.getBindGroupLayout(0),
entries: [
{
binding: 0,
resource:
texture instanceof GPUExternalTexture ? texture : texture.createView(viewDescriptor),
},
],
});
const renderTarget = t.createTextureTracked({
format: 'rgba32uint',
size: [expected.length, 1],
usage: GPUTextureUsage.COPY_SRC | GPUTextureUsage.RENDER_ATTACHMENT,
});
const resultBuffer = t.createBufferTracked({
label: 'executeAndExpectResult:resultBuffer',
size: align(expected.length * 4, 256),
usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.COPY_SRC,
});
let storageBuffer: GPUBuffer | undefined;
const encoder = device.createCommandEncoder({ label: 'executeAndExpectResult' });
if (stage === 'compute') {
storageBuffer = t.createBufferTracked({
label: 'executeAndExpectResult:storageBuffer',
size: resultBuffer.size,
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
});
const bindGroup1 = device.createBindGroup({
layout: pipeline!.getBindGroupLayout(1),
entries: [{ binding: 0, resource: { buffer: storageBuffer } }],
});
const pass = encoder.beginComputePass();
pass.setPipeline(pipeline! as GPUComputePipeline);
pass.setBindGroup(0, bindGroup0);
pass.setBindGroup(1, bindGroup1);
pass.dispatchWorkgroups(expected.length);
pass.end();
encoder.copyBufferToBuffer(storageBuffer, 0, resultBuffer, 0, storageBuffer.size);
} else {
const pass = encoder.beginRenderPass({
colorAttachments: [
{
view: renderTarget.createView(),
loadOp: 'clear',
storeOp: 'store',
},
],
});
pass.setPipeline(pipeline! as GPURenderPipeline);
pass.setBindGroup(0, bindGroup0);
for (let i = 0; i < expected.length; ++i) {
pass.setViewport(i, 0, 1, 1, 0, 1);
pass.draw(3, 1, 0, i);
}
pass.end();
encoder.copyTextureToBuffer(
{ texture: renderTarget },
{
buffer: resultBuffer,
bytesPerRow: resultBuffer.size,
},
[renderTarget.width, 1]
);
}
t.device.queue.submit([encoder.finish()]);
const e = new Uint32Array(4);
e.set(expected);
t.expectGPUBufferValuesEqual(resultBuffer, e);
}
* Used to specify a range from [0, num)
* The type is used to determine if values should be integers and if they can be negative.
*/
export type RangeDef = {
num: number;
type: 'f32' | 'i32' | 'u32';
};
function getLimitValue(v: number) {
switch (v) {
case Number.POSITIVE_INFINITY:
return 1000;
case Number.NEGATIVE_INFINITY:
return -1000;
default:
return v;
}
}
function getMinAndMaxTexelValueForComponent(
rep: TexelRepresentationInfo,
component: TexelComponent
) {
assert(!!rep.numericRange);
const perComponentRanges = rep.numericRange as PerComponentNumericRange;
const perComponentRange = perComponentRanges[component];
const range = rep.numericRange as NumericRange;
const { min, max } = perComponentRange ? perComponentRange : range;
return { min: getLimitValue(min), max: getLimitValue(max) };
}
* We need the software rendering to do the same interpolation as the hardware
* rendered so for -srgb formats we set the TexelView to an -srgb format as
* TexelView handles this case. Note: It might be nice to add rgba32float-srgb
* or something similar to TexelView.
*/
export function getTexelViewFormatForTextureFormat(format: GPUTextureFormat) {
if (format.endsWith('sint')) {
return 'rgba32sint';
} else if (format.endsWith('uint')) {
return 'rgba32uint';
}
return format.endsWith('-srgb') ? 'rgba8unorm-srgb' : 'rgba32float';
}
const kTextureTypeInfo = {
depth: {
componentType: 'f32',
resultType: 'vec4f',
resultFormat: 'rgba32float',
sampleType: 'depth',
},
float: {
componentType: 'f32',
resultType: 'vec4f',
resultFormat: 'rgba32float',
sampleType: 'float',
},
'unfilterable-float': {
componentType: 'f32',
resultType: 'vec4f',
resultFormat: 'rgba32float',
sampleType: 'unfilterable-float',
},
sint: {
componentType: 'i32',
resultType: 'vec4i',
resultFormat: 'rgba32sint',
sampleType: 'sint',
},
uint: {
componentType: 'u32',
resultType: 'vec4u',
resultFormat: 'rgba32uint',
sampleType: 'uint',
},
} as const;
export function getTextureFormatTypeInfo(
format: GPUTextureFormat,
aspect: GPUTextureAspect = 'all'
) {
const type = getTextureFormatType(format, aspect);
assert(!!type);
return kTextureTypeInfo[type];
}
* given a texture type 'base', returns the base with the correct component for the given texture format.
* eg: `getTextureType('texture_2d', someUnsignedIntTextureFormat)` -> `texture_2d<u32>`
*/
export function appendComponentTypeForFormatToTextureType(base: string, format: GPUTextureFormat) {
return base.includes('depth')
? base
: `${base}<${getTextureFormatTypeInfo(format).componentType}>`;
}
type RandomTextureOptions = {
generator: PerPixelAtLevel<PerTexelComponent<number>>;
};
* Gets the baseMipLevel, mipLevelCount, size of the baseMipLevel,
* baseArrayLayer, and arrayLayerCount
* taking into account the texture descriptor and the view descriptor.
*/
function getBaseMipLevelInfo(textureInfo: SoftwareTexture) {
const baseMipLevel = textureInfo.viewDescriptor.baseMipLevel ?? 0;
const mipLevelCount =
textureInfo.viewDescriptor.mipLevelCount ??
(textureInfo.descriptor.mipLevelCount ?? 1) - baseMipLevel;
const baseMipLevelSize = virtualMipSize(
textureInfo.descriptor.dimension ?? '2d',
textureInfo.descriptor.size,
baseMipLevel
);
const baseArrayLayer = textureInfo.viewDescriptor.baseArrayLayer ?? 0;
const arrayLayerCount =
textureInfo.viewDescriptor.arrayLayerCount ?? baseMipLevelSize[2] - baseArrayLayer;
baseMipLevelSize[2] = arrayLayerCount;
return { baseMipLevel, baseMipLevelSize, mipLevelCount, baseArrayLayer, arrayLayerCount };
}
* Make a generator for texels for depth comparison tests.
*/
export function makeRandomDepthComparisonTexelGenerator(
info: {
format: GPUTextureFormat;
size: GPUExtent3D;
},
comparison: GPUCompareFunction
) {
const format = isUnencodableDepthFormat(info.format) ? 'depth32float' : info.format;
const rep = kTexelRepresentationInfo[format as EncodableTextureFormat];
const size = reifyExtent3D(info.size);
const comparisonIsEqualOrNotEqual = comparison === 'equal' || comparison === 'not-equal';
const fixedValues = [0, 0.6, 1, 1];
const encode = comparisonIsEqualOrNotEqual
? (norm: number) => fixedValues[(norm * (fixedValues.length - 1)) | 0]
: (norm: number) => norm;
return (coords: SampleCoord): Readonly<PerTexelComponent<number>> => {
const texel: PerTexelComponent<number> = {};
for (const component of rep.componentOrder) {
const rnd = hashU32(
coords.x,
coords.y,
coords.z,
coords.sampleIndex ?? 0,
component.charCodeAt(0),
size.width,
size.height,
size.depthOrArrayLayers
);
const normalized = clamp(rnd / 0xffffffff, { min: 0, max: 1 });
texel[component] = encode(normalized);
}
return quantize(texel, rep);
};
}
function createRandomTexelViewViaColors(
info: {
format: GPUTextureFormat;
size: GPUExtent3D;
mipLevel: number;
},
options?: RandomTextureOptions | undefined
): TexelView {
const rep = kTexelRepresentationInfo[info.format as EncodableTextureFormat];
const size = reifyExtent3D(info.size);
const minMax = Object.fromEntries(
rep.componentOrder.map(component => [
component,
getMinAndMaxTexelValueForComponent(rep, component),
])
);
const generator = (coords: SampleCoord): Readonly<PerTexelComponent<number>> => {
const texel: PerTexelComponent<number> = {};
for (const component of rep.componentOrder) {
const rnd = hashU32(
coords.x,
coords.y,
coords.z,
coords.sampleIndex ?? 0,
component.charCodeAt(0),
info.mipLevel,
size.width,
size.height,
size.depthOrArrayLayers
);
const normalized = clamp(rnd / 0xffffffff, { min: 0, max: 1 });
const { min, max } = minMax[component];
texel[component] = lerp(min, max, normalized);
}
return quantize(texel, rep);
};
return TexelView.fromTexelsAsColors(
info.format as EncodableTextureFormat,
options?.generator ?? generator
);
}
function createRandomTexelViewViaBytes(info: {
format: GPUTextureFormat;
size: GPUExtent3D;
mipLevel: number;
sampleCount: number;
}): TexelView {
const { format } = info;
const formatInfo = getBlockInfoForTextureFormat(format);
const rep = kTexelRepresentationInfo[info.format as EncodableTextureFormat];
assert(!!rep);
const { bytesPerBlock } = formatInfo;
assert(bytesPerBlock !== undefined && bytesPerBlock > 0);
const size = physicalMipSize(reifyExtent3D(info.size), info.format, '2d', 0);
const blocksAcross = Math.ceil(size.width / formatInfo.blockWidth);
const blocksDown = Math.ceil(size.height / formatInfo.blockHeight);
const bytesPerRow = blocksAcross * bytesPerBlock * info.sampleCount;
const bytesNeeded = bytesPerRow * blocksDown * size.depthOrArrayLayers;
const data = new Uint8Array(bytesNeeded);
const hashBase =
sumOfCharCodesOfString(info.format) +
size.width +
size.height +
size.depthOrArrayLayers +
info.mipLevel +
info.sampleCount;
if (info.format.includes('32float') || info.format.includes('16float')) {
const { min, max } = getMinAndMaxTexelValueForComponent(rep, TexelComponent.R);
const asFloat = info.format.includes('32float')
? new Float32Array(data.buffer)
: new Float16Array(data.buffer);
for (let i = 0; i < asFloat.length; ++i) {
asFloat[i] = lerp(min, max, hashU32(hashBase + i) / 0xffff_ffff);
}
} else if (bytesNeeded % 4 === 0) {
const asU32 = new Uint32Array(data.buffer);
for (let i = 0; i < asU32.length; ++i) {
asU32[i] = hashU32(hashBase + i);
}
} else {
for (let i = 0; i < bytesNeeded; ++i) {
data[i] = hashU32(hashBase + i);
}
}
return TexelView.fromTextureDataByReference(info.format as EncodableTextureFormat, data, {
bytesPerRow,
rowsPerImage: size.height,
subrectOrigin: [0, 0, 0],
subrectSize: size,
});
}
* Creates a TexelView filled with random values.
*/
function createRandomTexelView(
info: {
format: GPUTextureFormat;
size: GPUExtent3D;
mipLevel: number;
sampleCount: number;
},
options?: RandomTextureOptions | undefined
): TexelView {
const { format } = info;
assert(!isCompressedTextureFormat(format));
const type = getTextureFormatType(format);
const canFillWithRandomTypedData =
!options &&
isEncodableTextureFormat(format) &&
((format.includes('norm') && type !== 'depth') ||
format.includes('16float') ||
(format.includes('32float') && type !== 'depth') ||
type === 'sint' ||
type === 'uint');
return canFillWithRandomTypedData
? createRandomTexelViewViaBytes(info)
: createRandomTexelViewViaColors(info, options);
}
* Creates a mip chain of TexelViews filled with random values
*/
function createRandomTexelViewMipmap(
info: {
format: GPUTextureFormat;
size: GPUExtent3D;
mipLevelCount?: number;
dimension?: GPUTextureDimension;
sampleCount?: number;
},
options?: RandomTextureOptions | undefined
): TexelView[] {
const mipLevelCount = info.mipLevelCount ?? 1;
const dimension = info.dimension ?? '2d';
return range(mipLevelCount, i =>
createRandomTexelView(
{
format: info.format,
size: virtualMipSize(dimension, info.size, i),
mipLevel: i,
sampleCount: info.sampleCount ?? 1,
},
options
)
);
}
export type vec1 = [number];
export type vec2 = [number, number];
export type vec3 = [number, number, number];
export type vec4 = [number, number, number, number];
export type Dimensionality = vec1 | vec2 | vec3;
type TextureCallArgKeys = keyof TextureCallArgs<vec1>;
const kTextureCallArgNames: readonly TextureCallArgKeys[] = [
'component',
'coords',
'derivativeMult',
'arrayIndex',
'bias',
'sampleIndex',
'mipLevel',
'ddx',
'ddy',
'depthRef',
'offset',
] as const;
export interface TextureCallArgs<T extends Dimensionality> {
component?: number;
coords?: T;
derivativeMult?: T;
mipLevel?: number;
arrayIndex?: number;
bias?: number;
sampleIndex?: number;
depthRef?: number;
ddx?: T;
ddy?: T;
offset?: T;
}
export type TextureBuiltin =
| 'textureGather'
| 'textureGatherCompare'
| 'textureLoad'
| 'textureSample'
| 'textureSampleBaseClampToEdge'
| 'textureSampleBias'
| 'textureSampleCompare'
| 'textureSampleCompareLevel'
| 'textureSampleGrad'
| 'textureSampleLevel';
export interface TextureCall<T extends Dimensionality> extends TextureCallArgs<T> {
builtin: TextureBuiltin;
coordType: 'f' | 'i' | 'u';
levelType?: 'i' | 'u' | 'f';
arrayIndexType?: 'i' | 'u';
sampleIndexType?: 'i' | 'u';
componentType?: 'i' | 'u';
}
export const isBuiltinComparison = (builtin: TextureBuiltin) =>
builtin === 'textureGatherCompare' ||
builtin === 'textureSampleCompare' ||
builtin === 'textureSampleCompareLevel';
export const isBuiltinGather = (builtin: TextureBuiltin | undefined) =>
builtin === 'textureGather' || builtin === 'textureGatherCompare';
export const builtinNeedsSampler = (builtin: TextureBuiltin) =>
builtin.startsWith('textureSample') || builtin.startsWith('textureGather');
export const builtinNeedsDerivatives = (builtin: TextureBuiltin) =>
builtin === 'textureSample' ||
builtin === 'textureSampleBias' ||
builtin === 'textureSampleCompare';
const isSingleChannelInput = (textureType: string) => textureType.startsWith('texture_depth');
const isCubeViewDimension = (viewDescriptor?: GPUTextureViewDescriptor) =>
viewDescriptor?.dimension === 'cube' || viewDescriptor?.dimension === 'cube-array';
const isViewDimensionCubeOrCubeArray = (viewDimension: GPUTextureViewDimension) =>
viewDimension === 'cube' || viewDimension === 'cube-array';
const s_u32 = new Uint32Array(1);
const s_f32 = new Float32Array(s_u32.buffer);
const s_i32 = new Int32Array(s_u32.buffer);
const kBitCastFunctions = {
f: (v: number) => {
s_f32[0] = v;
return s_u32[0];
},
i: (v: number) => {
s_i32[0] = v;
assert(s_i32[0] === v, 'check we are not casting non-int or out-of-range value');
return s_u32[0];
},
u: (v: number) => {
s_u32[0] = v;
assert(s_u32[0] === v, 'check we are not casting non-uint or out-of-range value');
return s_u32[0];
},
};
function getCallArgType<T extends Dimensionality>(
call: TextureCall<T>,
argName: (typeof kTextureCallArgNames)[number]
) {
switch (argName) {
case 'coords':
case 'derivativeMult':
return call.coordType;
case 'component':
assert(call.componentType !== undefined);
return call.componentType;
case 'mipLevel':
assert(call.levelType !== undefined);
return call.levelType;
case 'arrayIndex':
assert(call.arrayIndexType !== undefined);
return call.arrayIndexType;
case 'sampleIndex':
assert(call.sampleIndexType !== undefined);
return call.sampleIndexType;
case 'bias':
case 'depthRef':
case 'ddx':
case 'ddy':
return 'f';
default:
unreachable();
}
}
function toArray(coords: Dimensionality): number[] {
if (coords instanceof Array) {
return coords;
}
return [coords];
}
function quantize(texel: PerTexelComponent<number>, repl: TexelRepresentationInfo) {
return repl.bitsToNumber(repl.unpackBits(new Uint8Array(repl.pack(repl.encode(texel)))));
}
function apply(a: number[], b: number[], op: (x: number, y: number) => number) {
assert(a.length === b.length, `apply(${a}, ${b}): arrays must have same length`);
return a.map((v, i) => op(v, b[i]));
}
* At the corner of a cubemap we need to sample just 3 texels, not 4.
* The texels are in
*
* 0: (u,v)
* 1: (u + 1, v)
* 2: (u, v + 1)
* 3: (u + 1, v + 1)
*
* We pass in the original 2d (converted from cubemap) texture coordinate.
* If it's within half a pixel of the edge in both directions then it's
* a corner so we return the index of the one texel that's not needed.
* Otherwise we return -1.
*/
function getUnusedCubeCornerSampleIndex(textureSize: number, coords: vec3) {
const u = coords[0] * textureSize;
const v = coords[1] * textureSize;
if (v < 0.5) {
if (u < 0.5) {
return 0;
} else if (u >= textureSize - 0.5) {
return 1;
}
} else if (v >= textureSize - 0.5) {
if (u < 0.5) {
return 2;
} else if (u >= textureSize - 0.5) {
return 3;
}
}
return -1;
}
const add = (a: number[], b: number[]) => apply(a, b, (x, y) => x + y);
* The data needed by the software rendered to simulate a texture.
* In particular, it needs texels (the data), it needs a descriptor
* for the size, format, and dimension, and it needs a view descriptor
* for the viewDimension, baseMipLevel, mipLevelCount, baseArrayLayer,
* and arrayLayerCount.
*/
export interface SoftwareTexture {
texels: TexelView[];
descriptor: GPUTextureDescriptor;
viewDescriptor: GPUTextureViewDescriptor;
}
* Converts the src texel representation to an RGBA representation.
*/
export function convertPerTexelComponentToResultFormat(
src: PerTexelComponent<number>,
format: GPUTextureFormat,
aspect: GPUTextureAspect = 'all'
): PerTexelComponent<number> {
const effectiveFormat = aspect === 'stencil-only' ? 'stencil8' : format;
const components = isEncodableTextureFormat(effectiveFormat)
? kTexelRepresentationInfo[effectiveFormat as EncodableTextureFormat].componentOrder
: kRGBAComponents;
const out: PerTexelComponent<number> = { R: 0, G: 0, B: 0, A: 1 };
for (const component of components) {
switch (component) {
case 'Stencil':
case 'Depth':
out.R = src[component];
break;
default:
assert(out[component] !== undefined);
out[component] = src[component];
}
}
return out;
}
function swizzleComponentToTexelComponent(
src: PerTexelComponent<number>,
component: GPUComponentSwizzle
): number {
switch (component) {
case '0':
return 0;
case '1':
return 1;
case 'r':
return src.R!;
case 'g':
return src.G!;
case 'b':
return src.B!;
case 'a':
return src.A!;
}
}
export function swizzleTexel(
src: PerTexelComponent<number>,
swizzle: GPUTextureComponentSwizzle | undefined
): PerTexelComponent<number> {
swizzle = swizzle ?? 'rgba';
return {
R: swizzleComponentToTexelComponent(src, swizzle[0] as GPUComponentSwizzle),
G: swizzleComponentToTexelComponent(src, swizzle[1] as GPUComponentSwizzle),
B: swizzleComponentToTexelComponent(src, swizzle[2] as GPUComponentSwizzle),
A: swizzleComponentToTexelComponent(src, swizzle[3] as GPUComponentSwizzle),
};
}
* Convert RGBA result format to texel view format.
* Example, converts
* { R: 0.1, G: 0, B: 0, A: 1 } to { Depth: 0.1 }
* { R: 0.1 } to { R: 0.1, G: 0, B: 0, A: 1 }
*/
function convertToTexelViewFormat(src: PerTexelComponent<number>, format: GPUTextureFormat) {
const componentOrder = isDepthTextureFormat(format)
? [TexelComponent.Depth]
: isStencilTextureFormat(format)
? [TexelComponent.Stencil]
: [TexelComponent.R, TexelComponent.G, TexelComponent.B, TexelComponent.A];
const out: PerTexelComponent<number> = {};
for (const component of componentOrder) {
let v = src[component];
if (v === undefined) {
if (component === 'Depth' || component === 'Stencil') {
v = src.R;
} else if (component === 'G' || component === 'B') {
v = 0;
} else {
v = 1;
}
}
out[component] = v;
}
return out;
}
* Convert RGBA result format to texel view format of src texture.
* Effectively this converts something like { R: 0.1, G: 0, B: 0, A: 1 }
* to { Depth: 0.1 }
*/
function convertResultFormatToTexelViewFormat(
src: PerTexelComponent<number>,
format: EncodableTextureFormat
): PerTexelComponent<number> {
const rep = kTexelRepresentationInfo[format];
const out: PerTexelComponent<number> = {};
for (const component of rep.componentOrder) {
out[component] = src[component] ?? src.R;
}
return out;
}
function zeroValuePerTexelComponent(components: TexelComponent[]) {
const out: PerTexelComponent<number> = {};
for (const component of components) {
out[component] = 0;
}
return out;
}
const kSamplerFns: Record<GPUCompareFunction, (ref: number, v: number) => boolean> = {
never: (ref: number, v: number) => false,
less: (ref: number, v: number) => ref < v,
equal: (ref: number, v: number) => ref === v,
'less-equal': (ref: number, v: number) => ref <= v,
greater: (ref: number, v: number) => ref > v,
'not-equal': (ref: number, v: number) => ref !== v,
'greater-equal': (ref: number, v: number) => ref >= v,
always: (ref: number, v: number) => true,
} as const;
const kDefaultValueForDepthTextureComponents: Record<TexelComponent, number> = {
R: 0,
G: 0,
B: 0,
A: 1,
Depth: 0,
Stencil: 0,
} as const;
* Applies a comparison function the R or Depth component of a texel.
*/
export function applyCompareToTexel(
components: TexelComponent[],
src: PerTexelComponent<number>,
compare: GPUCompareFunction,
ref: number
): PerTexelComponent<number> {
const out: PerTexelComponent<number> = {};
const compareFn = kSamplerFns[compare];
for (const component of components) {
out[component] =
component === 'R' || component === 'Depth'
? compareFn(ref, src[component]!)
? 1
: 0
: kDefaultValueForDepthTextureComponents[component];
}
return out;
}
function applyCompare<T extends Dimensionality>(
call: TextureCall<T>,
sampler: GPUSamplerDescriptor | undefined,
components: TexelComponent[],
src: PerTexelComponent<number>
): PerTexelComponent<number> {
if (isBuiltinComparison(call.builtin)) {
assert(sampler !== undefined);
assert(call.depthRef !== undefined);
return applyCompareToTexel(components, src, sampler.compare!, call.depthRef);
} else {
return src;
}
}
function getEffectiveLodClamp(
builtin: TextureBuiltin,
sampler: GPUSamplerDescriptor | undefined,
softwareTexture: SoftwareTexture
) {
const { mipLevelCount } = getBaseMipLevelInfo(softwareTexture);
const lodMinClamp =
isBuiltinGather(builtin) || sampler?.lodMinClamp === undefined ? 0 : sampler.lodMinClamp;
const lodMaxClamp =
isBuiltinGather(builtin) || sampler?.lodMaxClamp === undefined
? mipLevelCount - 1
: sampler.lodMaxClamp;
assert(lodMinClamp >= 0 && lodMinClamp < mipLevelCount, 'lodMinClamp in range');
assert(lodMaxClamp >= 0 && lodMaxClamp < mipLevelCount, 'lodMaxClamp in range');
assert(lodMinClamp <= lodMinClamp, 'lodMinClamp <= lodMaxClamp');
return { min: lodMinClamp, max: lodMaxClamp };
}
* Returns the expect value for a WGSL builtin texture function for a single
* mip level
*/
function softwareTextureReadMipLevel<T extends Dimensionality>(
t: GPUTest,
call: TextureCall<T>,
softwareTexture: SoftwareTexture,
sampler: GPUSamplerDescriptor | undefined,
mipLevel: number
): PerTexelComponent<number> {
assert(mipLevel % 1 === 0);
const { format } = softwareTexture.texels[0];
const rep = kTexelRepresentationInfo[format];
const { baseMipLevel, baseMipLevelSize, baseArrayLayer, arrayLayerCount } =
getBaseMipLevelInfo(softwareTexture);
const mipLevelSize = virtualMipSize(
softwareTexture.descriptor.dimension || '2d',
baseMipLevelSize,
mipLevel
);
const swizzle = softwareTexture.viewDescriptor.swizzle;
const addressMode: GPUAddressMode[] =
call.builtin === 'textureSampleBaseClampToEdge'
? ['clamp-to-edge', 'clamp-to-edge', 'clamp-to-edge']
: [
sampler?.addressModeU ?? 'clamp-to-edge',
sampler?.addressModeV ?? 'clamp-to-edge',
sampler?.addressModeW ?? 'clamp-to-edge',
];
const isCube = isCubeViewDimension(softwareTexture.viewDescriptor);
const arrayIndexMult = isCube ? 6 : 1;
const numLayers = arrayLayerCount / arrayIndexMult;
assert(numLayers % 1 === 0);
const textureSizeForCube = [mipLevelSize[0], mipLevelSize[1], 6];
const load = (at: number[]) => {
const zFromArrayIndex =
call.arrayIndex !== undefined
? clamp(call.arrayIndex, { min: 0, max: numLayers - 1 }) * arrayIndexMult
: 0;
return softwareTexture.texels[mipLevel + baseMipLevel].color({
x: Math.floor(at[0]),
y: Math.floor(at[1] ?? 0),
z: Math.floor(at[2] ?? 0) + zFromArrayIndex + baseArrayLayer,
sampleIndex: call.sampleIndex,
});
};
switch (call.builtin) {
case 'textureGather':
case 'textureGatherCompare':
case 'textureSample':
case 'textureSampleBias':
case 'textureSampleBaseClampToEdge':
case 'textureSampleCompare':
case 'textureSampleCompareLevel':
case 'textureSampleGrad':
case 'textureSampleLevel': {
let coords = toArray(call.coords!);
if (isCube) {
coords = convertCubeCoordToNormalized3DTextureCoord(coords as vec3);
}
let at = coords.map((v, i) => v * (isCube ? textureSizeForCube : mipLevelSize)[i] - 0.5);
if (call.offset !== undefined) {
at = add(at, toArray(call.offset));
}
const samples: { at: number[]; weight: number }[] = [];
const filter = isBuiltinGather(call.builtin) ? 'linear' : sampler?.minFilter ?? 'nearest';
switch (filter) {
case 'linear': {
const p0 = at.map(v => Math.floor(v));
const p1 = p0.map((v, i) => v + (isCube ? (i === 2 ? 0 : 1) : 1));
const p1W = at.map((v, i) => v - p0[i]);
const p0W = p1W.map(v => 1 - v);
switch (coords.length) {
case 1:
samples.push({ at: p0, weight: p0W[0] });
samples.push({ at: p1, weight: p1W[0] });
break;
case 2: {
samples.push({ at: [p0[0], p1[1]], weight: p0W[0] * p1W[1] });
samples.push({ at: p1, weight: p1W[0] * p1W[1] });
samples.push({ at: [p1[0], p0[1]], weight: p1W[0] * p0W[1] });
samples.push({ at: p0, weight: p0W[0] * p0W[1] });
break;
}
case 3: {
if (isCube) {
samples.push({ at: [p0[0], p1[1], p0[2]], weight: p0W[0] * p1W[1] });
samples.push({ at: p1, weight: p1W[0] * p1W[1] });
samples.push({ at: [p1[0], p0[1], p0[2]], weight: p1W[0] * p0W[1] });
samples.push({ at: p0, weight: p0W[0] * p0W[1] });
const ndx = getUnusedCubeCornerSampleIndex(mipLevelSize[0], coords as vec3);
if (ndx >= 0) {
unreachable(
`corners of cubemaps are not testable:\n ${describeTextureCall(call)}`
);
}
} else {
const p = [p0, p1];
const w = [p0W, p1W];
for (let z = 0; z < 2; ++z) {
for (let y = 0; y < 2; ++y) {
for (let x = 0; x < 2; ++x) {
samples.push({
at: [p[x][0], p[y][1], p[z][2]],
weight: w[x][0] * w[y][1] * w[z][2],
});
}
}
}
}
break;
}
}
break;
}
case 'nearest': {
const p = at.map(v => Math.round(quantizeToF32(v)));
samples.push({ at: p, weight: 1 });
break;
}
default:
unreachable();
}
if (isBuiltinGather(call.builtin)) {
const componentNdx = call.component ?? 0;
assert(componentNdx >= 0 && componentNdx < 4);
assert(samples.length === 4);
const component = kRGBAComponents[componentNdx];
const out: PerTexelComponent<number> = {};
samples.forEach((sample, i) => {
const c = isCube
? wrapFaceCoordToCubeFaceAtEdgeBoundaries(mipLevelSize[0], sample.at as vec3)
: applyAddressModesToCoords(addressMode, mipLevelSize, sample.at);
const v = load(c);
const postV = applyCompare(call, sampler, rep.componentOrder, v);
const rgba = convertPerTexelComponentToResultFormat(postV, format);
const swizzled = swizzleTexel(rgba, swizzle);
out[kRGBAComponents[i]] = swizzled[component];
});
return out;
}
const out: PerTexelComponent<number> = {};
for (const sample of samples) {
const c = isCube
? wrapFaceCoordToCubeFaceAtEdgeBoundaries(mipLevelSize[0], sample.at as vec3)
: applyAddressModesToCoords(addressMode, mipLevelSize, sample.at);
const v = load(c);
const postV = applyCompare(call, sampler, rep.componentOrder, v);
for (const component of rep.componentOrder) {
out[component] = (out[component] ?? 0) + postV[component]! * sample.weight;
}
}
const rgba = convertPerTexelComponentToResultFormat(out, format);
return swizzleTexel(rgba, swizzle);
}
case 'textureLoad': {
const out: PerTexelComponent<number> = isOutOfBoundsCall(softwareTexture, call)
? zeroValuePerTexelComponent(rep.componentOrder)
: load(call.coords!);
const rgba = convertPerTexelComponentToResultFormat(out, format);
return swizzleTexel(rgba, swizzle);
}
default:
unreachable();
}
}
* Reads a texture, optionally sampling between 2 mipLevels
*/
function softwareTextureReadLevel<T extends Dimensionality>(
t: GPUTest,
stage: ShaderStage,
call: TextureCall<T>,
softwareTexture: SoftwareTexture,
sampler: GPUSamplerDescriptor | undefined,
mipLevel: number
): PerTexelComponent<number> {
if (!sampler) {
return softwareTextureReadMipLevel<T>(t, call, softwareTexture, sampler, mipLevel);
}
const { mipLevelCount } = getBaseMipLevelInfo(softwareTexture);
const lodClampMinMax = getEffectiveLodClamp(call.builtin, sampler, softwareTexture);
const effectiveMipmapFilter = isBuiltinGather(call.builtin) ? 'nearest' : sampler.mipmapFilter;
switch (effectiveMipmapFilter) {
case 'linear': {
const clampedMipLevel = clamp(mipLevel, lodClampMinMax);
const rootMipLevel = Math.floor(clampedMipLevel);
const nextMipLevel = Math.ceil(clampedMipLevel);
const t0 = softwareTextureReadMipLevel<T>(t, call, softwareTexture, sampler, rootMipLevel);
const t1 = softwareTextureReadMipLevel<T>(t, call, softwareTexture, sampler, nextMipLevel);
const weightType = call.builtin === 'textureSampleLevel' ? 'sampleLevelWeights' : 'identity';
const mix = getWeightForMipLevel(t, stage, weightType, mipLevelCount, clampedMipLevel);
assert(mix >= 0 && mix <= 1);
const values = [
{ v: t0, weight: 1 - mix },
{ v: t1, weight: mix },
];
const out: PerTexelComponent<number> = {};
for (const { v, weight } of values) {
for (const component of kRGBAComponents) {
out[component] = (out[component] ?? 0) + v[component]! * weight;
}
}
return out;
}
default: {
const baseMipLevel = Math.floor(clamp(mipLevel, lodClampMinMax) + 0.5);
return softwareTextureReadMipLevel<T>(t, call, softwareTexture, sampler, baseMipLevel);
}
}
}
function computeMipLevelFromGradients(
ddx: readonly number[],
ddy: readonly number[],
baseMipLevelSize: GPUExtent3D
) {
const texSize = reifyExtent3D(baseMipLevelSize);
const textureSize = [texSize.width, texSize.height, texSize.depthOrArrayLayers];
const scaledDdx = ddx.map((v, i) => v * textureSize[i]);
const scaledDdy = ddy.map((v, i) => v * textureSize[i]);
const dotDDX = dotProduct(scaledDdx, scaledDdx);
const dotDDY = dotProduct(scaledDdy, scaledDdy);
const deltaMax = Math.max(dotDDX, dotDDY);
const mipLevel = 0.5 * Math.log2(deltaMax);
return mipLevel;
}
function computeMipLevelFromGradientsForCall<T extends Dimensionality>(
call: TextureCall<T>,
baseMipLevelSize: GPUExtent3D
) {
assert(!!call.ddx);
assert(!!call.ddy);
const ddx: readonly number[] = typeof call.ddx === 'number' ? [call.ddx] : call.ddx;
const ddy: readonly number[] = typeof call.ddy === 'number' ? [call.ddy] : call.ddy;
return computeMipLevelFromGradients(ddx, ddy, baseMipLevelSize);
}
* The software version of textureSampleGrad except with optional level.
*/
function softwareTextureReadGrad<T extends Dimensionality>(
t: GPUTest,
stage: ShaderStage,
call: TextureCall<T>,
softwareTexture: SoftwareTexture,
sampler?: GPUSamplerDescriptor
): PerTexelComponent<number> {
const bias = call.bias === undefined ? 0 : clamp(call.bias, { min: -16.0, max: 15.99 });
if (call.ddx) {
const { mipLevelCount, baseMipLevelSize } = getBaseMipLevelInfo(softwareTexture);
const mipLevel = computeMipLevelFromGradientsForCall(call, baseMipLevelSize);
const clampedMipLevel = clamp(mipLevel + bias, { min: 0, max: mipLevelCount - 1 });
const weightMipLevel = mapSoftwareMipLevelToGPUMipLevel(t, stage, clampedMipLevel);
return softwareTextureReadLevel(t, stage, call, softwareTexture, sampler, weightMipLevel);
} else {
return softwareTextureReadLevel(
t,
stage,
call,
softwareTexture,
sampler,
(call.mipLevel ?? 0) + bias
);
}
}
* This must match the code in doTextureCalls for derivativeBase
*
* Note: normal implicit derivatives are computed like this
*
* fn textureSample(T, S, coord) -> vec4f {
* return textureSampleGrad(T, S, dpdx(coord), dpdy(coord));
* }
*
* dpdx and dpdy are effectively computed by,
* getting the values of coord for 2x2 adjacent texels.
*
* p0 = coord value at x, y
* p1 = coord value at x + 1, y
* p2 = coord value at x, y + 1
* p3 = coord value at x + 1, y + 1
*
* dpdx is the average delta in x and dpdy is the average delta in y
*
* dpdx = (p1 - p0 + p3 - p2) / 2 // average of horizontal change
* dpdy = (p2 - p0 + p3 - p1) / 2 // average of vertical change
*
* derivativeBase is
*
* '1d' '2d' '3d'
* p0 = [0] [0, 0] [0, 0, 0]
* p1 = [1] [1, 0] [1, 0, 0]
* p2 = [0] [0, 1] [0, 1, 0]
* p3 = [1] [1, 1] [1, 1, 0]
*
* But, these values are normalized texels coords so if the src texture
* is 8x8 these would be * 0.125
*
* Note: to test other derivatives we add in a multiplier but,
* this base gives us something to add that starts at 0,0 at the call
* but who's derivatives we can easily set. We need the default
* derivativeBase to be 1 otherwise it's 0 which makes the computed mip level
* be -Infinity which means bias in `textureSampleBias` has no meaning.
*/
function derivativeBaseForCall<T extends Dimensionality>(
softwareTexture: SoftwareTexture,
isDDX: boolean
) {
const { baseMipLevelSize } = getBaseMipLevelInfo(softwareTexture);
if (isCubeViewDimension(softwareTexture.viewDescriptor)) {
return (isDDX ? [1 / baseMipLevelSize[0], 0, 1] : [0, 1 / baseMipLevelSize[1], 1]) as T;
} else if (softwareTexture.descriptor.dimension === '3d') {
return (isDDX ? [1 / baseMipLevelSize[0], 0, 0] : [0, 1 / baseMipLevelSize[1], 0]) as T;
} else if (softwareTexture.descriptor.dimension === '1d') {
return [1 / baseMipLevelSize[0]] as T;
} else {
return (isDDX ? [1 / baseMipLevelSize[0], 0] : [0, 1 / baseMipLevelSize[1]]) as T;
}
}
* Multiplies derivativeBase by derivativeMult or 1
*/
function derivativeForCall<T extends Dimensionality>(
softwareTexture: SoftwareTexture,
call: TextureCall<T>,
isDDX: boolean
) {
const dd = derivativeBaseForCall(softwareTexture, isDDX);
return dd.map((v, i) => v * (call.derivativeMult?.[i] ?? 1)) as T;
}
function softwareTextureRead<T extends Dimensionality>(
t: GPUTest,
stage: ShaderStage,
call: TextureCall<T>,
softwareTexture: SoftwareTexture,
sampler?: GPUSamplerDescriptor
): PerTexelComponent<number> {
if (builtinNeedsDerivatives(call.builtin) && !call.ddx) {
const newCall: TextureCall<T> = {
...call,
ddx: call.ddx ?? derivativeForCall<T>(softwareTexture, call, true),
ddy: call.ddy ?? derivativeForCall<T>(softwareTexture, call, false),
};
call = newCall;
}
return softwareTextureReadGrad(t, stage, call, softwareTexture, sampler);
}
export type TextureTestOptions<T extends Dimensionality> = {
ddx?: number;
ddy?: number;
uvwStart?: Readonly<T>;
offset?: Readonly<T>;
depthTexture?: boolean;
arrayIndexType?: 'i' | 'u';
};
* out of bounds is defined as any of the following being true
*
* * coords is outside the range [0, textureDimensions(t, level))
* * array_index is outside the range [0, textureNumLayers(t))
* * level is outside the range [0, textureNumLevels(t))
* * sample_index is outside the range [0, textureNumSamples(s))
*/
function isOutOfBoundsCall<T extends Dimensionality>(
softwareTexture: SoftwareTexture,
call: TextureCall<T>
) {
assert(call.coords !== undefined);
const desc = reifyTextureDescriptor(softwareTexture.descriptor);
const { coords, mipLevel: callMipLevel, arrayIndex, sampleIndex } = call;
const { baseMipLevelSize, mipLevelCount, arrayLayerCount } = getBaseMipLevelInfo(softwareTexture);
if (callMipLevel !== undefined && (callMipLevel < 0 || callMipLevel >= mipLevelCount)) {
return true;
}
const size = virtualMipSize(
softwareTexture.descriptor.dimension || '2d',
baseMipLevelSize,
callMipLevel ?? 0
);
for (let i = 0; i < coords.length; ++i) {
const v = coords[i];
if (v < 0 || v >= size[i]) {
return true;
}
}
if (arrayIndex !== undefined) {
if (arrayIndex < 0 || arrayIndex >= arrayLayerCount) {
return true;
}
}
if (sampleIndex !== undefined) {
if (sampleIndex < 0 || sampleIndex >= desc.sampleCount) {
return true;
}
}
return false;
}
function isValidOutOfBoundsValue(
device: GPUDevice,
softwareTexture: SoftwareTexture,
builtin: TextureBuiltin,
textureType: string,
gotRGBA: PerTexelComponent<number>,
maxFractionalDiff: number
) {
if (softwareTexture.descriptor.format.includes('depth')) {
if (gotRGBA.R === 0) {
return true;
}
} else {
if (
gotRGBA.R === 0 &&
gotRGBA.B === 0 &&
gotRGBA.G === 0 &&
(gotRGBA.A === 0 || gotRGBA.A === 1)
) {
return true;
}
}
const swizzle = softwareTexture.viewDescriptor.swizzle;
for (let mipLevel = 0; mipLevel < softwareTexture.texels.length; ++mipLevel) {
const mipTexels = softwareTexture.texels[mipLevel];
const size = virtualMipSize(
softwareTexture.descriptor.dimension || '2d',
softwareTexture.descriptor.size,
mipLevel
);
const sampleCount = softwareTexture.descriptor.sampleCount ?? 1;
for (let z = 0; z < size[2]; ++z) {
for (let y = 0; y < size[1]; ++y) {
for (let x = 0; x < size[0]; ++x) {
for (let sampleIndex = 0; sampleIndex < sampleCount; ++sampleIndex) {
const texel = mipTexels.color({ x, y, z, sampleIndex });
const rgba = swizzleTexel(
convertPerTexelComponentToResultFormat(texel, mipTexels.format),
swizzle
);
if (
texelsApproximatelyEqual(
device,
builtin,
textureType,
gotRGBA,
softwareTexture.descriptor.format,
rgba,
mipTexels.format,
maxFractionalDiff
)
) {
return true;
}
}
}
}
}
}
return false;
}
* For a texture builtin with no sampler (eg textureLoad),
* any out of bounds access is allowed to return one of:
*
* * the value of any texel in the texture
* * 0,0,0,0 or 0,0,0,1 if not a depth texture
* * 0 if a depth texture
*/
function okBecauseOutOfBounds<T extends Dimensionality>(
device: GPUDevice,
softwareTexture: SoftwareTexture,
call: TextureCall<T>,
textureType: string,
gotRGBA: PerTexelComponent<number>,
maxFractionalDiff: number
) {
if (!isOutOfBoundsCall(softwareTexture, call)) {
return false;
}
return isValidOutOfBoundsValue(
device,
softwareTexture,
call.builtin,
textureType,
gotRGBA,
maxFractionalDiff
);
}
const kRGBAComponents = [
TexelComponent.R,
TexelComponent.G,
TexelComponent.B,
TexelComponent.A,
] as const;
const kRComponent = [TexelComponent.R] as const;
* Compares two Texels
*
* Note: device is needed because if texture-component-swizzle
* is enabled then depth and stencil textures are required to
* return specific values for g, b, and a where is if it's not
* enabled then they are implementation defined.
*/
export function texelsApproximatelyEqual(
device: GPUDevice,
builtin: TextureBuiltin,
textureType: string,
gotRGBA: PerTexelComponent<number>,
gotFormat: GPUTextureFormat,
expectRGBA: PerTexelComponent<number>,
expectedFormat: EncodableTextureFormat,
maxFractionalDiff: number
) {
const rep = kTexelRepresentationInfo[expectedFormat];
const got = convertResultFormatToTexelViewFormat(gotRGBA, expectedFormat);
const expect = convertResultFormatToTexelViewFormat(expectRGBA, expectedFormat);
const gULP = convertPerTexelComponentToResultFormat(
rep.bitsToULPFromZero(rep.numberToBits(got)),
expectedFormat
);
const eULP = convertPerTexelComponentToResultFormat(
rep.bitsToULPFromZero(rep.numberToBits(expect)),
expectedFormat
);
const rgbaComponentsToCheck = getComponentsToCheck(device, gotFormat, builtin, textureType);
for (const component of rgbaComponentsToCheck) {
const g = gotRGBA[component]!;
const e = expectRGBA[component]!;
assert(!isNaN(g), () => `got component is NaN: ${g}`);
assert(!isNaN(e), () => `expected component is NaN: ${e}`);
const absDiff = Math.abs(g - e);
const ulpDiff = Math.abs(gULP[component]! - eULP[component]!);
if (ulpDiff > 3 && absDiff > maxFractionalDiff) {
return false;
}
}
return true;
}
function getULPFromZeroForComponents(
rgba: PerTexelComponent<number>,
format: EncodableTextureFormat,
builtin: TextureBuiltin,
componentNdx?: number
): PerTexelComponent<number> {
const rep = kTexelRepresentationInfo[format];
if (isBuiltinGather(builtin)) {
const out: PerTexelComponent<number> = {};
const component = kRGBAComponents[componentNdx ?? 0];
const temp: PerTexelComponent<number> = { R: 0, G: 0, B: 0, A: 1 };
for (const comp of kRGBAComponents) {
temp[component] = rgba[comp];
const texel = convertResultFormatToTexelViewFormat(temp, format);
const ulp = convertPerTexelComponentToResultFormat(
rep.bitsToULPFromZero(rep.numberToBits(texel)),
format
);
out[comp] = ulp[component];
}
return out;
} else {
const texel = convertResultFormatToTexelViewFormat(rgba, format);
return convertPerTexelComponentToResultFormat(
rep.bitsToULPFromZero(rep.numberToBits(texel)),
format
);
}
}
function getTextureViewDescription(softwareTexture: SoftwareTexture) {
const size = reifyExtent3D(softwareTexture.descriptor.size);
const { baseMipLevel, mipLevelCount, baseArrayLayer, arrayLayerCount, baseMipLevelSize } =
getBaseMipLevelInfo(softwareTexture);
const physicalMipLevelCount = softwareTexture.descriptor.mipLevelCount ?? 1;
return `
physical size: [${size.width}, ${size.height}, ${size.depthOrArrayLayers}]
baseMipLevel: ${baseMipLevel}
mipLevelCount: ${mipLevelCount}
baseMipLevelSize: [${baseMipLevelSize.join(', ')}]
baseArrayLayer: ${baseArrayLayer}
arrayLayerCount: ${arrayLayerCount}
physicalMipCount: ${physicalMipLevelCount}
`;
}
function getComponentsToCheck(
device: GPUDevice,
format: GPUTextureFormat,
builtin: TextureBuiltin,
textureType: string
) {
const returnsOneComponent = !isBuiltinGather(builtin) && isSingleChannelInput(textureType);
if (returnsOneComponent) {
return kRComponent;
}
const gbaUndefined =
isDepthOrStencilTextureFormat(format) &&
!hasFeature(device.features, 'texture-component-swizzle');
if (gbaUndefined) {
return kRComponent;
}
return kRGBAComponents;
}
* Checks the result of each call matches the expected result.
*/
export async function checkCallResults<T extends Dimensionality>(
t: GPUTest,
softwareTexture: SoftwareTexture,
textureType: string,
sampler: GPUSamplerDescriptor | undefined,
calls: TextureCall<T>[],
results: Awaited<ReturnType<typeof doTextureCalls<T>>>,
shortShaderStage: ShortShaderStage,
gpuTexture?: GPUTexture
) {
const stage = kShortShaderStageToShaderStage[shortShaderStage];
if (builtinNeedsMipLevelWeights(calls[0].builtin)) {
await initMipLevelWeightsForDevice(t, stage);
}
let haveComparisonCheckInfo = false;
let checkInfo = {
runner: results.runner,
calls,
sampler,
};
let gpuTexels: TexelView[] | undefined;
const errs: string[] = [];
const format = softwareTexture.texels[0].format;
const size = reifyExtent3D(softwareTexture.descriptor.size);
const maxFractionalDiff =
sampler?.minFilter === 'linear' ||
sampler?.magFilter === 'linear' ||
sampler?.mipmapFilter === 'linear'
? getMaxFractionalDiffForTextureFormat(softwareTexture.descriptor.format)
: 0;
t.debug(() => getTextureViewDescription(softwareTexture));
for (let callIdx = 0; callIdx < calls.length; callIdx++) {
const call = calls[callIdx];
t.debug(`#${callIdx}: ${describeTextureCall(call)}`);
const gotRGBA = results.results[callIdx];
const expectRGBA = softwareTextureRead(t, stage, call, softwareTexture, sampler);
const callSpecificMaxFractionalDiff =
call.bias! >= 12 ? maxFractionalDiff * (2 + call.bias! - 12) : maxFractionalDiff;
if (
isDepthOrStencilTextureFormat(format) &&
isBuiltinGather(call.builtin) &&
call.component! > 0 &&
!hasFeature(t.device.features, 'texture-component-swizzle')
) {
continue;
}
if (
texelsApproximatelyEqual(
t.device,
call.builtin,
textureType,
gotRGBA,
softwareTexture.descriptor.format,
expectRGBA,
format,
callSpecificMaxFractionalDiff
)
) {
continue;
}
if (
!sampler &&
okBecauseOutOfBounds(
t.device,
softwareTexture,
call,
textureType,
gotRGBA,
callSpecificMaxFractionalDiff
)
) {
continue;
}
const gULP = getULPFromZeroForComponents(gotRGBA, format, call.builtin, call.component);
const eULP = getULPFromZeroForComponents(expectRGBA, format, call.builtin, call.component);
const rgbaComponentsToCheck = getComponentsToCheck(t.device, format, call.builtin, textureType);
let bad = false;
const diffs = rgbaComponentsToCheck.map(component => {
const g = gotRGBA[component]!;
const e = expectRGBA[component]!;
const absDiff = Math.abs(g - e);
const ulpDiff = Math.abs(gULP[component]! - eULP[component]!);
assert(!Number.isNaN(ulpDiff));
const maxAbs = Math.max(Math.abs(g), Math.abs(e));
const relDiff = maxAbs > 0 ? absDiff / maxAbs : 0;
if (ulpDiff > 3 && absDiff > callSpecificMaxFractionalDiff) {
bad = true;
}
return { absDiff, relDiff, ulpDiff };
});
const isFloatType = (format: GPUTextureFormat) => {
const type = getTextureFormatType(format);
return type === 'float' || type === 'depth';
};
const fix5 = (n: number) => (isFloatType(format) ? n.toFixed(5) : n.toString());
const fix5v = (arr: number[]) => arr.map(v => fix5(v)).join(', ');
const rgbaToArray = (p: PerTexelComponent<number>): number[] =>
rgbaComponentsToCheck.map(component => p[component]!);
if (bad) {
const { baseMipLevelSize } = getBaseMipLevelInfo(softwareTexture);
const physicalMipLevelCount = softwareTexture.descriptor.mipLevelCount ?? 1;
const lodClamp = getEffectiveLodClamp(call.builtin, sampler, softwareTexture);
const desc = describeTextureCall(call);
errs.push(`result was not as expected:${getTextureViewDescription(softwareTexture)}
lodMinClamp: ${lodClamp.min} (effective)
lodMaxClamp: ${lodClamp.max} (effective)
call: ${desc} // #${callIdx}`);
if (isCubeViewDimension(softwareTexture.viewDescriptor)) {
const coord = convertCubeCoordToNormalized3DTextureCoord(call.coords as vec3);
const faceNdx = Math.floor(coord[2] * 6);
errs.push(` : as 3D texture coord: (${coord[0]}, ${coord[1]}, ${coord[2]})`);
for (let mipLevel = 0; mipLevel < physicalMipLevelCount; ++mipLevel) {
const mipSize = virtualMipSize(
softwareTexture.descriptor.dimension ?? '2d',
softwareTexture.descriptor.size,
mipLevel
);
const t = coord.slice(0, 2).map((v, i) => (v * mipSize[i]).toFixed(3));
errs.push(
` : as texel coord mip level[${mipLevel}]: (${t[0]}, ${t[1]}), face: ${faceNdx}(${kFaceNames[faceNdx]})`
);
}
} else if (call.coordType === 'f') {
for (let mipLevel = 0; mipLevel < physicalMipLevelCount; ++mipLevel) {
const mipSize = virtualMipSize(
softwareTexture.descriptor.dimension ?? '2d',
softwareTexture.descriptor.size,
mipLevel
);
const t = call.coords!.map((v, i) => (v * mipSize[i]).toFixed(3));
errs.push(` : as texel coord @ mip level[${mipLevel}]: (${t.join(', ')})`);
}
}
if (builtinNeedsDerivatives(call.builtin)) {
const ddx = derivativeForCall<T>(softwareTexture, call, true);
const ddy = derivativeForCall<T>(softwareTexture, call, false);
const mipLevel = computeMipLevelFromGradients(ddx, ddy, baseMipLevelSize);
const biasStr = call.bias === undefined ? '' : ' (without bias)';
errs.push(`implicit derivative based mip level: ${fix5(mipLevel)}${biasStr}`);
if (call.bias) {
const clampedBias = clamp(call.bias ?? 0, { min: -16.0, max: 15.99 });
errs.push(`\
clamped bias: ${fix5(clampedBias)}
mip level with bias: ${fix5(mipLevel + clampedBias)}`);
}
} else if (call.ddx) {
const mipLevel = computeMipLevelFromGradientsForCall(call, size);
errs.push(`gradient based mip level: ${mipLevel}`);
}
errs.push(`\
got: ${fix5v(rgbaToArray(gotRGBA))}
expected: ${fix5v(rgbaToArray(expectRGBA))}
max diff: ${callSpecificMaxFractionalDiff}
abs diffs: ${fix5v(diffs.map(({ absDiff }) => absDiff))}
rel diffs: ${diffs.map(({ relDiff }) => `${(relDiff * 100).toFixed(2)}%`).join(', ')}
ulp diffs: ${diffs.map(({ ulpDiff }) => ulpDiff).join(', ')}
`);
if (sampler) {
if (t.rec.debugging) {
if (isBuiltinComparison(call.builtin)) {
if (!haveComparisonCheckInfo) {
const debugCalls = calls.map(call => {
const debugCall = { ...call };
debugCall.depthRef = undefined;
switch (call.builtin) {
case 'textureGatherCompare':
debugCall.builtin = 'textureGather';
break;
case 'textureSampleCompare':
debugCall.builtin = 'textureSample';
break;
case 'textureSampleCompareLevel':
debugCall.builtin = 'textureSampleLevel';
debugCall.levelType = 'u';
debugCall.mipLevel = 0;
break;
default:
unreachable();
}
return debugCall;
});
const debugSampler = { ...sampler };
delete debugSampler.compare;
const debugRunner = createTextureCallsRunner(
t,
{
format,
dimension: softwareTexture.descriptor.dimension ?? '2d',
sampleCount: softwareTexture.descriptor.sampleCount ?? 1,
depthOrArrayLayers: size.depthOrArrayLayers,
},
softwareTexture.viewDescriptor,
textureType,
debugSampler,
debugCalls,
stage
);
checkInfo = {
runner: debugRunner,
sampler: debugSampler,
calls: debugCalls,
};
haveComparisonCheckInfo = true;
}
}
if (!gpuTexels && gpuTexture) {
gpuTexels = await readTextureToTexelViews(
t,
gpuTexture,
softwareTexture.descriptor,
getTexelViewFormatForTextureFormat(gpuTexture.format)
);
}
const callForSamplePoints = checkInfo.calls[callIdx];
const useTexelFormatForGPUTexture = isCompressedTextureFormat(
softwareTexture.descriptor.format
);
if (useTexelFormatForGPUTexture) {
errs.push(`
### WARNING: sample points are derived from un-compressed textures and may not match the
actual GPU results of sampling a compressed texture. The test itself failed at this point
(see expected: and got: above). We're only trying to determine what the GPU sampled, but
we can not do that easily with compressed textures. ###
`);
}
const expectedSamplePoints = [
'expected:',
...(await identifySamplePoints(
softwareTexture,
sampler,
callForSamplePoints,
call,
softwareTexture.texels,
(texels: TexelView[]) => {
return Promise.resolve(
softwareTextureRead(
t,
stage,
callForSamplePoints,
{
texels,
descriptor: softwareTexture.descriptor,
viewDescriptor: softwareTexture.viewDescriptor,
},
checkInfo.sampler
)
);
}
)),
];
const gotSamplePoints = [
'got:',
...(await identifySamplePoints(
softwareTexture,
sampler,
callForSamplePoints,
call,
gpuTexels,
async (texels: TexelView[]) => {
const descriptor = { ...softwareTexture.descriptor };
if (useTexelFormatForGPUTexture) {
descriptor.format = texels[0].format;
}
const gpuTexture = createTextureFromTexelViewsLocal(t, [texels], descriptor);
const result = (await checkInfo.runner.run(gpuTexture))[callIdx];
gpuTexture.destroy();
return result;
}
)),
];
errs.push(' sample points:');
errs.push(layoutTwoColumns(expectedSamplePoints, gotSamplePoints).join('\n'));
errs.push('', '');
}
if (!t.rec.debugging) {
errs.push('### turn on debugging to see sample points ###');
}
}
break;
}
}
results.runner.destroy();
checkInfo.runner.destroy();
return errs.length > 0 ? new Error(errs.join('\n')) : undefined;
}
function getMaxFractionalDiffForTextureFormat(format: GPUTextureFormat) {
if (format.includes('depth')) {
return 3 / 100;
} else if (format.includes('8unorm')) {
return 7 / 255;
} else if (format.includes('2unorm')) {
return 13 / 512;
} else if (format.includes('unorm')) {
return 7 / 255;
} else if (format.includes('8snorm')) {
return 7.9 / 128;
} else if (format.includes('snorm')) {
return 7.9 / 128;
} else if (format.endsWith('ufloat')) {
return 156.249;
} else if (format.endsWith('float')) {
return 44;
} else {
return 0;
}
}
const sumOfCharCodesOfString = (s: unknown) =>
String(s)
.split('')
.reduce((sum, c) => sum + c.charCodeAt(0), 0);
* Makes a function that fills a block portion of a Uint8Array with random valid data
* for an astc block.
*
* The astc format is fairly complicated. For now we do the simplest thing.
* which is to set the block as a "void-extent" block (a solid color).
* This makes our test have far less precision.
*
* MAINTENANCE_TODO: generate other types of astc blocks. One option would
* be to randomly select from set of pre-made blocks.
*
* See Spec:
* https://registry.khronos.org/OpenGL/extensions/KHR/KHR_texture_compression_astc_hdr.txt
*/
function makeAstcBlockFiller(format: ColorTextureFormat) {
const { bytesPerBlock } = getBlockInfoForColorTextureFormat(format);
return (data: Uint8Array, offset: number, hashBase: number) => {
data.set(
[
0b1111_1100,
0b1111_1101,
0b1111_1111,
0b1111_1111,
0b1111_1111,
0b1111_1111,
0b1111_1111,
0b1111_1111,
],
offset
);
const end = offset + bytesPerBlock;
for (let i = offset + 8; i < end; ++i) {
data[i] = hashU32(hashBase, i);
}
};
}
* Makes a function that fills a block portion of a Uint8Array with random bytes.
*/
function makeRandomBytesBlockFiller(format: ColorTextureFormat) {
const { bytesPerBlock } = getBlockInfoForColorTextureFormat(format);
return (data: Uint8Array, offset: number, hashBase: number) => {
const end = offset + bytesPerBlock;
for (let i = offset; i < end; ++i) {
data[i] = hashU32(hashBase, i);
}
};
}
function getBlockFiller(format: ColorTextureFormat) {
if (format.startsWith('astc')) {
return makeAstcBlockFiller(format);
} else {
return makeRandomBytesBlockFiller(format);
}
}
* Fills a texture with random data.
*/
function fillTextureWithRandomData(device: GPUDevice, texture: GPUTexture) {
assert(isColorTextureFormat(texture.format));
assert(!isCompressedFloatTextureFormat(texture.format));
const info = getBlockInfoForColorTextureFormat(texture.format as ColorTextureFormat);
const hashBase =
sumOfCharCodesOfString(texture.format) +
sumOfCharCodesOfString(texture.dimension) +
texture.width +
texture.height +
texture.depthOrArrayLayers +
texture.mipLevelCount;
const bytesPerBlock = info.bytesPerBlock;
const fillBlock = getBlockFiller(texture.format as ColorTextureFormat);
for (let mipLevel = 0; mipLevel < texture.mipLevelCount; ++mipLevel) {
const size = physicalMipSizeFromTexture(texture, mipLevel);
const blocksAcross = Math.ceil(size[0] / info.blockWidth);
const blocksDown = Math.ceil(size[1] / info.blockHeight);
const bytesPerRow = blocksAcross * bytesPerBlock;
const bytesNeeded = bytesPerRow * blocksDown * size[2];
const data = new Uint8Array(bytesNeeded);
for (let offset = 0; offset < bytesNeeded; offset += bytesPerBlock) {
fillBlock(data, offset, hashBase);
}
device.queue.writeTexture(
{ texture, mipLevel },
data,
{ bytesPerRow, rowsPerImage: blocksDown },
size
);
}
}
const s_readTextureToRGBA32DeviceToPipeline = new WeakMap<
GPUDevice,
Map<string, GPUComputePipeline>
>();
function getEffectiveViewDimension(
t: GPUTest,
descriptor: Omit<GPUTextureDescriptor, 'format' | 'usage'>
): GPUTextureViewDimension {
const size = reifyExtent3D(descriptor.size);
return effectiveViewDimensionForDimension(
descriptor.textureBindingViewDimension,
descriptor.dimension,
size.depthOrArrayLayers
);
}
* Reads a texture to an array of TexelViews, one per mip level.
* format is the format of the TexelView you want. Often this is
* same as the texture.format but if the texture.format is not
* "Encodable" then you need to choose a different format.
* Example: depth24plus -> r32float, bc1-rgba-unorm to rgba32float
*/
export async function readTextureToTexelViews(
t: GPUTest,
texture: GPUTexture,
descriptor: Omit<GPUTextureDescriptor, 'format' | 'usage'>,
format: EncodableTextureFormat
) {
const device = t.device;
const viewDimensionToPipelineMap =
s_readTextureToRGBA32DeviceToPipeline.get(device) ??
new Map<GPUTextureViewDimension, GPUComputePipeline>();
s_readTextureToRGBA32DeviceToPipeline.set(device, viewDimensionToPipelineMap);
const { componentType, resultType } = getTextureFormatTypeInfo(texture.format);
const viewDimension = getEffectiveViewDimension(t, descriptor);
const id = `${texture.format}:${viewDimension}:${texture.sampleCount}`;
let pipeline = viewDimensionToPipelineMap.get(id);
if (!pipeline) {
let textureWGSL;
let loadWGSL;
let dimensionWGSL = 'textureDimensions(tex, 0)';
switch (viewDimension) {
case '2d':
if (texture.sampleCount > 1) {
textureWGSL = `texture_multisampled_2d<${componentType}>`;
loadWGSL = 'textureLoad(tex, coord.xy, sampleIndex)';
dimensionWGSL = 'textureDimensions(tex)';
} else {
textureWGSL = `texture_2d<${componentType}>`;
loadWGSL = 'textureLoad(tex, coord.xy, 0)';
}
break;
case 'cube-array':
case '2d-array':
textureWGSL = `texture_2d_array<${componentType}>`;
loadWGSL = `
textureLoad(
tex,
coord.xy,
coord.z,
0)`;
break;
case '3d':
textureWGSL = `texture_3d<${componentType}>`;
loadWGSL = 'textureLoad(tex, coord.xyz, 0)';
break;
case 'cube':
textureWGSL = `texture_cube<${componentType}>`;
loadWGSL = `
textureLoadCubeAs2DArray(tex, coord.xy, coord.z);
`;
break;
case '1d':
textureWGSL = `texture_1d<${componentType}>`;
loadWGSL = `textureLoad(tex, coord.x, 0)`;
dimensionWGSL = `vec2u(textureDimensions(tex), 1)`;
break;
default:
unreachable(`unsupported view: ${viewDimension}`);
}
const textureLoadCubeWGSL = `
const faceMat = array(
mat3x3f( 0, 0, -2, 0, -2, 0, 1, 1, 1), // pos-x
mat3x3f( 0, 0, 2, 0, -2, 0, -1, 1, -1), // neg-x
mat3x3f( 2, 0, 0, 0, 0, 2, -1, 1, -1), // pos-y
mat3x3f( 2, 0, 0, 0, 0, -2, -1, -1, 1), // neg-y
mat3x3f( 2, 0, 0, 0, -2, 0, -1, 1, 1), // pos-z
mat3x3f(-2, 0, 0, 0, -2, 0, 1, 1, -1)); // neg-z
// needed for compat mode.
fn textureLoadCubeAs2DArray(tex: texture_cube<${componentType}>, coord: vec2u, layer: u32) -> ${resultType} {
// convert texel coord normalized coord
let size = textureDimensions(tex, 0);
// Offset by 0.75 instead of the more common 0.5 for converting from texel to normalized texture coordinate
// because we're using textureGather. 0.5 would indicate the center of a texel but based on precision issues
// the "gather" could go in any direction from that center. Off center it should go in an expected direction.
let uv = (vec2f(coord) + 0.75) / vec2f(size.xy);
// convert uv + layer into cube coord
let cubeCoord = faceMat[layer] * vec3f(uv, 1.0);
// We have to use textureGather as it's the only texture builtin that works on cubemaps
// with integer texture formats.
let r = textureGather(0, tex, smp, cubeCoord);
let g = textureGather(1, tex, smp, cubeCoord);
let b = textureGather(2, tex, smp, cubeCoord);
let a = textureGather(3, tex, smp, cubeCoord);
// element 3 is the texel corresponding to cubeCoord
return ${resultType}(r[3], g[3], b[3], a[3]);
}
`;
const module = device.createShaderModule({
code: `
${isViewDimensionCubeOrCubeArray(viewDimension) ? textureLoadCubeWGSL : ''}
struct Uniforms {
sampleCount: u32,
};
@group(0) @binding(0) var<uniform> uni: Uniforms;
@group(0) @binding(1) var tex: ${textureWGSL};
@group(0) @binding(2) var smp: sampler;
@group(0) @binding(3) var<storage, read_write> data: array<${resultType}>;
@compute @workgroup_size(1) fn cs(
@builtin(global_invocation_id) global_invocation_id : vec3<u32>) {
_ = smp;
let size = ${dimensionWGSL};
let ndx = global_invocation_id.z * size.x * size.y * uni.sampleCount +
global_invocation_id.y * size.x * uni.sampleCount +
global_invocation_id.x;
let coord = vec3u(global_invocation_id.x / uni.sampleCount, global_invocation_id.yz);
let sampleIndex = global_invocation_id.x % uni.sampleCount;
data[ndx] = ${loadWGSL};
}
`,
});
const type = getTextureFormatType(texture.format);
const sampleType = isDepthTextureFormat(texture.format)
? 'unfilterable-float'
: isStencilTextureFormat(texture.format)
? 'uint'
: type === 'float'
? 'unfilterable-float'
: type;
const bindGroupLayout = device.createBindGroupLayout({
entries: [
{
binding: 0,
visibility: GPUShaderStage.COMPUTE,
buffer: {
type: 'uniform',
},
},
{
binding: 1,
visibility: GPUShaderStage.COMPUTE,
texture: {
sampleType,
viewDimension,
multisampled: texture.sampleCount > 1,
},
},
{
binding: 2,
visibility: GPUShaderStage.COMPUTE,
sampler: {
type: 'non-filtering',
},
},
{
binding: 3,
visibility: GPUShaderStage.COMPUTE,
buffer: {
type: 'storage',
},
},
],
});
const layout = device.createPipelineLayout({
bindGroupLayouts: [bindGroupLayout],
});
pipeline = device.createComputePipeline({ layout, compute: { module } });
viewDimensionToPipelineMap.set(id, pipeline);
}
const encoder = device.createCommandEncoder({ label: 'readTextureToTexelViews' });
const readBuffers = [];
for (let mipLevel = 0; mipLevel < texture.mipLevelCount; ++mipLevel) {
const size = virtualMipSize(texture.dimension, texture, mipLevel);
const uniformValues = new Uint32Array([texture.sampleCount, 0, 0, 0]);
const uniformBuffer = t.createBufferTracked({
label: 'readTextureToTexelViews:uniformBuffer',
size: uniformValues.byteLength,
usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
});
device.queue.writeBuffer(uniformBuffer, 0, uniformValues);
const storageBuffer = t.createBufferTracked({
label: 'readTextureToTexelViews:storageBuffer',
size: size[0] * size[1] * size[2] * 4 * 4 * texture.sampleCount,
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
});
const readBuffer = t.createBufferTracked({
label: 'readTextureToTexelViews:readBuffer',
size: storageBuffer.size,
usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST,
});
readBuffers.push({ size, readBuffer });
const sampler = device.createSampler();
const aspect = getAspectForTexture(texture);
const bindGroup = device.createBindGroup({
layout: pipeline.getBindGroupLayout(0),
entries: [
{ binding: 0, resource: { buffer: uniformBuffer } },
{
binding: 1,
resource: texture.createView({
dimension: viewDimension,
aspect,
baseMipLevel: mipLevel,
mipLevelCount: 1,
}),
},
{ binding: 2, resource: sampler },
{ binding: 3, resource: { buffer: storageBuffer } },
],
});
const pass = encoder.beginComputePass();
pass.setPipeline(pipeline);
pass.setBindGroup(0, bindGroup);
pass.dispatchWorkgroups(size[0] * texture.sampleCount, size[1], size[2]);
pass.end();
encoder.copyBufferToBuffer(storageBuffer, 0, readBuffer, 0, readBuffer.size);
}
device.queue.submit([encoder.finish()]);
const texelViews: TexelView[] = [];
for (const { readBuffer, size } of readBuffers) {
await readBuffer.mapAsync(GPUMapMode.READ);
const Ctor =
componentType === 'i32' ? Int32Array : componentType === 'u32' ? Uint32Array : Float32Array;
const data = new Ctor(readBuffer.getMappedRange()).slice();
readBuffer.unmap();
const { sampleCount } = texture;
texelViews.push(
TexelView.fromTexelsAsColors(format, coord => {
const offset =
((coord.z * size[0] * size[1] + coord.y * size[0] + coord.x) * sampleCount +
(coord.sampleIndex ?? 0)) *
4;
return convertResultFormatToTexelViewFormat(
{
R: data[offset + 0],
G: data[offset + 1],
B: data[offset + 2],
A: data[offset + 3],
},
format
);
})
);
}
return texelViews;
}
* Creates a texture from an array of TexelViews.
* @param t the current test
* @param texelViews Array of TexelViews per aspect per mip level. Note that only depth-stencil textures
* have 2 aspects, in which case it's assumed the texelViews are in the order [depth, stencil]
* @param desc description for the texture to be created
* @returns created texture
*/
function createTextureFromTexelViewsLocal(
t: GPUTest,
texelViews: TexelView[][],
desc: GPUTextureDescriptor
): GPUTexture {
const modifiedDescriptor = { ...desc };
if (isDepthOrStencilTextureFormat(desc.format) || desc.sampleCount! > 1) {
modifiedDescriptor.usage = desc.usage | GPUTextureUsage.RENDER_ATTACHMENT;
}
const texture = createTextureFromTexelViews(t, texelViews[0], modifiedDescriptor);
if (texelViews.length > 1) {
copyTexelViewsToTexture(t, texture, 'stencil-only', texelViews[1]);
}
return texture;
}
* Fills a texture with random data and returns that data as
* an array of arrays of TexelView. The top level arrays are
* per aspect. For textures with 1 aspect it would just be
* `[[textureViewsByMipLevel]]`. For a depth-stencil textures
* it would be [[depthViewsByMipLevel], [stencilViewsByMipLevel]]
*
* For compressed textures the texture is filled with random bytes
* and then read back from the GPU by sampling so the GPU decompressed
* the texture.
*
* For uncompressed textures the TexelViews are generated and then
* copied to the texture.
*/
export async function createTextureWithRandomDataAndGetTexelsForEachAspect(
t: GPUTest,
descriptor: GPUTextureDescriptor,
options?: RandomTextureOptions
) {
if (isCompressedTextureFormat(descriptor.format)) {
assert(!options, 'options not supported for compressed textures');
const texture = t.createTextureTracked(descriptor);
fillTextureWithRandomData(t.device, texture);
const texels = await readTextureToTexelViews(
t,
texture,
descriptor,
getTexelViewFormatForTextureFormat(texture.format)
);
return { texture, texels: [texels] };
} else if (isUnencodableDepthFormat(descriptor.format)) {
const d32Descriptor = {
...descriptor,
format: 'depth32float' as GPUTextureFormat,
};
const stencilTexels = isStencilTextureFormat(descriptor.format)
? [createRandomTexelViewMipmap({ ...descriptor, format: 'stencil8' })]
: [];
const texture = createTextureFromTexelViewsLocal(
t,
[createRandomTexelViewMipmap(d32Descriptor, options), ...stencilTexels],
descriptor
);
const texels = await readTextureToTexelViews(
t,
texture,
descriptor,
getTexelViewFormatForTextureFormat(texture.format)
);
const texelsPerAspect = [texels, ...stencilTexels];
return { texture, texels: texelsPerAspect };
} else {
const texels = createRandomTexelViewMipmap(descriptor, options);
const texture = createTextureFromTexelViewsLocal(t, [texels], descriptor);
return { texture, texels: [texels] };
}
}
* Fills a texture with random data and returns that data an arrays of TexelView.
* Note: If the texture has multiple aspects it only returns the `depth-only`
* TexelViews. If you need views for both aspects use
* @see {@link createTextureWithRandomDataAndGetTexelsForEachAspect}.
*/
export async function createTextureWithRandomDataAndGetTexels(
t: GPUTest,
descriptor: GPUTextureDescriptor,
options?: RandomTextureOptions
) {
const { texture, texels } = await createTextureWithRandomDataAndGetTexelsForEachAspect(
t,
descriptor,
options
);
return { texture, texels: texels[0] };
}
function valueIfAllComponentsAreEqual(
c: PerTexelComponent<number>,
componentOrder: readonly TexelComponent[]
) {
const s = new Set(componentOrder.map(component => c[component]!));
return s.size === 1 ? s.values().next().value : undefined;
}
* Creates a Canvas with random data and a TexelView with the same data.
*/
export function createCanvasWithRandomDataAndGetTexels(textureSize: GPUExtent3D) {
const size = reifyExtent3D(textureSize);
assert(size.depthOrArrayLayers === 1);
const imageData = new ImageData(size.width, size.height);
const data = imageData.data;
const asU32 = new Uint32Array(data.buffer);
for (let i = 0; i < asU32.length; ++i) {
asU32[i] = hashU32(i);
}
const canvas = new OffscreenCanvas(size.width, size.height);
const ctx = canvas.getContext('2d')!;
ctx.putImageData(imageData, 0, 0);
for (let i = 0; i < data.length; i += 4) {
const alpha = data[i + 3] / 255;
data[i + 0] = data[i + 0] * alpha;
data[i + 1] = data[i + 1] * alpha;
data[i + 2] = data[i + 2] * alpha;
}
const texels = [
TexelView.fromTextureDataByReference('rgba8unorm', data, {
bytesPerRow: size.width * 4,
rowsPerImage: size.height,
subrectOrigin: [0, 0, 0],
subrectSize: size,
}),
];
return { canvas, texels };
}
const kFaceNames = ['+x', '-x', '+y', '-y', '+z', '-z'] as const;
* Generates a text art grid showing which texels were sampled
* followed by a list of the samples and the weights used for each
* component.
*
* It works by making a set of indices for every texel in the texture.
* It splits the set into 2. It picks one set and generates texture data
* using TexelView.fromTexelsAsColor with [1, 1, 1, 1] texels for members
* of the current set.
*
* In then calls 'run' which renders a single `call`. `run` uses either
* the software renderer or WebGPU. It then checks the results. If the
* result is zero, all texels in the current had no influence when sampling
* and can be discarded.
*
* If the result is > 0 then, if the set has more than one member, the
* set is split and added to the list to sets to test. If the set only
* had one member then the result is the weight used when sampling that texel.
*
* This lets you see if the weights from the software renderer match the
* weights from WebGPU.
*
* Example:
*
* 0 1 2 3 4 5 6 7
* +---+---+---+---+---+---+---+---+
* 0 | | | | | | | | |
* +---+---+---+---+---+---+---+---+
* 1 | | | | | | | | a |
* +---+---+---+---+---+---+---+---+
* 2 | | | | | | | | b |
* +---+---+---+---+---+---+---+---+
* 3 | | | | | | | | |
* +---+---+---+---+---+---+---+---+
* 4 | | | | | | | | |
* +---+---+---+---+---+---+---+---+
* 5 | | | | | | | | |
* +---+---+---+---+---+---+---+---+
* 6 | | | | | | | | |
* +---+---+---+---+---+---+---+---+
* 7 | | | | | | | | |
* +---+---+---+---+---+---+---+---+
* a: at: [7, 1], weights: [R: 0.75000]
* b: at: [7, 2], weights: [R: 0.25000]
*/
async function identifySamplePoints<T extends Dimensionality>(
softwareTexture: SoftwareTexture,
sampler: GPUSamplerDescriptor,
callForSamples: TextureCall<T>,
originalCall: TextureCall<T>,
texels: TexelView[] | undefined,
run: (texels: TexelView[]) => Promise<PerTexelComponent<number>>
) {
const info = softwareTexture.descriptor;
const isCube = isCubeViewDimension(softwareTexture.viewDescriptor);
const mipLevelCount = softwareTexture.descriptor.mipLevelCount ?? 1;
const mipLevelSizes = range(mipLevelCount, mipLevel =>
virtualMipSize(
softwareTexture.descriptor.dimension ?? '2d',
softwareTexture.descriptor.size,
mipLevel
)
);
const numTexelsPerLevel = mipLevelSizes.map(size => size.reduce((s, v) => s * v));
const numTexelsOfPrecedingLevels = (() => {
let total = 0;
return numTexelsPerLevel.map(v => {
const num = total;
total += v;
return num;
});
})();
const numTexels = numTexelsPerLevel.reduce((sum, v) => sum + v);
const getMipLevelFromTexelId = (texelId: number) => {
for (let mipLevel = mipLevelCount - 1; mipLevel > 0; --mipLevel) {
if (texelId - numTexelsOfPrecedingLevels[mipLevel] >= 0) {
return mipLevel;
}
}
return 0;
};
const getTexelCoordFromTexelId = (texelId: number) => {
const mipLevel = getMipLevelFromTexelId(texelId);
const size = mipLevelSizes[mipLevel];
const texelsPerSlice = size[0] * size[1];
const id = texelId - numTexelsOfPrecedingLevels[mipLevel];
const layer = Math.floor(id / texelsPerSlice);
const xyId = id - layer * texelsPerSlice;
const y = (xyId / size[0]) | 0;
const x = xyId % size[0];
return { x, y, z: layer, mipLevel, xyId };
};
const format = (
kEncodableTextureFormats.includes(info.format as EncodableTextureFormat)
? info.format
: isDepthTextureFormat(info.format)
? 'depth16unorm'
: 'rgba8unorm'
) as EncodableTextureFormat;
const rep = kTexelRepresentationInfo[format];
const components = isBuiltinGather(callForSamples.builtin) ? kRGBAComponents : rep.componentOrder;
const convertResultAsAppropriate = isBuiltinGather(callForSamples.builtin)
? <T>(v: T) => v
: convertResultFormatToTexelViewFormat;
const sampledTexelWeights = new Map<number, PerTexelComponent<number>>();
const unclassifiedStack = [new Set<number>(range(numTexels, v => v))];
while (unclassifiedStack.length > 0) {
const unclassified = unclassifiedStack.pop()!;
const setA = new Set<number>();
const setB = new Set<number>();
[...unclassified.keys()].forEach((t, i) => ((i & 1) === 0 ? setA : setB).add(t));
if (setB.size > 0) {
unclassifiedStack.push(setB);
}
const results = convertResultAsAppropriate(
await run(
range(mipLevelCount, mipLevel =>
TexelView.fromTexelsAsColors(
format,
(coords: Required<GPUOrigin3DDict>): Readonly<PerTexelComponent<number>> => {
const size = mipLevelSizes[mipLevel];
const texelsPerSlice = size[0] * size[1];
const texelsPerRow = size[0];
const texelId =
numTexelsOfPrecedingLevels[mipLevel] +
coords.x +
coords.y * texelsPerRow +
coords.z * texelsPerSlice;
const isCandidate = setA.has(texelId);
const texel: PerTexelComponent<number> = {};
for (const component of rep.componentOrder) {
texel[component] = isCandidate ? 1 : 0;
}
return texel;
}
)
)
),
format
);
if (components.some(c => results[c] !== 0)) {
if (setA.size === 1) {
setA.forEach(texel => sampledTexelWeights.set(texel, results));
} else {
unclassifiedStack.push(setA);
}
}
}
const levels: Map<number, PerTexelComponent<number>>[][] = [];
for (const [texelId, weight] of sampledTexelWeights.entries()) {
const { xyId, z, mipLevel } = getTexelCoordFromTexelId(texelId);
const level = levels[mipLevel] ?? [];
levels[mipLevel] = level;
const layerEntries = level[z] ?? new Map();
level[z] = layerEntries;
layerEntries.set(xyId, weight);
}
const blockParts = {
top: { left: '╔', fill: '═══', right: '╗', block: '╦', texel: '╤' },
mid: { left: '╠', fill: '═══', right: '╣', block: '╬', texel: '╪' },
bot: { left: '╚', fill: '═══', right: '╝', block: '╩', texel: '╧' },
texelMid: { left: '╟', fill: '───', right: '╢', block: '╫', texel: '┼' },
value: { left: '║', fill: ' ', right: '║', block: '║', texel: '│' },
} as const;
const nonBlockParts = {
top: { left: '┌', fill: '───', right: '┐', block: '┬', texel: '┬' },
mid: { left: '├', fill: '───', right: '┤', block: '┼', texel: '┼' },
bot: { left: '└', fill: '───', right: '┘', block: '┴', texel: '┴' },
texelMid: { left: '├', fill: '───', right: '┤', block: '┼', texel: '┼' },
value: { left: '│', fill: ' ', right: '│', block: '│', texel: '│' },
} as const;
const lines: string[] = [];
const letter = (idx: number) => String.fromCodePoint(idx < 30 ? 97 + idx : idx + 9600 - 30);
let idCount = 0;
const { blockWidth, blockHeight } = getBlockInfoForTextureFormat(
softwareTexture.descriptor.format
);
const rangeCat = <T>(num: number, fn: (i: number) => T) => range(num, fn).join('');
const joinFn = (arr: string[], fn: (i: number) => string) => {
const joins = range(arr.length - 1, fn);
return arr.map((s, i) => `${s}${joins[i] ?? ''}`).join('');
};
const parts = Math.max(blockWidth, blockHeight) > 1 ? blockParts : nonBlockParts;
* Makes a row that's [left, fill, texel, fill, block, fill, texel, fill, right]
* except if `contents` is supplied then it would be
* [left, contents[0], texel, contents[1], block, contents[2], texel, contents[3], right]
*/
const makeRow = (
blockPaddedWidth: number,
width: number,
{
left,
fill,
right,
block,
texel,
}: {
left: string;
fill: string;
right: string;
block: string;
texel: string;
},
contents?: string[]
) => {
return `${left}${joinFn(contents ?? range(blockPaddedWidth, x => fill), x => {
return (x + 1) % blockWidth === 0 ? block : texel;
})}${right}`;
};
for (let mipLevel = 0; mipLevel < mipLevelCount; ++mipLevel) {
const level = levels[mipLevel];
if (!level) {
continue;
}
const [width, height, depthOrArrayLayers] = mipLevelSizes[mipLevel];
const texelsPerRow = width;
for (let layer = 0; layer < depthOrArrayLayers; ++layer) {
const layerEntries = level[layer];
const orderedTexelIndices: number[] = [];
lines.push('');
const unSampled = layerEntries ? '' : 'un-sampled';
if (isCube) {
const face = kFaceNames[layer % 6];
lines.push(
`layer: ${layer} mip(${mipLevel}), cube-layer: ${(layer / 6) | 0} (${face}) ${unSampled}`
);
} else {
lines.push(`layer: ${layer} mip(${mipLevel}) ${unSampled}`);
}
if (!layerEntries) {
continue;
}
const blockPaddedHeight = align(height, blockHeight);
const blockPaddedWidth = align(width, blockWidth);
lines.push(` ${rangeCat(width, x => ` ${x.toString().padEnd(2)}`)}`);
lines.push(` ${makeRow(blockPaddedWidth, width, parts.top)}`);
for (let y = 0; y < blockPaddedHeight; y++) {
lines.push(
`${y.toString().padStart(2)} ${makeRow(
blockPaddedWidth,
width,
parts.value,
range(blockPaddedWidth, x => {
const texelIdx = x + y * texelsPerRow;
const weight = layerEntries.get(texelIdx);
const outside = y >= height || x >= width;
if (outside || weight === undefined) {
return outside ? '░░░' : ' ';
} else {
const id = letter(idCount + orderedTexelIndices.length);
orderedTexelIndices.push(texelIdx);
return ` ${id} `;
}
})
)}`
);
const end = y < blockPaddedHeight - 1;
const lineParts = end
? (y + 1) % blockHeight === 0
? parts.mid
: parts.texelMid
: parts.bot;
lines.push(` ${makeRow(blockPaddedWidth, width, lineParts)}`);
}
const pad2 = (n: number) => n.toString().padStart(2);
const pad3 = (n: number) => n.toString().padStart(3);
const fix5 = (n: number) => {
const s = n.toFixed(5);
return s === '0.00000' && n !== 0 ? n.toString() : s;
};
const formatValue = isSintOrUintFormat(format) ? pad3 : fix5;
const formatTexel = (texel: PerTexelComponent<number> | undefined) =>
texel
? Object.entries(texel)
.map(([k, v]) => `${k}: ${formatValue(v)}`)
.join(', ')
: '*texel values unavailable*';
const colorLines: string[] = [];
const compareLines: string[] = [];
let levelWeight = 0;
orderedTexelIndices.forEach((texelIdx, i) => {
const weights = layerEntries.get(texelIdx)!;
const y = Math.floor(texelIdx / texelsPerRow);
const x = texelIdx % texelsPerRow;
const singleWeight = valueIfAllComponentsAreEqual(weights, components)!;
levelWeight += singleWeight;
const w =
singleWeight !== undefined
? `weight: ${fix5(singleWeight)}`
: `weights: [${components.map(c => `${c}: ${fix5(weights[c]!)}`).join(', ')}]`;
const coord = `${pad2(x)}, ${pad2(y)}, ${pad2(layer)}`;
const texel =
texels &&
convertToTexelViewFormat(
texels[mipLevel].color({ x, y, z: layer }),
softwareTexture.descriptor.format
);
const texelStr = formatTexel(texel);
const id = letter(idCount + i);
lines.push(`${id}: mip(${mipLevel}) at: [${coord}], ${w}`);
colorLines.push(`${id}: value: ${texelStr}`);
if (isBuiltinComparison(originalCall.builtin)) {
assert(!!texel);
const compareTexel = applyCompare(originalCall, sampler, [TexelComponent.Depth], texel);
compareLines.push(
`${id}: compare(${sampler.compare}) result with depthRef(${fix5(
originalCall.depthRef!
)}): ${fix5(compareTexel.Depth!)}`
);
}
});
lines.push(...colorLines);
lines.push(...compareLines);
if (!isNaN(levelWeight)) {
lines.push(`mip level (${mipLevel}) weight: ${fix5(levelWeight)}`);
}
idCount += orderedTexelIndices.length;
}
}
return lines;
}
function layoutTwoColumns(columnA: string[], columnB: string[]) {
const widthA = Math.max(...columnA.map(l => l.length));
const lines = Math.max(columnA.length, columnB.length);
const out: string[] = new Array<string>(lines);
for (let line = 0; line < lines; line++) {
const a = columnA[line] ?? '';
const b = columnB[line] ?? '';
out[line] = `${a}${' '.repeat(widthA - a.length)} | ${b}`;
}
return out;
}
* Returns the number of layers ot test for a given view dimension
*/
export function getDepthOrArrayLayersForViewDimension(viewDimension?: GPUTextureViewDimension) {
switch (viewDimension) {
case '1d':
return 1;
case undefined:
case '2d':
return 1;
case '2d-array':
return 4;
case '3d':
return 8;
case 'cube':
return 6;
default:
unreachable();
}
}
* Choose a texture size based on the given parameters.
* The size will be in a multiple of blocks. If it's a cube
* map the size will so be square.
*/
export function chooseTextureSize({
minSize,
minBlocks,
format,
viewDimension,
}: {
minSize: number;
minBlocks: number;
format: GPUTextureFormat;
viewDimension?: GPUTextureViewDimension;
}) {
const { blockWidth, blockHeight } = getBlockInfoForTextureFormat(format);
const width = align(Math.max(minSize, blockWidth * minBlocks), blockWidth);
const height =
viewDimension === '1d' ? 1 : align(Math.max(minSize, blockHeight * minBlocks), blockHeight);
if (viewDimension === 'cube' || viewDimension === 'cube-array') {
const blockLCM = lcm(blockWidth, blockHeight);
const largest = Math.max(width, height);
const size = align(largest, blockLCM);
return [size, size, viewDimension === 'cube-array' ? 24 : 6];
}
const depthOrArrayLayers = getDepthOrArrayLayersForViewDimension(viewDimension);
return [width, height, depthOrArrayLayers];
}
export const kSamplePointMethods = ['texel-centre', 'spiral'] as const;
export type SamplePointMethods = (typeof kSamplePointMethods)[number];
export const kCubeSamplePointMethods = ['cube-edges', 'texel-centre', 'spiral'] as const;
export type CubeSamplePointMethods = (typeof kSamplePointMethods)[number];
type TextureBuiltinInputArgs = {
textureBuiltin?: TextureBuiltin;
descriptor?: GPUTextureDescriptor;
softwareTexture?: SoftwareTexture;
sampler?: GPUSamplerDescriptor;
derivatives?: boolean;
mipLevel?: RangeDef;
sampleIndex?: RangeDef;
arrayIndex?: RangeDef;
grad?: boolean;
bias?: boolean;
component?: boolean;
depthRef?: boolean;
offset?: boolean;
hashInputs: (number | string | boolean)[];
};
* Generates an array of coordinates at which to sample a texture.
*/
function generateTextureBuiltinInputsImpl<T extends Dimensionality>(
makeValue: (x: number, y: number, z: number) => T,
n: number,
args:
| (TextureBuiltinInputArgs & {
method: 'texel-centre';
})
| (TextureBuiltinInputArgs & {
method: 'spiral';
radius?: number;
loops?: number;
})
): {
coords: T;
derivativeMult?: T;
ddx?: T;
ddy?: T;
mipLevel: number;
sampleIndex?: number;
arrayIndex?: number;
bias?: number;
offset?: T;
component?: number;
depthRef?: number;
}[] {
const { method, descriptor, softwareTexture: info } = args;
assert(!!descriptor !== !!info, 'must pass descriptor or textureInfo');
const textureInfo: SoftwareTexture = info ?? {
descriptor: descriptor!,
texels: [],
viewDescriptor: {},
};
const { mipLevelCount, baseMipLevelSize } = getBaseMipLevelInfo(textureInfo);
const dimension = textureInfo.descriptor.dimension ?? '2d';
const coords: T[] = [];
switch (method) {
case 'texel-centre': {
for (let i = 0; i < n; i++) {
const r = hashU32(i);
const x = Math.floor(lerp(0, baseMipLevelSize[0] - 1, (r & 0xff) / 0xff)) + 0.5;
const y = Math.floor(lerp(0, baseMipLevelSize[1] - 1, ((r >> 8) & 0xff) / 0xff)) + 0.5;
const z = Math.floor(lerp(0, baseMipLevelSize[2] - 1, ((r >> 16) & 0xff) / 0xff)) + 0.5;
coords.push(
makeValue(x / baseMipLevelSize[0], y / baseMipLevelSize[1], z / baseMipLevelSize[2])
);
}
break;
}
case 'spiral': {
const { radius = 1.5, loops = 2 } = args;
for (let i = 0; i < n; i++) {
const f = i / (Math.max(n, 2) - 1);
const r = radius * f;
const a = loops * 2 * Math.PI * f;
coords.push(makeValue(0.5 + r * Math.cos(a), 0.5 + r * Math.sin(a), 0));
}
break;
}
}
const _hashInputs = args.hashInputs.map(v =>
typeof v === 'string' ? sumOfCharCodesOfString(v) : typeof v === 'boolean' ? (v ? 1 : 0) : v
);
const makeRandValue = ({ num, type }: RangeDef, ...hashInputs: number[]) => {
const range = num;
const number = (hashU32(..._hashInputs, ...hashInputs) / 0x1_0000_0000) * range;
return type === 'f32' ? number : Math.floor(number);
};
const makeRangeValue = ({ num, type }: RangeDef, ...hashInputs: number[]) => {
const range = num + (type === 'u32' ? 1 : 2);
const number =
(hashU32(..._hashInputs, ...hashInputs) / 0x1_0000_0000) * range - (type === 'u32' ? 0 : 1);
return type === 'f32' ? number : Math.floor(number);
};
const makeIntHashValueRepeatable = (min: number, max: number, ...hashInputs: number[]) => {
const range = max - min;
return min + Math.floor((hashU32(...hashInputs) / 0x1_0000_0000) * range);
};
const kSubdivisionsPerTexel = 4;
const avoidEdgeCase =
!args.sampler || args.sampler.minFilter === 'nearest' || isBuiltinGather(args.textureBuiltin);
const edgeRemainder = isBuiltinGather(args.textureBuiltin) ? kSubdivisionsPerTexel / 2 : 0;
const euclideanModulo = (n: number, m: number) => ((n % m) + m) % m;
const addressMode: GPUAddressMode[] =
args.textureBuiltin === 'textureSampleBaseClampToEdge'
? ['clamp-to-edge', 'clamp-to-edge', 'clamp-to-edge']
: [
args.sampler?.addressModeU ?? 'clamp-to-edge',
args.sampler?.addressModeV ?? 'clamp-to-edge',
args.sampler?.addressModeW ?? 'clamp-to-edge',
];
const avoidTextureEdge = (axis: number, textureDimensionUnits: number, v: number) => {
assert(isBuiltinGather(args.textureBuiltin));
if (addressMode[axis] === 'repeat') {
return v;
}
const inside = euclideanModulo(v, textureDimensionUnits);
const outside = v - inside;
return outside + clamp(inside, { min: 1, max: textureDimensionUnits - 1 });
};
const numComponents = isDepthOrStencilTextureFormat(textureInfo.descriptor.format) ? 1 : 4;
return coords.map((c, i) => {
const mipLevel = args.mipLevel
? quantizeMipLevel(makeRangeValue(args.mipLevel, i), args.sampler?.mipmapFilter ?? 'nearest')
: 0;
const clampedMipLevel = clamp(mipLevel, { min: 0, max: mipLevelCount - 1 });
const mipSize = virtualMipSize(dimension, baseMipLevelSize, clampedMipLevel);
const q = mipSize.map(v => v * kSubdivisionsPerTexel);
const coords = c.map((v, i) => {
const v1 = Math.floor(v * q[i]);
const isTexelEdgeCase = Math.abs(v1 % kSubdivisionsPerTexel) === edgeRemainder;
const v2 = isTexelEdgeCase && avoidEdgeCase ? v1 + 1 : v1;
const v3 = isBuiltinGather(args.textureBuiltin) ? avoidTextureEdge(i, q[i], v2) : v2;
return v3 / q[i];
}) as T;
const makeGradient = <T>(hashInput: number): T => {
return coords.map((_, i) => {
const intPart = makeRangeValue({ num: 8, type: 'u32' }, i, hashInput) - 4;
const fractPart = makeRangeValue({ num: 0, type: 'f32' }, i, hashInput + 1) * 0.25;
assert(fractPart >= -0.25 && fractPart <= 0.25);
return intPart + fractPart;
}) as T;
};
const makeDerivativeMult = (coords: T, mipLevel: number): T => {
const mult = new Array(coords.length).fill(0);
const ndx = makeRangeValue({ num: coords.length - 1, type: 'u32' }, i, 8);
assert(ndx < coords.length);
mult[ndx] = Math.pow(2, mipLevel);
return mult as T;
};
const chooseMipLevel = () => {
const innerLevelR = makeRandValue({ num: 9, type: 'u32' }, i, 11);
const innerLevel =
args?.sampler?.mipmapFilter === 'linear'
? innerLevelR + 1
: innerLevelR < 5
? innerLevelR
: innerLevelR + 1;
const outerLevel = makeRangeValue({ num: mipLevelCount - 1, type: 'i32' }, i, 11);
return outerLevel + innerLevel / 10;
};
const makeDerivativeMultForTextureSample = (coords: T): T => {
const mipLevel = chooseMipLevel();
return makeDerivativeMult(coords, mipLevel);
};
const makeBiasAndDerivativeMult = (coords: T): [number, T] => {
const testType = makeRandValue({ num: 4, type: 'u32' }, i, 11);
let mipLevel;
let bias;
switch (testType) {
case 0:
mipLevel = mipLevelCount + 1;
bias = -25;
break;
case 1:
mipLevel = -1;
bias = 25;
break;
default:
mipLevel = chooseMipLevel();
bias = makeRangeValue({ num: 6, type: 'f32' }, i, 9) - 3;
break;
}
const clampedBias = clamp(bias, { min: -16, max: 15.99 });
const derivativeBasedMipLevel = mipLevel - clampedBias;
const derivativeMult = makeDerivativeMult(coords, derivativeBasedMipLevel);
return [bias, derivativeMult];
};
const [bias, derivativeMult] = args.bias
? makeBiasAndDerivativeMult(coords)
: args.derivatives
? [undefined, makeDerivativeMultForTextureSample(coords)]
: [];
return {
coords,
derivativeMult,
mipLevel,
sampleIndex: args.sampleIndex ? makeRangeValue(args.sampleIndex, i, 1) : undefined,
arrayIndex: args.arrayIndex ? makeRangeValue(args.arrayIndex, i, 2) : undefined,
depthRef: args.depthRef ? makeRandValue({ num: 3, type: 'u32' }, i, 5) / 2 : undefined,
ddx: args.grad ? makeGradient(7) : undefined,
ddy: args.grad ? makeGradient(8) : undefined,
bias,
offset: args.offset
? (coords.map((_, j) => makeIntHashValueRepeatable(-8, 8, i, 3 + j)) as T)
: undefined,
component: args.component ? makeIntHashValueRepeatable(0, numComponents, i, 4) : undefined,
};
});
}
* When mipmapFilter === 'nearest' we need to stay away from 0.5
* because the GPU could decide to choose one mip or the other.
*
* Some example transition values, the value at which the GPU chooses
* mip level 1 over mip level 0:
*
* M1 Mac: 0.515381
* Intel Mac: 0.49999
* AMD Mac: 0.5
*/
const kMipEpsilon = 0.02;
function quantizeMipLevel(mipLevel: number, mipmapFilter: GPUMipmapFilterMode) {
if (mipmapFilter === 'linear') {
return mipLevel;
}
const intMip = Math.floor(mipLevel);
const fractionalMip = mipLevel - intMip;
if (fractionalMip < 0.5 - kMipEpsilon || fractionalMip > 0.5 + kMipEpsilon) {
return mipLevel;
} else {
return intMip + 0.5 + (fractionalMip < 0.5 ? -kMipEpsilon : +kMipEpsilon);
}
}
type FilterFirstElement<T extends unknown[]> = T extends [unknown, ...infer R] ? R : [];
type GenerateTextureBuiltinInputsImplArgs = FilterFirstElement<
Parameters<typeof generateTextureBuiltinInputsImpl>
>;
export function generateTextureBuiltinInputs1D(...args: GenerateTextureBuiltinInputsImplArgs) {
return generateTextureBuiltinInputsImpl<vec1>((x: number) => [x], ...args);
}
export function generateTextureBuiltinInputs2D(...args: GenerateTextureBuiltinInputsImplArgs) {
return generateTextureBuiltinInputsImpl<vec2>((x: number, y: number) => [x, y], ...args);
}
export function generateTextureBuiltinInputs3D(...args: GenerateTextureBuiltinInputsImplArgs) {
return generateTextureBuiltinInputsImpl<vec3>(
(x: number, y: number, z: number) => [x, y, z],
...args
);
}
type mat3 =
[
number, number, number,
number, number, number,
number, number, number,
];
const kFaceUVMatrices: mat3[] =
[
[ 0, 0, -2, 0, -2, 0, 1, 1, 1],
[ 0, 0, 2, 0, -2, 0, -1, 1, -1],
[ 2, 0, 0, 0, 0, 2, -1, 1, -1],
[ 2, 0, 0, 0, 0, -2, -1, -1, 1],
[ 2, 0, 0, 0, -2, 0, -1, 1, 1],
[-2, 0, 0, 0, -2, 0, 1, 1, -1],
];
function transformMat3(v: vec3, m: mat3): vec3 {
const x = v[0];
const y = v[1];
const z = v[2];
return [
x * m[0] + y * m[3] + z * m[6],
x * m[1] + y * m[4] + z * m[7],
x * m[2] + y * m[5] + z * m[8],
];
}
function normalize(v: vec3): vec3 {
const length = Math.sqrt(v[0] * v[0] + v[1] * v[1] + v[2] * v[2]);
assert(length > 0);
return v.map(v => v / length) as vec3;
}
* Converts a cube map coordinate to a uv coordinate (0 to 1) and layer (0.5/6.0 to 5.5/6.0).
*/
function convertCubeCoordToNormalized3DTextureCoord(v: vec3): vec3 {
let uvw;
let layer;
const r = normalize(v);
const absR = r.map(v => Math.abs(v));
if (absR[0] > absR[1] && absR[0] > absR[2]) {
const negX = r[0] < 0.0 ? 1 : 0;
uvw = [negX ? r[2] : -r[2], -r[1], absR[0]];
layer = negX;
} else if (absR[1] > absR[2]) {
const negY = r[1] < 0.0 ? 1 : 0;
uvw = [r[0], negY ? -r[2] : r[2], absR[1]];
layer = 2 + negY;
} else {
const negZ = r[2] < 0.0 ? 1 : 0;
uvw = [negZ ? -r[0] : r[0], -r[1], absR[2]];
layer = 4 + negZ;
}
return [(uvw[0] / uvw[2] + 1) * 0.5, (uvw[1] / uvw[2] + 1) * 0.5, (layer + 0.5) / 6];
}
* Convert a 3d texcoord into a cube map coordinate.
*/
function convertNormalized3DTexCoordToCubeCoord(uvLayer: vec3) {
const [u, v, faceLayer] = uvLayer;
return normalize(transformMat3([u, v, 1], kFaceUVMatrices[Math.min(5, faceLayer * 6) | 0]));
}
* Wrap a texel based face coord across cube faces
*
* We have a face texture in texels coord where U/V choose a texel and W chooses the face.
* If U/V are outside the size of the texture then, when normalized and converted
* to a cube map coordinate, they'll end up pointing to a different face.
*
* addressMode is effectively ignored for cube
*
* By converting from a texel based coord to a normalized coord and then to a cube map coord,
* if the texel was outside of the face, the cube map coord will end up pointing to a different
* face. We then convert back cube coord -> normalized face coord -> texel based coord
*/
function wrapFaceCoordToCubeFaceAtEdgeBoundaries(mipLevelSize: number, faceCoord: vec3) {
const nc0: vec3 = [
(faceCoord[0] + 0.5) / mipLevelSize,
(faceCoord[1] + 0.5) / mipLevelSize,
(faceCoord[2] + 0.5) / 6,
];
const cc = convertNormalized3DTexCoordToCubeCoord(nc0);
const nc1 = convertCubeCoordToNormalized3DTextureCoord(cc);
const fc = [
Math.floor(nc1[0] * mipLevelSize),
Math.floor(nc1[1] * mipLevelSize),
Math.floor(nc1[2] * 6),
];
return fc;
}
function applyAddressModesToCoords(
addressMode: GPUAddressMode[],
mipLevelSize: number[],
coord: number[]
) {
return coord.map((v, i) => {
switch (addressMode[i]) {
case 'clamp-to-edge':
return clamp(v, { min: 0, max: mipLevelSize[i] - 1 });
case 'mirror-repeat': {
const n = Math.floor(v / mipLevelSize[i]);
v = v - n * mipLevelSize[i];
return (n & 1) !== 0 ? mipLevelSize[i] - v - 1 : v;
}
case 'repeat':
return v - Math.floor(v / mipLevelSize[i]) * mipLevelSize[i];
default:
unreachable();
}
});
}
* Generates an array of coordinates at which to sample a texture for a cubemap
*/
export function generateSamplePointsCube(
n: number,
args:
| (TextureBuiltinInputArgs & {
method: 'texel-centre';
})
| (TextureBuiltinInputArgs & {
method: 'spiral';
radius?: number;
loops?: number;
})
| (TextureBuiltinInputArgs & {
method: 'cube-edges';
})
): {
coords: vec3;
derivativeMult?: vec3;
ddx?: vec3;
ddy?: vec3;
mipLevel: number;
arrayIndex?: number;
bias?: number;
offset?: undefined;
component?: number;
depthRef?: number;
}[] {
const { method, descriptor, softwareTexture: info } = args;
assert(!!descriptor !== !!info, 'must pass descriptor or textureInfo');
const textureInfo: SoftwareTexture = info ?? {
descriptor: descriptor!,
texels: [],
viewDescriptor: {},
};
const { mipLevelCount, baseMipLevelSize } = getBaseMipLevelInfo(textureInfo);
const textureWidth = baseMipLevelSize[0];
const coords: vec3[] = [];
switch (method) {
case 'texel-centre': {
for (let i = 0; i < n; i++) {
const r = hashU32(i);
const u = (Math.floor(lerp(0, textureWidth - 1, (r & 0xff) / 0xff)) + 0.5) / textureWidth;
const v =
(Math.floor(lerp(0, textureWidth - 1, ((r >> 8) & 0xff) / 0xff)) + 0.5) / textureWidth;
const face = Math.floor(lerp(0, 6, ((r >> 16) & 0xff) / 0x100));
coords.push(convertNormalized3DTexCoordToCubeCoord([u, v, face]));
}
break;
}
case 'spiral': {
const { radius = 1.5, loops = 2 } = args;
for (let i = 0; i < n; i++) {
const f = (i + 1) / (Math.max(n, 2) - 1);
const r = radius * f;
const theta = loops * 2 * Math.PI * f;
const phi = loops * 1.3 * Math.PI * f;
const sinTheta = Math.sin(theta);
const cosTheta = Math.cos(theta);
const sinPhi = Math.sin(phi);
const cosPhi = Math.cos(phi);
const ux = cosTheta * sinPhi;
const uy = cosPhi;
const uz = sinTheta * sinPhi;
coords.push([ux * r, uy * r, uz * r]);
}
break;
}
case 'cube-edges': {
coords.push(
[ 1 , -1.01, 0 ],
[ 1 , +1.01, 0 ],
[ 1 , 0 , -1.01 ],
[ 1 , 0 , +1.01 ],
[ -1 , -1.01, 0 ],
[ -1 , +1.01, 0 ],
[ -1 , 0 , -1.01 ],
[ -1 , 0 , +1.01 ],
[ -1.01, 1 , 0 ],
[ +1.01, 1 , 0 ],
[ 0 , 1 , -1.01 ],
[ 0 , 1 , +1.01 ],
[ -1.01, -1 , 0 ],
[ +1.01, -1 , 0 ],
[ 0 , -1 , -1.01 ],
[ 0 , -1 , +1.01 ],
[ -1.01, 0 , 1 ],
[ +1.01, 0 , 1 ],
[ 0 , -1.01, 1 ],
[ 0 , +1.01, 1 ],
[ -1.01, 0 , -1 ],
[ +1.01, 0 , -1 ],
[ 0 , -1.01, -1 ],
[ 0 , +1.01, -1 ],
);
break;
}
}
const _hashInputs = args.hashInputs.map(v =>
typeof v === 'string' ? sumOfCharCodesOfString(v) : typeof v === 'boolean' ? (v ? 1 : 0) : v
);
const makeRandValue = ({ num, type }: RangeDef, ...hashInputs: number[]) => {
const range = num;
const number = (hashU32(..._hashInputs, ...hashInputs) / 0x1_0000_0000) * range;
return type === 'f32' ? number : Math.floor(number);
};
const makeRangeValue = ({ num, type }: RangeDef, ...hashInputs: number[]) => {
const range = num + (type === 'u32' ? 1 : 2);
const number =
(hashU32(..._hashInputs, ...hashInputs) / 0x1_0000_0000) * range - (type === 'u32' ? 0 : 1);
return type === 'f32' ? number : Math.floor(number);
};
const makeIntHashValue = (min: number, max: number, ...hashInputs: number[]) => {
const range = max - min;
return min + Math.floor((hashU32(..._hashInputs, ...hashInputs) / 0x1_0000_0000) * range);
};
const kSubdivisionsPerTexel = 4;
const avoidEdgeCase =
!args.sampler || args.sampler.minFilter === 'nearest' || isBuiltinGather(args.textureBuiltin);
const edgeRemainder = isBuiltinGather(args.textureBuiltin) ? kSubdivisionsPerTexel / 2 : 0;
return coords.map((c, i) => {
const mipLevel = args.mipLevel
? quantizeMipLevel(makeRangeValue(args.mipLevel, i), args.sampler?.mipmapFilter ?? 'nearest')
: 0;
const clampedMipLevel = clamp(mipLevel, { min: 0, max: mipLevelCount - 1 });
const mipSize = virtualMipSize('2d', baseMipLevelSize, Math.ceil(clampedMipLevel));
const q = [
mipSize[0] * kSubdivisionsPerTexel,
mipSize[0] * kSubdivisionsPerTexel,
6 * kSubdivisionsPerTexel,
];
const uvw = convertCubeCoordToNormalized3DTextureCoord(c);
const ndx = getUnusedCubeCornerSampleIndex(mipSize[0], uvw);
if (ndx >= 0) {
const halfTexel = 0.5 / mipSize[0];
uvw[0] = clamp(uvw[0], { min: halfTexel, max: 1 - halfTexel });
}
const quantizedUVW = uvw.map((v, i) => {
const v1 = Math.floor(v * q[i]);
const isEdgeCase = Math.abs(v1 % kSubdivisionsPerTexel) === edgeRemainder;
const v2 = isEdgeCase && avoidEdgeCase ? v1 + 1 : v1;
return (v2 + 1 / 16) / q[i];
}) as vec3;
const quantize = (v: number, units: number) => Math.floor(v * units) * units;
const makeGradient = <T>(hashInput: number): T => {
return coords.map((_, i) =>
quantize(makeRangeValue({ num: 8, type: 'f32' }, i, hashInput) - 4, 1 / 3)
) as T;
};
const coords = convertNormalized3DTexCoordToCubeCoord(quantizedUVW);
const makeDerivativeMult = (coords: vec3, mipLevel: number): vec3 => {
const mult = new Array(coords.length).fill(0);
const ndx = makeRangeValue({ num: coords.length - 1, type: 'u32' }, i, 8);
assert(ndx < coords.length);
mult[ndx] = Math.pow(2, mipLevel);
return mult as vec3;
};
const chooseMipLevel = () => {
const innerLevelR = makeRandValue({ num: 9, type: 'u32' }, i, 11);
const innerLevel =
args?.sampler?.mipmapFilter === 'linear'
? innerLevelR + 1
: innerLevelR < 4
? innerLevelR
: innerLevelR + 1;
const outerLevel = makeRangeValue({ num: mipLevelCount - 1, type: 'i32' }, i, 11);
return outerLevel + innerLevel / 10;
};
const makeDerivativeMultForTextureSample = (coords: vec3): vec3 => {
const mipLevel = chooseMipLevel();
return makeDerivativeMult(coords, mipLevel);
};
const makeBiasAndDerivativeMult = (coords: vec3): [number, vec3] => {
const testType = makeRandValue({ num: 4, type: 'u32' }, i, 11);
let mipLevel;
let bias;
switch (testType) {
case 0:
mipLevel = mipLevelCount + 1;
bias = -25;
break;
case 1:
mipLevel = -1;
bias = 25;
break;
default:
mipLevel = chooseMipLevel();
bias = makeRangeValue({ num: 6, type: 'f32' }, i, 9) - 3;
break;
}
const clampedBias = clamp(bias, { min: -16, max: 15.99 });
const derivativeBasedMipLevel = mipLevel - clampedBias;
const derivativeMult = makeDerivativeMult(coords, derivativeBasedMipLevel);
return [bias, derivativeMult];
};
const [bias, derivativeMult] = args.bias
? makeBiasAndDerivativeMult(coords)
: args.derivatives
? [undefined, makeDerivativeMultForTextureSample(coords)]
: [];
return {
coords,
derivativeMult,
ddx: args.grad ? makeGradient(7) : undefined,
ddy: args.grad ? makeGradient(8) : undefined,
mipLevel,
arrayIndex: args.arrayIndex ? makeRangeValue(args.arrayIndex, i, 2) : undefined,
bias,
depthRef: args.depthRef ? makeRandValue({ num: 3, type: 'u32' }, i, 5) / 2 : undefined,
component: args.component ? makeIntHashValue(0, 4, i, 4) : undefined,
};
});
}
function wgslTypeFor(data: number | Dimensionality, type: 'f' | 'i' | 'u'): string {
if (Array.isArray(data)) {
switch (data.length) {
case 1:
return `${type}32`;
case 2:
return `vec2${type}`;
case 3:
return `vec3${type}`;
default:
unreachable();
}
}
return `${type}32`;
}
function wgslExpr(
data: number | Readonly<vec1> | Readonly<vec2> | Readonly<vec3> | Readonly<vec4>
): string {
if (Array.isArray(data)) {
switch (data.length) {
case 1:
return data[0].toString();
case 2:
return `vec2(${data.map(v => v.toString()).join(', ')})`;
case 3:
return `vec3(${data.map(v => v.toString()).join(', ')})`;
default:
unreachable();
}
}
return data.toString();
}
function wgslExprFor(data: number | vec1 | vec2 | vec3 | vec4, type: 'f' | 'i' | 'u'): string {
if (Array.isArray(data)) {
switch (data.length) {
case 1:
return `${type}(${data[0].toString()})`;
case 2:
return `vec2${type}(${data.map(v => v.toString()).join(', ')})`;
case 3:
return `vec3${type}(${data.map(v => v.toString()).join(', ')})`;
default:
unreachable();
}
}
return `${type}32(${data.toString()})`;
}
function binKey<T extends Dimensionality>(call: TextureCall<T>): string {
const keys: string[] = [];
for (const name of kTextureCallArgNames) {
const value = call[name];
if (value !== undefined) {
if (name === 'offset' || name === 'component') {
keys.push(`${name}: ${wgslExpr(value)}`);
} else {
keys.push(`${name}: ${wgslTypeFor(value, call.coordType)}`);
}
}
}
return `${call.builtin}(${keys.join(', ')})`;
}
function buildBinnedCalls<T extends Dimensionality>(calls: TextureCall<T>[]) {
const args: string[] = [];
const fields: string[] = [];
const data: number[] = [];
const prototype = calls[0];
if (isBuiltinGather(prototype.builtin) && prototype['componentType']) {
args.push(`/* component */ ${wgslExpr(prototype['component']!)}`);
}
args.push('T');
if (builtinNeedsSampler(prototype.builtin)) {
args.push('S');
}
for (const name of kTextureCallArgNames) {
const value = prototype[name];
if (value !== undefined) {
if (name === 'offset') {
args.push(`/* offset */ ${wgslExpr(value)}`);
} else if (name === 'component') {
} else {
const type =
name === 'mipLevel'
? prototype.levelType!
: name === 'arrayIndex'
? prototype.arrayIndexType!
: name === 'sampleIndex'
? prototype.sampleIndexType!
: name === 'bias' || name === 'depthRef' || name === 'ddx' || name === 'ddy'
? 'f'
: prototype.coordType;
if (name !== 'derivativeMult') {
args.push(
`args.${name}${
name === 'coords' && builtinNeedsDerivatives(prototype.builtin)
? ' + derivativeBase * args.derivativeMult'
: ''
}`
);
}
fields.push(`@align(16) ${name} : ${wgslTypeFor(value, type)}`);
}
}
}
for (const call of calls) {
for (const name of kTextureCallArgNames) {
const value = call[name];
assert(
(prototype[name] === undefined) === (value === undefined),
'texture calls are not binned correctly'
);
if (value !== undefined && name !== 'offset' && name !== 'component') {
const type = getCallArgType<T>(call, name);
const bitcastToU32 = kBitCastFunctions[type];
if (value instanceof Array) {
for (const c of value) {
data.push(bitcastToU32(c));
}
} else {
data.push(bitcastToU32(value));
}
while ((data.length & 3) !== 0) {
data.push(0);
}
}
}
}
const expr = `${prototype.builtin}(${args.join(', ')})`;
return { expr, fields, data };
}
function binCalls<T extends Dimensionality>(calls: TextureCall<T>[]): number[][] {
const map = new Map<string, number>();
const bins: number[][] = [];
calls.forEach((call, callIdx) => {
const key = binKey(call);
const binIdx = map.get(key);
if (binIdx === undefined) {
map.set(key, bins.length);
bins.push([callIdx]);
} else {
bins[binIdx].push(callIdx);
}
});
return bins;
}
function describeTextureCall<T extends Dimensionality>(call: TextureCall<T>): string {
const args: string[] = [];
if (isBuiltinGather(call.builtin) && call.componentType) {
args.push(`component: ${wgslExprFor(call.component!, call.componentType)}`);
}
args.push('texture: T');
if (builtinNeedsSampler(call.builtin)) {
args.push('sampler: S');
}
for (const name of kTextureCallArgNames) {
const value = call[name];
if (value !== undefined && name !== 'component') {
if (name === 'coords') {
const derivativeWGSL = builtinNeedsDerivatives(call.builtin)
? ` + derivativeBase * derivativeMult(${
call.derivativeMult ? wgslExprFor(call.derivativeMult, call.coordType) : '1'
})`
: '';
args.push(`${name}: ${wgslExprFor(value, call.coordType)}${derivativeWGSL}`);
} else if (name === 'derivativeMult') {
} else if (name === 'ddx' || name === 'ddy') {
args.push(`${name}: ${wgslExprFor(value, call.coordType)}`);
} else if (name === 'mipLevel') {
args.push(`${name}: ${wgslExprFor(value, call.levelType!)}`);
} else if (name === 'arrayIndex') {
args.push(`${name}: ${wgslExprFor(value, call.arrayIndexType!)}`);
} else if (name === 'bias') {
args.push(`${name}: ${wgslExprFor(value, 'f')}`);
} else if (name === 'sampleIndex') {
args.push(`${name}: ${wgslExprFor(value, call.sampleIndexType!)}`);
} else if (name === 'depthRef') {
args.push(`${name}: ${wgslExprFor(value, 'f')}`);
} else {
args.push(`${name}: ${wgslExpr(value)}`);
}
}
}
return `${call.builtin}(${args.join(', ')})`;
}
const getAspectForTexture = (texture: GPUTexture | GPUExternalTexture): GPUTextureAspect =>
texture instanceof GPUExternalTexture
? 'all'
: isDepthTextureFormat(texture.format)
? 'depth-only'
: isStencilTextureFormat(texture.format)
? 'stencil-only'
: 'all';
const s_deviceToPipelines = new WeakMap<
GPUDevice,
Map<string, GPURenderPipeline | GPUComputePipeline>
>();
* Given a list of "calls", each one of which has a texture coordinate,
* generates a fragment shader that uses the instance_index as an index. That
* index is then used to look up a coordinate from a storage buffer which is
* used to call the WGSL texture function to read/sample the texture, and then
* write to a storage buffer. We then read the storage buffer for the per "call"
* results.
*
* We use a 1x1 target and use instance drawing, once instance per call.
* This allows use to more easily adjust derivatives per call.
*
* An issue we ran into before this "one draw call per instance" change;
* Before we had a single draw call and wrote the result of one call per
* pixel rendered.
*
* Imagine we have code like this:
*
* ```
* @group(0) @binding(0) var T: texture_2d<f32>;
* @group(0) @binding(1) var S: sampler;
* @group(0) @binding(2) var<storage> coords: array<vec4f>;
* @fragment fn fs(@builtin(position) pos: vec4f) -> vec4f {
* let ndx = u32(pos.x) * u32(pos.y) * targetWidth;
* return textureSample(T, S, coords[ndx].xy);
* }
* ```
*
* T points to 8x8 pixel texture with 3 mip levels
* S is 'nearest'
* coords: is a storage buffer, 16 bytes long [0,0,0,0], one vec4f.
* our render target is 1x1 pixels
*
* Looking above it appears `ndx` will only ever be 0 but that's
* not what happens. Instead, the GPU will run the fragment shader for
* a 2x2 area. It does this to compute derivatives by running the code
* above and looking at what values it gets passed as coords to
* textureSample. When it does this it ends up with
*
* ndx = 0 for invocation 0
* ndx = 1 for invocation 1
* ndx = 0 + 1 * targetWidth for invocation 2
* ndx = 1 + 1 * targetWidth for invocation 3
*
* In 3 of those cases `ndx` is out of bounds with respect to `coords`.
* Out of bounds access is indeterminate. That means the derivatives are
* indeterminate so what lod it tries to read is indeterminate.
*
* By using instance_index for ndx we avoid this issue. ndx is the same
* on all 4 executions.
*
* Calls are "binned" by call parameters. Each bin has its own structure and
* field in the storage buffer. This allows the calls to be non-homogenous and
* each have their own data type for coordinates.
*
* Note: this function returns:
*
* 'results': an array of results, one for each call.
*
* 'run': a function that accepts a texture and runs the same class pipeline with
* that texture as input, returning an array of results. This can be used by
* identifySamplePoints to query the mix-weights used. We do this so we're
* using the same shader that generated the original results when querying
* the weights.
*
* 'destroy': a function that cleans up the buffers used by `run`.
*/
function createTextureCallsRunner<T extends Dimensionality>(
t: GPUTest,
{
format,
dimension,
sampleCount,
depthOrArrayLayers,
}: {
format: GPUTextureFormat;
dimension: GPUTextureDimension;
sampleCount: number;
depthOrArrayLayers: number;
},
viewDescriptor: GPUTextureViewDescriptor,
textureType: string,
sampler: GPUSamplerDescriptor | undefined,
calls: TextureCall<T>[],
stage: ShaderStage
) {
let structs = '';
let body = '';
let dataFields = '';
const data: number[] = [];
let callCount = 0;
const binned = binCalls(calls);
binned.forEach((binCalls, binIdx) => {
const b = buildBinnedCalls(binCalls.map(callIdx => calls[callIdx]));
structs += `struct Args${binIdx} {
${b.fields.join(',\n ')}
}
`;
dataFields += ` args${binIdx} : array<Args${binIdx}, ${binCalls.length}>,
`;
body += `
{
let is_active = (idx >= ${callCount}) & (idx < ${callCount + binCalls.length});
let args = data.args${binIdx}[idx - ${callCount}];
let call = ${b.expr};
result = select(result, call, is_active);
}
`;
callCount += binCalls.length;
data.push(...b.data);
});
const dataBuffer = t.createBufferTracked({
label: 'createTextureCallsRunner:dataBuffer',
size: data.length * 4,
usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.UNIFORM,
});
t.device.queue.writeBuffer(dataBuffer, 0, new Uint32Array(data));
const builtin = calls[0].builtin;
const isCompare = isBuiltinComparison(builtin);
const { resultType, resultFormat, componentType } = isBuiltinGather(builtin)
? getTextureFormatTypeInfo(format)
: textureType === 'texture_external'
? ({ resultType: 'vec4f', resultFormat: 'rgba32float', componentType: 'f32' } as const)
: textureType.includes('depth')
? ({ resultType: 'f32', resultFormat: 'rgba32float', componentType: 'f32' } as const)
: getTextureFormatTypeInfo(format);
const returnType = `vec4<${componentType}>`;
const samplerType = isCompare ? 'sampler_comparison' : 'sampler';
const renderTarget = t.createTextureTracked({
format: 'rgba32uint',
size: [calls.length, 1],
usage: GPUTextureUsage.COPY_SRC | GPUTextureUsage.RENDER_ATTACHMENT,
});
const derivativeBaseWGSL = `
let derivativeBase = ${
isCubeViewDimension(viewDescriptor)
? '(v.pos.xyx - 0.5 - vec3f(f32(v.ndx), 0, f32(v.ndx))) / vec3f(vec2f(textureDimensions(T)), 1.0)'
: dimension === '1d'
? 'f32(v.pos.x - 0.5 - f32(v.ndx)) / f32(textureDimensions(T))'
: dimension === '3d'
? 'vec3f(v.pos.xy - 0.5 - vec2f(f32(v.ndx), 0), 0) / vec3f(textureDimensions(T))'
: '(v.pos.xy - 0.5 - vec2f(f32(v.ndx), 0)) / vec2f(textureDimensions(T))'
};`;
const derivativeType =
isCubeViewDimension(viewDescriptor) || dimension === '3d'
? 'vec3f'
: dimension === '1d'
? 'f32'
: 'vec2f';
const stageWGSL =
stage === 'vertex'
? `
// --------------------------- vertex stage shaders --------------------------------
@vertex fn vsVertex(
@builtin(vertex_index) vertex_index : u32,
@builtin(instance_index) instance_index : u32) -> VOut {
let positions = array(vec2f(-1, 3), vec2f(3, -1), vec2f(-1, -1));
return VOut(vec4f(positions[vertex_index], 0, 1),
instance_index,
getResult(instance_index, ${derivativeType}(0)));
}
@fragment fn fsVertex(v: VOut) -> @location(0) vec4u {
return bitcast<vec4u>(v.result);
}
`
: stage === 'fragment'
? `
// --------------------------- fragment stage shaders --------------------------------
@vertex fn vsFragment(
@builtin(vertex_index) vertex_index : u32,
@builtin(instance_index) instance_index : u32) -> VOut {
let positions = array(vec2f(-1, 3), vec2f(3, -1), vec2f(-1, -1));
return VOut(vec4f(positions[vertex_index], 0, 1), instance_index, ${returnType}(0));
}
@fragment fn fsFragment(v: VOut) -> @location(0) vec4u {
${derivativeBaseWGSL}
return bitcast<vec4u>(getResult(v.ndx, derivativeBase));
}
`
: `
// --------------------------- compute stage shaders --------------------------------
@group(1) @binding(0) var<storage, read_write> results: array<${returnType}>;
@compute @workgroup_size(1) fn csCompute(@builtin(global_invocation_id) id: vec3u) {
results[id.x] = getResult(id.x, ${derivativeType}(0));
}
`;
const code = `
${structs}
struct Data {
${dataFields}
}
struct VOut {
@builtin(position) pos: vec4f,
@location(0) @interpolate(flat, either) ndx: u32,
@location(1) @interpolate(flat, either) result: ${returnType},
};
@group(0) @binding(0) var T : ${textureType};
${sampler ? `@group(0) @binding(1) var S : ${samplerType}` : ''};
@group(0) @binding(2) var<uniform> data : Data;
fn getResult(idx: u32, derivativeBase: ${derivativeType}) -> ${returnType} {
var result : ${resultType};
${body}
return ${returnType}(result);
}
${stageWGSL}
`;
const pipelines =
s_deviceToPipelines.get(t.device) ?? new Map<string, GPURenderPipeline | GPUComputePipeline>();
s_deviceToPipelines.set(t.device, pipelines);
const type = getTextureFormatType(format ?? 'rgba8unorm');
const isFiltering =
!!sampler &&
(sampler.minFilter === 'linear' ||
sampler.magFilter === 'linear' ||
sampler.mipmapFilter === 'linear');
let sampleType: GPUTextureSampleType = isSingleChannelInput(textureType)
? 'depth'
: isDepthTextureFormat(format)
? 'unfilterable-float'
: isStencilTextureFormat(format)
? 'uint'
: type ?? 'float';
if (isFiltering && sampleType === 'unfilterable-float') {
assert(isTextureFormatFilterable(t.device.features, format));
sampleType = 'float';
}
if (sampleCount > 1 && sampleType === 'float') {
sampleType = 'unfilterable-float';
}
const visibility =
stage === 'compute'
? GPUShaderStage.COMPUTE
: stage === 'fragment'
? GPUShaderStage.FRAGMENT
: GPUShaderStage.VERTEX;
const entries: GPUBindGroupLayoutEntry[] = [
{
binding: 2,
visibility,
buffer: {
type: 'uniform',
},
},
];
const viewDimension = effectiveViewDimensionForDimension(
viewDescriptor.dimension,
dimension,
depthOrArrayLayers
);
if (textureType.includes('storage')) {
entries.push({
binding: 0,
visibility,
storageTexture: {
access: 'read-only',
viewDimension,
format,
},
});
} else if (textureType === 'texture_external') {
entries.push({
binding: 0,
visibility,
externalTexture: {},
});
} else {
entries.push({
binding: 0,
visibility,
texture: {
sampleType,
viewDimension,
multisampled: sampleCount > 1,
},
});
}
if (sampler) {
const type = isCompare ? 'comparison' : isFiltering ? 'filtering' : 'non-filtering';
entries.push({
binding: 1,
visibility,
sampler: { type },
});
}
const id = `${resultType}:${stage}:${JSON.stringify(entries)}:${code}`;
let pipeline = pipelines.get(id);
if (!pipeline) {
const module = t.device.createShaderModule({ code });
const bindGroupLayout0 = t.device.createBindGroupLayout({ entries });
const bindGroupLayouts = [bindGroupLayout0];
if (stage === 'compute') {
const bindGroupLayout1 = t.device.createBindGroupLayout({
entries: [
{
binding: 0,
visibility: GPUShaderStage.COMPUTE,
buffer: {
type: 'storage',
},
},
],
});
bindGroupLayouts.push(bindGroupLayout1);
}
const layout = t.device.createPipelineLayout({
bindGroupLayouts,
});
switch (stage) {
case 'compute':
pipeline = t.device.createComputePipeline({
layout,
compute: { module },
});
break;
case 'fragment':
case 'vertex':
pipeline = t.device.createRenderPipeline({
layout,
vertex: { module },
fragment: {
module,
targets: [{ format: 'rgba32uint' }],
},
});
break;
}
pipelines.set(id, pipeline);
}
const gpuSampler = sampler ? t.device.createSampler(sampler) : undefined;
const run = async (gpuTexture: GPUTexture | GPUExternalTexture) => {
const resultBuffer = t.createBufferTracked({
label: 'createTextureCallsRunner:resultBuffer',
size: align(calls.length * 16, 256),
usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ,
});
const aspect = getAspectForTexture(gpuTexture);
const runViewDescriptor = {
...viewDescriptor,
aspect,
};
const bindGroup0 = t.device.createBindGroup({
layout: pipeline!.getBindGroupLayout(0),
entries: [
{
binding: 0,
resource:
gpuTexture instanceof GPUExternalTexture
? gpuTexture
: gpuTexture.createView(runViewDescriptor),
},
...(sampler ? [{ binding: 1, resource: gpuSampler! }] : []),
{ binding: 2, resource: { buffer: dataBuffer } },
],
});
let storageBuffer: GPUBuffer | undefined;
const encoder = t.device.createCommandEncoder({ label: 'createTextureCallsRunner' });
if (stage === 'compute') {
storageBuffer = t.createBufferTracked({
label: 'createTextureCallsRunner:storageBuffer',
size: resultBuffer.size,
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
});
const bindGroup1 = t.device.createBindGroup({
layout: pipeline!.getBindGroupLayout(1),
entries: [{ binding: 0, resource: { buffer: storageBuffer } }],
});
const pass = encoder.beginComputePass();
pass.setPipeline(pipeline! as GPUComputePipeline);
pass.setBindGroup(0, bindGroup0);
pass.setBindGroup(1, bindGroup1);
pass.dispatchWorkgroups(calls.length);
pass.end();
encoder.copyBufferToBuffer(storageBuffer, 0, resultBuffer, 0, storageBuffer.size);
} else {
const pass = encoder.beginRenderPass({
colorAttachments: [
{
view: renderTarget.createView(),
loadOp: 'clear',
storeOp: 'store',
},
],
});
pass.setPipeline(pipeline! as GPURenderPipeline);
pass.setBindGroup(0, bindGroup0);
for (let i = 0; i < calls.length; ++i) {
pass.setViewport(i, 0, 1, 1, 0, 1);
pass.draw(3, 1, 0, i);
}
pass.end();
encoder.copyTextureToBuffer(
{ texture: renderTarget },
{
buffer: resultBuffer,
bytesPerRow: resultBuffer.size,
},
[renderTarget.width, 1]
);
}
t.device.queue.submit([encoder.finish()]);
await resultBuffer.mapAsync(GPUMapMode.READ);
const view = TexelView.fromTextureDataByReference(
resultFormat,
new Uint8Array(resultBuffer.getMappedRange()),
{
bytesPerRow: calls.length * 16,
rowsPerImage: 1,
subrectOrigin: [0, 0, 0],
subrectSize: [calls.length, 1],
}
);
let outIdx = 0;
const out = new Array<PerTexelComponent<number>>(calls.length);
for (const bin of binned) {
for (const callIdx of bin) {
const x = outIdx;
out[callIdx] = view.color({ x, y: 0, z: 0 });
outIdx++;
}
}
storageBuffer?.destroy();
resultBuffer.destroy();
return out;
};
return {
run,
destroy() {
dataBuffer.destroy();
renderTarget.destroy();
},
};
}
export async function doTextureCalls<T extends Dimensionality>(
t: GPUTest,
gpuTexture: GPUTexture | GPUExternalTexture,
viewDescriptor: GPUTextureViewDescriptor,
textureType: string,
sampler: GPUSamplerDescriptor | undefined,
calls: TextureCall<T>[],
shortShaderStage: ShortShaderStage
) {
const stage = kShortShaderStageToShaderStage[shortShaderStage];
const runner = createTextureCallsRunner(
t,
gpuTexture instanceof GPUExternalTexture
? { format: 'rgba8unorm', dimension: '2d', depthOrArrayLayers: 1, sampleCount: 1 }
: gpuTexture,
viewDescriptor,
textureType,
sampler,
calls,
stage
);
const results = await runner.run(gpuTexture);
return {
runner,
results,
};
}
function texelFormat(texel: Readonly<PerTexelComponent<number>>, rep: TexelRepresentationInfo) {
return rep.componentOrder.map(component => `${component}: ${texel[component]}`).join(', ');
}
function compareTexelViewsImpl(
device: GPUDevice,
{
actualTexelView,
expectedTexelView,
origin,
size,
sampleCount,
}: {
actualTexelView: TexelView;
expectedTexelView: TexelView;
origin: Required<GPUOrigin3DDict>;
size: Required<GPUExtent3DDict>;
sampleCount: number;
}
) {
const errors: string[] = [];
const actualRep = kTexelRepresentationInfo[actualTexelView.format];
const expectedRep = kTexelRepresentationInfo[expectedTexelView.format];
for (let z = origin.z; z < size.depthOrArrayLayers; ++z) {
for (let y = origin.y; y < size.height; ++y) {
for (let x = origin.x; x < size.width; ++x) {
for (let sampleIndex = 0; sampleIndex < sampleCount; ++sampleIndex) {
const actual = actualTexelView.color({ x, y, z, sampleIndex });
const expected = expectedTexelView.color({ x, y, z, sampleIndex });
const actualRGBA = convertPerTexelComponentToResultFormat(actual, actualTexelView.format);
const expectedRGBA = convertPerTexelComponentToResultFormat(
expected,
expectedTexelView.format
);
const maxFractionalDiff = 0;
if (
!texelsApproximatelyEqual(
device,
'textureLoad',
'unused',
actualRGBA,
actualTexelView.format,
expectedRGBA,
expectedTexelView.format,
maxFractionalDiff
)
) {
const actualStr = texelFormat(actual, actualRep);
const expectedStr = texelFormat(expected, expectedRep);
errors.push(
`texel at ${x}, ${y}, ${z}, sampleIndex: ${sampleIndex} expected: ${expectedStr}, actual: ${actualStr}`
);
}
}
}
}
}
return errors;
}
* Compares 1 mip level of two arrays of TexelViews where each element in the array is the next mip level.
*/
export function compareTexelViewsMipLevel(
device: GPUDevice,
{
actualTexelViews,
expectedTexelViews,
mipLevel = 0,
origin,
size,
sampleCount = 1,
}: {
actualTexelViews: TexelView[];
expectedTexelViews: TexelView[];
origin?: GPUOrigin3D;
size: GPUExtent3D;
mipLevel?: number;
sampleCount?: number;
}
) {
const errors = compareTexelViewsImpl(device, {
actualTexelView: actualTexelViews[mipLevel ?? 0],
expectedTexelView: expectedTexelViews[mipLevel ?? 0],
origin: reifyOrigin3D(origin ?? [0]),
size: reifyExtent3D(size),
sampleCount: sampleCount ?? 1,
});
if (errors.length > 0) {
errors.unshift(`errors in mipLevel: ${mipLevel}`);
}
return errors;
}
* Compares two arrays of TexelViews where each element in the array is the next mip level.
*/
export function compareTexelViews(
device: GPUDevice,
{
actualTexelViews,
expectedTexelViews,
size,
dimension = '2d',
sampleCount = 1,
}: {
actualTexelViews: TexelView[];
expectedTexelViews: TexelView[];
dimension?: GPUTextureDimension;
size: GPUExtent3D;
sampleCount?: number;
}
) {
assert(actualTexelViews.length === expectedTexelViews.length);
const errors: string[] = [];
for (let mipLevel = 0; mipLevel < actualTexelViews.length; ++mipLevel) {
errors.push(
...compareTexelViewsMipLevel(device, {
actualTexelViews,
expectedTexelViews,
size: virtualMipSize(dimension, size, mipLevel),
mipLevel,
sampleCount,
})
);
}
return errors;
}
