diff --git a/.eslintrc.cjs b/.eslintrc.cjs index c2f7d3b8c633..ec939e9ea56c 100644 --- a/.eslintrc.cjs +++ b/.eslintrc.cjs @@ -14,6 +14,7 @@ const kKnownWGSLLanguageFeatures = [ 'immediate_address_space', 'fragment_depth', 'buffer_view', + 'unrestricted_aliasing', // IMPORTANT: Always add features to capability_info.ts before adding them here. // This ensures that they're tested properly. ]; diff --git a/src/webgpu/capability_info.ts b/src/webgpu/capability_info.ts index 5bcf21086c99..b0b3debe47d4 100644 --- a/src/webgpu/capability_info.ts +++ b/src/webgpu/capability_info.ts @@ -994,6 +994,7 @@ export const kKnownWGSLLanguageFeatures = [ 'immediate_address_space', 'fragment_depth', 'buffer_view', + 'unrestricted_aliasing', // This list should be kept in sync with .eslintrc.cjs. ] as const; diff --git a/src/webgpu/shader/execution/expression/call/builtin/bufferArrayView.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/bufferArrayView.spec.ts index 992ca756e352..c3ed7e91d9e4 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/bufferArrayView.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/bufferArrayView.spec.ts @@ -7,6 +7,17 @@ import { keysOf } from '../../../../../../common/util/data_tables.js'; import { assert } from '../../../../../../common/util/util.js'; import { AllFeaturesMaxLimitsGPUTest } from '../../../../../gpu_test.js'; import { Type } from '../../../../../util/conversion.js'; +import { + kMixedTypeBuffers, + kMixedTypeIdx, + kMixedTypeOps, + kMixedTypeOverlaps, + kMixedTypePairs, + MixedType, + mixedTypeName, + mixedTypeWords, + runMixedTypeAliasingTest, +} from '../mixed_type_aliasing_utils.js'; import { kBufferSizes, @@ -22,15 +33,6 @@ import { runReadLayoutTest, runWriteLayoutTest, runReadWriteTest, - kMixedTypeOps, - kMixedTypeIdx, - kMixedTypePairs, - kMixedTypeOverlaps, - kMixedTypeBuffers, - MixedType, - mixedTypeName, - mixedTypeWords, - runMixedTypeAliasingTest, } from './buffer_view_utils.js'; export const g = makeTestGroup(AllFeaturesMaxLimitsGPUTest); diff --git a/src/webgpu/shader/execution/expression/call/builtin/bufferView.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/bufferView.spec.ts index 97c2bfcca459..f4240c555ad1 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/bufferView.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/bufferView.spec.ts @@ -7,6 +7,17 @@ import { keysOf } from '../../../../../../common/util/data_tables.js'; import { assert } from '../../../../../../common/util/util.js'; import { AllFeaturesMaxLimitsGPUTest } from '../../../../../gpu_test.js'; import { Type } from '../../../../../util/conversion.js'; +import { + kMixedTypeBuffers, + kMixedTypeIdx, + kMixedTypeOps, + kMixedTypeOverlaps, + kMixedTypePairs, + MixedType, + mixedTypeName, + mixedTypeWords, + runMixedTypeAliasingTest, +} from '../mixed_type_aliasing_utils.js'; import { kBufferSizes, @@ -21,15 +32,6 @@ import { runReadLayoutTest, runWriteLayoutTest, runReadWriteTest, - kMixedTypeOps, - kMixedTypeIdx, - kMixedTypePairs, - kMixedTypeOverlaps, - kMixedTypeBuffers, - MixedType, - mixedTypeName, - mixedTypeWords, - runMixedTypeAliasingTest, } from './buffer_view_utils.js'; export const g = makeTestGroup(AllFeaturesMaxLimitsGPUTest); diff --git a/src/webgpu/shader/execution/expression/call/builtin/buffer_view_utils.ts b/src/webgpu/shader/execution/expression/call/builtin/buffer_view_utils.ts index 6635f19e992f..8622b98a1f8d 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/buffer_view_utils.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/buffer_view_utils.ts @@ -1,5 +1,4 @@ import { - assert, iterRange, typedArrayParam, typedArrayFromParam, @@ -12,7 +11,6 @@ import { ArrayType, MatrixType, VectorType, - float32ToFloat16Bits, } from '../../../../../util/conversion.js'; export const kBufferSizes = [128, 256, 512, 1024] as const; @@ -975,665 +973,3 @@ export function runReadWriteTest( t.expectGPUBufferValuesEqual(outputBuffer, outputData); } - -/** @returns the bit pattern of the f32 value 'f' as an i32. */ -export function f32Bits(f: number): number { - return new Int32Array(new Float32Array([f]).buffer)[0]; -} - -/** The scalar types that may be used by the lanes of a mixed type. */ -export type MixedScalar = 'i32' | 'u32' | 'f32' | 'f16'; - -/** @returns the size in bytes of the scalar type 's'. */ -function mixedScalarSize(s: MixedScalar): number { - return s === 'f16' ? 2 : 4; -} - -/** A scalar lane of a mixed type. */ -interface MixedLane { - /** The scalar type of the lane. */ - elem: MixedScalar; - /** The byte offset of the lane from the start of the value. */ - offset: number; - /** The WGSL accessor suffix for the lane. e.g. '', '[1]', '.a' or '.v[1]' */ - access: string; -} - -/** - * A scalar, vector or structure type used by the mixed type aliasing tests. - * A value of the type is treated as a flat list of scalar lanes. Structures may contain padding, - * which is not part of any lane. - */ -export interface MixedType { - /** The WGSL type name. */ - name: string; - /** The size in bytes. This is also the array element stride. */ - size: number; - /** The alignment in bytes. */ - align: number; - /** The scalar lanes, in declaration order. */ - lanes: MixedLane[]; - /** Present if the type is a structure. */ - struct?: { - /** The structure members, with their byte offsets. */ - members: { name: string; type: MixedType; offset: number }[]; - }; -} - -function scalar(elem: MixedScalar): MixedType { - const size = mixedScalarSize(elem); - return { name: elem, size, align: size, lanes: [{ elem, offset: 0, access: '' }] }; -} - -function vec(n: 2 | 4, elem: MixedScalar): MixedType { - const size = n * mixedScalarSize(elem); - const lanes = [...Array(n).keys()].map(k => ({ - elem, - offset: k * mixedScalarSize(elem), - access: `[${k}]`, - })); - return { name: `vec${n}<${elem}>`, size, align: size, lanes }; -} - -/** @returns a structure type, laid out with the WGSL layout rules. */ -function structType(name: string, members: [string, MixedType][]): MixedType { - const roundUp = (n: number, k: number) => Math.ceil(n / k) * k; - let end = 0; - let align = 1; - const laidOut = members.map(([memberName, type]) => { - const offset = roundUp(end, type.align); - end = offset + type.size; - align = Math.max(align, type.align); - return { name: memberName, type, offset }; - }); - const lanes = laidOut.flatMap(m => - m.type.lanes.map(l => ({ - elem: l.elem, - offset: m.offset + l.offset, - access: `.${m.name}${l.access}`, - })) - ); - return { name, size: roundUp(end, align), align, lanes, struct: { members: laidOut } }; -} - -const kI32 = scalar('i32'); -const kU32 = scalar('u32'); -const kF32 = scalar('f32'); -const kF16 = scalar('f16'); -const kVec2I = vec(2, 'i32'); -const kVec4I = vec(4, 'i32'); -const kVec4U = vec(4, 'u32'); -const kVec2F = vec(2, 'f32'); -const kVec4F = vec(4, 'f32'); -const kVec2H = vec(2, 'f16'); -const kVec4H = vec(4, 'f16'); -// Structures without padding. -const kStructI = structType('SI', [ - ['a', kI32], - ['b', kI32], - ['c', kI32], - ['d', kI32], -]); -const kStructF = structType('SF', [ - ['a', kF32], - ['b', kF32], - ['c', kF32], - ['d', kF32], -]); -const kStructV = structType('SV', [ - ['v', kVec2I], - ['s', kI32], - ['t', kI32], -]); -// Structures with padding. -/** 4 bytes of padding between 'a' and 'b'. */ -const kPaddedI = structType('PI', [ - ['a', kI32], - ['b', kVec2I], -]); -/** 4 bytes of padding between 'a' and 'b'. */ -const kPaddedF = structType('PF', [ - ['a', kF32], - ['b', kVec2F], -]); -/** 4 bytes of trailing padding. */ -const kPaddedTrailingF = structType('PT', [ - ['v', kVec2F], - ['s', kF32], -]); -/** 2 bytes of padding between 'a' and 'b', within the first 4-byte word. */ -const kPaddedH = structType('PH', [ - ['a', kF16], - ['b', kF32], -]); - -/** @returns true if the type 't' has any f16 lanes. */ -export function mixedTypeUsesF16(t: MixedType): boolean { - return t.lanes.some(l => l.elem === 'f16'); -} - -/** @returns the WGSL name of the type 't'. */ -export function mixedTypeName(t: MixedType): string { - return t.name; -} - -/** @returns a WGSL constructor expression for a value of type 't', given an expression per lane. */ -function mixedTypeCtorFromLanes(t: MixedType, lanes: string[]): string { - if (t.struct) { - let next = 0; - const members = t.struct.members.map(m => { - const n = m.type.lanes.length; - const expr = mixedTypeCtorFromLanes(m.type, lanes.slice(next, next + n)); - next += n; - return expr; - }); - return `${t.name}(${members.join(', ')})`; - } - return t.lanes.length === 1 ? lanes[0] : `${t.name}(${lanes.join(', ')})`; -} - -/** - * @returns a WGSL constructor expression for a value of type 't', where lane 'k' has the value - * 'base + k', and 'base' is a WGSL scalar expression. - */ -function mixedTypeCtor(t: MixedType, base: string): string { - const lanes = t.lanes.map(({ elem }, k) => - k === 0 ? `${elem}(${base})` : `${elem}(${base}) + ${elem}(${k})` - ); - return mixedTypeCtorFromLanes(t, lanes); -} - -/** @returns the WGSL declarations of the structures used by the types 'types', including nested. */ -function mixedTypeStructDecls(types: MixedType[]): string { - const decls = new Map(); - const visit = (t: MixedType) => { - if (!t.struct || decls.has(t.name)) { - return; - } - t.struct.members.forEach(m => visit(m.type)); - const members = t.struct.members.map(m => `${m.name} : ${m.type.name}`).join(', '); - decls.set(t.name, `struct ${t.name} { ${members} }`); - }; - types.forEach(visit); - return [...decls.values()].join('\n'); -} - -/** - * @returns WGSL declarations required by the mixed type ops, where 'int' is the type of the loaded - * view, and 'other' is the type of the stored view. This declares the structures used by the types - * and the functions: - * 'mixed_sub(a, b)' - lane-wise 'a - b' - * 'mixed_xor(a, b)' - lane-wise 'a ^ b' - * If either type uses f16, the shader must also 'enable f16;'. - */ -export function mixedTypeDecls(int: MixedType, other: MixedType): string { - const ty = mixedTypeName(int); - const laneWise = (name: string, op: string) => { - if (!int.struct) { - return `fn ${name}(a : ${ty}, b : ${ty}) -> ${ty} { return a ${op} b; }`; - } - const lanes = int.lanes.map(({ access: l }) => ` r${l} = a${l} ${op} b${l};`); - return `fn ${name}(a : ${ty}, b : ${ty}) -> ${ty} {\n var r : ${ty};\n${lanes.join( - '\n' - )}\n return r;\n}`; - }; - return `${mixedTypeStructDecls([int, other])} -${laneWise('mixed_sub', '-')} -${laneWise('mixed_xor', '^')} -`; -} - -/** - * @returns WGSL statements that write each lane of the value 'value' of type 't', as an i32 bit - * pattern, to consecutive elements of the i32 array 'array'. All lanes of 't' must be 32-bit. - */ -export function mixedTypeWriteLanes(t: MixedType, value: string, array: string): string { - return t.lanes - .map(({ access }, k) => `${array}[${k}] = bitcast(${value}${access});`) - .join('\n '); -} - -/** - * The pairs of view types tested. - * 'int' is a type with only i32 or u32 lanes, and at most 4 lanes. It is the only view that is - * loaded from, which keeps the expected results exact: float values are never loaded, so denormal - * flushing and NaN canonicalization cannot affect the results. - * 'other' is a type that is only stored to. - */ -export const kMixedTypePairs: Record = { - // scalar - scalar - i32_f32: { int: kI32, other: kF32 }, - u32_f32: { int: kU32, other: kF32 }, - i32_i32: { int: kI32, other: kI32 }, - // vector - scalar - vec2i_f32: { int: kVec2I, other: kF32 }, - vec4i_f32: { int: kVec4I, other: kF32 }, - vec4u_f32: { int: kVec4U, other: kF32 }, - vec4u_u32: { int: kVec4U, other: kU32 }, - // scalar - vector - i32_vec2f: { int: kI32, other: kVec2F }, - i32_vec4f: { int: kI32, other: kVec4F }, - i32_vec4u: { int: kI32, other: kVec4U }, - i32_vec4i: { int: kI32, other: kVec4I }, - // vector - vector - vec2i_vec2f: { int: kVec2I, other: kVec2F }, - vec4i_vec4f: { int: kVec4I, other: kVec4F }, - vec4u_vec4f: { int: kVec4U, other: kVec4F }, - vec4i_vec2f: { int: kVec4I, other: kVec2F }, - vec2i_vec4f: { int: kVec2I, other: kVec4F }, - vec4i_vec4u: { int: kVec4I, other: kVec4U }, - vec4i_vec4i: { int: kVec4I, other: kVec4I }, - // struct - scalar - structi_f32: { int: kStructI, other: kF32 }, - structv_f32: { int: kStructV, other: kF32 }, - i32_structf: { int: kI32, other: kStructF }, - // struct - vector - structi_vec4f: { int: kStructI, other: kVec4F }, - structi_vec2f: { int: kStructI, other: kVec2F }, - vec4i_structf: { int: kVec4I, other: kStructF }, - // struct - struct - structi_structf: { int: kStructI, other: kStructF }, - structv_structf: { int: kStructV, other: kStructF }, - // f16. The f16 stores only write part of a 4-byte word. - i32_f16: { int: kI32, other: kF16 }, - u32_f16: { int: kU32, other: kF16 }, - i32_vec2h: { int: kI32, other: kVec2H }, - i32_vec4h: { int: kI32, other: kVec4H }, - vec4i_f16: { int: kVec4I, other: kF16 }, - vec2i_vec4h: { int: kVec2I, other: kVec4H }, - vec4i_vec4h: { int: kVec4I, other: kVec4H }, - // padded struct. Stores of padded structures must not write the padding. - paddedi_f32: { int: kPaddedI, other: kF32 }, - paddedi_vec4f: { int: kPaddedI, other: kVec4F }, - i32_paddedf: { int: kI32, other: kPaddedF }, - vec4i_paddedf: { int: kVec4I, other: kPaddedF }, - i32_paddedtf: { int: kI32, other: kPaddedTrailingF }, - vec4i_paddedtf: { int: kVec4I, other: kPaddedTrailingF }, - paddedi_paddedf: { int: kPaddedI, other: kPaddedF }, - paddedi_structf: { int: kPaddedI, other: kStructF }, - structi_paddedf: { int: kStructI, other: kPaddedF }, - // padded struct with f16 - i32_paddedh: { int: kI32, other: kPaddedH }, - vec4i_paddedh: { int: kVec4I, other: kPaddedH }, - paddedi_paddedh: { int: kPaddedI, other: kPaddedH }, -}; - -/** - * How the two views overlap. Views always start on a 4-byte boundary. - * 'same_start' - Both views start at the same byte. - * 'offset' - The views start at different bytes, and a lane of one view overlaps a lane of - * the other. Not all pairs of types can do this while respecting alignment. - * 'padding' - The views overlap, but only lanes of one view overlap padding of the other. - * Only possible for structures with padding. - */ -export const kMixedTypeOverlaps = ['same_start', 'offset', 'padding'] as const; -export type MixedTypeOverlap = (typeof kMixedTypeOverlaps)[number]; - -/** - * @returns the word offsets of the 'int' and 'other' views for the given overlap, or undefined if - * the types cannot be positioned to satisfy the overlap. Word offsets are always at least 4. - */ -export function mixedTypeWords( - int: MixedType, - other: MixedType, - overlap: MixedTypeOverlap -): { intWord: number; otherWord: number } | undefined { - const lanesOverlap = (a: number, b: number) => - int.lanes.some(li => - other.lanes.some(lo => { - const i0 = a * 4 + li.offset; - const o0 = b * 4 + lo.offset; - return i0 < o0 + mixedScalarSize(lo.elem) && o0 < i0 + mixedScalarSize(li.elem); - }) - ); - const viewsOverlap = (a: number, b: number) => - a * 4 < b * 4 + other.size && b * 4 < a * 4 + int.size; - for (let a = 4; a < 12; a++) { - if ((a * 4) % int.align !== 0) continue; - for (let b = 4; b < 12; b++) { - if ((b * 4) % other.align !== 0) continue; - let ok = false; - switch (overlap) { - case 'same_start': - ok = a === b; - break; - case 'offset': - ok = a !== b && lanesOverlap(a, b); - break; - case 'padding': - ok = viewsOverlap(a, b) && !lanesOverlap(a, b); - break; - } - if (ok) { - return { intWord: a, otherWord: b }; - } - } - } - return undefined; -} - -/** A simulated memory of 4-byte words, used to calculate the expected results of mixed type tests. */ -export interface MixedTypeMemory { - ld(ptr: number): number; - st(ptr: number, value: number): void; -} - -/** The runtime arguments passed to a mixed type op. */ -export interface MixedTypeArgs { - /** A value to write through the 'int' view. Also used as a second value for the 'other' view. */ - va: number; - /** A value to write through the 'other' view. */ - vf: number; - /** A loop count. */ - n: number; -} - -/** - * Stores a value of type 't' where lane 'k' is 'base + k', at word 'ptr'. - * Only the bytes of the lanes are written. Padding is left untouched. - */ -function simStore(m: MixedTypeMemory, t: MixedType, ptr: number, base: number) { - t.lanes.forEach(({ elem, offset }, k) => { - const v = base + k; - const word = ptr + Math.floor(offset / 4); - switch (elem) { - case 'i32': - case 'u32': - m.st(word, v | 0); - break; - case 'f32': - m.st(word, f32Bits(v)); - break; - case 'f16': { - const shift = (offset % 4) * 8; - const mask = 0xffff << shift; - m.st(word, (m.ld(word) & ~mask) | ((float32ToFloat16Bits(v) << shift) & mask)); - break; - } - } - }); -} - -/** @returns the lanes of a value of type 't' at word 'ptr', as i32 bit patterns. */ -function simLoad(m: MixedTypeMemory, t: MixedType, ptr: number): number[] { - return t.lanes.map(({ elem, offset }) => { - assert(elem === 'i32' || elem === 'u32', 'only integer views are loaded'); - return m.ld(ptr + offset / 4); - }); -} - -/** - * An operation on a pointer 'pi' to an integer type, and a pointer 'pf' to another type, that may - * refer to overlapping memory. Shaders using these ops must include mixedTypeDecls(int, other). - */ -export interface MixedTypeOp { - /** - * @returns a WGSL function body that returns a value of type 'int', and may use: - * 'pi' - a pointer to 'int' - * 'pf' - a pointer to 'other' - * 'va' - an i32 value - * 'vf' - an f32 value, whose value is exactly representable as an f16 - * 'n' - an i32 loop count - */ - wgsl: (int: MixedType, other: MixedType) => string; - /** - * Simulates 'wgsl', where 'pi' and 'pf' are word offsets into 'm'. - * @returns the lanes of the returned value, as i32 bit patterns. - */ - sim: ( - m: MixedTypeMemory, - int: MixedType, - other: MixedType, - pi: number, - pf: number, - args: MixedTypeArgs - ) => number[]; -} - -/** - * Operations that produce different results if an implementation assumes that differently typed - * pointers do not alias (i.e. applies C/C++ style type-based alias analysis). - */ -export const kMixedTypeOps: Record = { - int_store_other_store_int_load: { - wgsl: (int, other) => ` - *pi = ${mixedTypeCtor(int, 'va')}; - *pf = ${mixedTypeCtor(other, 'vf')}; - return *pi;`, - sim: (m, int, other, pi, pf, { va, vf }) => { - simStore(m, int, pi, va); - simStore(m, other, pf, vf); - return simLoad(m, int, pi); - }, - }, - other_store_int_load: { - wgsl: (int, other) => ` - *pf = ${mixedTypeCtor(other, 'vf')}; - return *pi;`, - sim: (m, int, other, pi, pf, { vf }) => { - simStore(m, other, pf, vf); - return simLoad(m, int, pi); - }, - }, - int_load_other_store_int_load: { - wgsl: (int, other) => ` - let before = *pi; - *pf = ${mixedTypeCtor(other, 'vf')}; - return mixed_sub(*pi, before);`, - sim: (m, int, other, pi, pf, { vf }) => { - const before = simLoad(m, int, pi); - simStore(m, other, pf, vf); - return simLoad(m, int, pi).map((v, k) => (v - before[k]) | 0); - }, - }, - dead_other_store: { - // The first store to '*pf' is only dead if '*pi' does not alias '*pf'. - wgsl: (int, other) => ` - *pf = ${mixedTypeCtor(other, 'vf')}; - let tmp = *pi; - *pf = ${mixedTypeCtor(other, 'va')}; - return tmp;`, - sim: (m, int, other, pi, pf, { va, vf }) => { - simStore(m, other, pf, vf); - const tmp = simLoad(m, int, pi); - simStore(m, other, pf, va); - return tmp; - }, - }, - loop_other_store_int_load: { - // The load of '*pi' must not be hoisted out of the loop. - wgsl: (int, other) => ` - var acc = ${mixedTypeName(int)}(); - for (var i = 0; i < n; i++) { - *pf = ${mixedTypeCtor(other, 'i')}; - acc = mixed_xor(acc, *pi); - } - return acc;`, - sim: (m, int, other, pi, pf, { n }) => { - const acc = new Array(int.lanes.length).fill(0); - for (let i = 0; i < n; i++) { - simStore(m, other, pf, i); - simLoad(m, int, pi).forEach((v, k) => { - acc[k] ^= v; - }); - } - return acc; - }, - }, -}; - -/** The number of 4-byte words in each of the buffers 'A' and 'B' used by runMixedTypeAliasingTest. */ -export const kMixedTypeBufferWords = 32; - -/** - * The runtime parameters passed to the shader of runMixedTypeAliasingTest in 'input.p'. - * p[0]: 'va' - an i32 value to write. - * p[1]: 'vf' - an f32 value to write, passed as an i32 and converted to f32. - * p[2]: 'n' - a loop count. - * p[3]: 'idx' - a runtime value that views may use to form offsets and indices. - */ -export const kMixedTypeParams = [1000, 2000, 4, 1] as const; - -/** The value of 'input.p[3]', which may be used by views to form offsets and indices. */ -export const kMixedTypeIdx = kMixedTypeParams[3]; - -/** - * The kinds of buffer that the views are formed on: - * 'workgroup' - var of type buffer - * 'storage' - var of type buffer - * 'storage_unsized' - var of type buffer - */ -export const kMixedTypeBuffers = ['workgroup', 'storage', 'storage_unsized'] as const; -export type MixedTypeBuffer = (typeof kMixedTypeBuffers)[number]; - -interface MixedTypeAliasingParams { - /** The kind of buffer the views are formed on. */ - buffer: MixedTypeBuffer; - /** The type of the view that is loaded from. */ - int: MixedType; - /** The type of the view that is only stored to. */ - other: MixedType; - /** The word offset of the 'int' view. */ - intWord: number; - /** The word offset of the 'other' view. */ - otherWord: number; - /** - * @returns a WGSL expression for a pointer to a value of type 'type' at word offset 'word' of the - * buffer variable named 'buffer'. May use 'input.p[3]' as a runtime value. - */ - view: (type: MixedType, buffer: string, word: number) => string; - /** - * If true, the 'other' view is formed in the same buffer as the 'int' view. - * Otherwise it is formed in a different buffer. - */ - aliased: boolean; - /** The operation to perform. */ - op: MixedTypeOp; -} - -/** - * Runs a test where two differently typed pointers are formed from buffer views and used within a - * single function. If 'aliased' is true, then both views refer to the same buffer, otherwise the - * 'other' view refers to a different buffer. - * - * The shader initializes buffers 'A' and 'B' from 'input', calls the op, and writes the lanes of - * the result of the op, followed by the contents of 'A' and 'B', to 'output'. - */ -export function runMixedTypeAliasingTest(t: GPUTest, params: MixedTypeAliasingParams) { - t.skipIfLanguageFeatureNotSupported('buffer_view'); - const usesF16 = mixedTypeUsesF16(params.int) || mixedTypeUsesF16(params.other); - if (usesF16) { - t.skipIfDeviceDoesNotHaveFeature('shader-f16'); - } - - const N = kMixedTypeBufferWords; - let decls = ''; - switch (params.buffer) { - case 'workgroup': - decls = `var A : buffer<${N * 4}>;\nvar B : buffer<${N * 4}>;`; - break; - case 'storage': - decls = `@group(0) @binding(2) var A : buffer<${N * 4}>; -@group(0) @binding(3) var B : buffer<${N * 4}>;`; - break; - case 'storage_unsized': - decls = `@group(0) @binding(2) var A : buffer; -@group(0) @binding(3) var B : buffer;`; - break; - } - - const intTy = mixedTypeName(params.int); - const writeResult = mixedTypeWriteLanes(params.int, 'r', 'output.r'); - - const wgsl = `${usesF16 ? 'enable f16;' : ''} -struct In { - a : array, - b : array, - p : array, -} - -struct Out { - r : array, - a : array, - b : array, -} - -@group(0) @binding(0) var input : In; -@group(0) @binding(1) var output : Out; - -${decls} - -${mixedTypeDecls(params.int, params.other)} - -fn f(va : i32, vf : f32, n : i32) -> ${intTy} { - let pi = ${params.view(params.int, 'A', params.intWord)}; - let pf = ${params.view(params.other, params.aliased ? 'A' : 'B', params.otherWord)}; - ${params.op.wgsl(params.int, params.other)} -} - -@compute @workgroup_size(1) -fn main() { - *bufferView>(&A, 0) = input.a; - *bufferView>(&B, 0) = input.b; - let r = f(input.p[0], f32(input.p[1]), input.p[2]); - ${writeResult} - output.a = *bufferView>(&A, 0); - output.b = *bufferView>(&B, 0); -} -`; - - // Simulate the op. 'A' is at word 0, and 'B' is at word N. - const initA = Array.from({ length: N }, (_, i) => 10 + i); - const initB = Array.from({ length: N }, (_, i) => 100 + i); - const mem = [...initA, ...initB]; - const memory: MixedTypeMemory = { - ld: ptr => mem[ptr], - st: (ptr, value) => { - mem[ptr] = value | 0; - }, - }; - const [va, vf, n] = kMixedTypeParams; - const pi = params.intWord; - const pf = (params.aliased ? 0 : N) + params.otherWord; - const r = params.op.sim(memory, params.int, params.other, pi, pf, { va, vf, n }); - const rOut = [0, 0, 0, 0]; - r.forEach((v, k) => { - rOut[k] = v; - }); - const expected = new Int32Array([...rOut, ...mem]); - - const pipeline = t.device.createComputePipeline({ - layout: 'auto', - compute: { module: t.device.createShaderModule({ code: wgsl }) }, - }); - - const inputBuffer = t.makeBufferWithContents( - new Int32Array([...initA, ...initB, ...kMixedTypeParams]), - GPUBufferUsage.STORAGE - ); - const outputBuffer = t.createBufferTracked({ - size: expected.byteLength, - usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, - }); - const entries: GPUBindGroupEntry[] = [ - { binding: 0, resource: { buffer: inputBuffer } }, - { binding: 1, resource: { buffer: outputBuffer } }, - ]; - if (params.buffer !== 'workgroup') { - for (const binding of [2, 3]) { - const buffer = t.createBufferTracked({ size: N * 4, usage: GPUBufferUsage.STORAGE }); - entries.push({ binding, resource: { buffer } }); - } - } - const bg = t.device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries }); - - const encoder = t.device.createCommandEncoder(); - const pass = encoder.beginComputePass(); - pass.setPipeline(pipeline); - pass.setBindGroup(0, bg); - pass.dispatchWorkgroups(1); - pass.end(); - t.queue.submit([encoder.finish()]); - - t.expectGPUBufferValuesEqual(outputBuffer, expected); -} diff --git a/src/webgpu/shader/execution/expression/call/mixed_type_aliasing_utils.ts b/src/webgpu/shader/execution/expression/call/mixed_type_aliasing_utils.ts new file mode 100644 index 000000000000..5264b2ab8353 --- /dev/null +++ b/src/webgpu/shader/execution/expression/call/mixed_type_aliasing_utils.ts @@ -0,0 +1,698 @@ +/** + * Utilities for testing that differently typed pointers to overlapping bytes of a buffer, formed + * with bufferView or bufferArrayView, behave correctly. i.e. that implementations do not apply + * type-based alias analysis. Used by the bufferView, bufferArrayView and pointer aliasing tests. + */ + +import { assert } from '../../../../../common/util/util.js'; +import { GPUTest } from '../../../../gpu_test.js'; +import { float32ToFloat16Bits } from '../../../../util/conversion.js'; + +/** @returns the bit pattern of the f32 value 'f' as an i32. */ +export function f32Bits(f: number): number { + return new Int32Array(new Float32Array([f]).buffer)[0]; +} + +/** The scalar types that may be used by the lanes of a mixed type. */ +export type MixedScalar = 'i32' | 'u32' | 'f32' | 'f16'; + +/** @returns the size in bytes of the scalar type 's'. */ +function mixedScalarSize(s: MixedScalar): number { + return s === 'f16' ? 2 : 4; +} + +/** A scalar lane of a mixed type. */ +interface MixedLane { + /** The scalar type of the lane. */ + elem: MixedScalar; + /** The byte offset of the lane from the start of the value. */ + offset: number; + /** The WGSL accessor suffix for the lane. e.g. '', '[1]', '.a' or '.v[1]' */ + access: string; +} + +/** + * A scalar, vector or structure type used by the mixed type aliasing tests. + * A value of the type is treated as a flat list of scalar lanes. Structures may contain padding, + * which is not part of any lane. + */ +export interface MixedType { + /** The WGSL type name. */ + name: string; + /** The size in bytes. This is also the array element stride. */ + size: number; + /** The alignment in bytes. */ + align: number; + /** The scalar lanes, in declaration order. */ + lanes: MixedLane[]; + /** Present if the type is a structure. */ + struct?: { + /** The structure members, with their byte offsets. */ + members: { name: string; type: MixedType; offset: number }[]; + }; +} + +function scalar(elem: MixedScalar): MixedType { + const size = mixedScalarSize(elem); + return { name: elem, size, align: size, lanes: [{ elem, offset: 0, access: '' }] }; +} + +function vec(n: 2 | 4, elem: MixedScalar): MixedType { + const size = n * mixedScalarSize(elem); + const lanes = [...Array(n).keys()].map(k => ({ + elem, + offset: k * mixedScalarSize(elem), + access: `[${k}]`, + })); + return { name: `vec${n}<${elem}>`, size, align: size, lanes }; +} + +/** @returns a structure type, laid out with the WGSL layout rules. */ +function structType(name: string, members: [string, MixedType][]): MixedType { + const roundUp = (n: number, k: number) => Math.ceil(n / k) * k; + let end = 0; + let align = 1; + const laidOut = members.map(([memberName, type]) => { + const offset = roundUp(end, type.align); + end = offset + type.size; + align = Math.max(align, type.align); + return { name: memberName, type, offset }; + }); + const lanes = laidOut.flatMap(m => + m.type.lanes.map(l => ({ + elem: l.elem, + offset: m.offset + l.offset, + access: `.${m.name}${l.access}`, + })) + ); + return { name, size: roundUp(end, align), align, lanes, struct: { members: laidOut } }; +} + +const kI32 = scalar('i32'); +const kU32 = scalar('u32'); +const kF32 = scalar('f32'); +const kF16 = scalar('f16'); +const kVec2I = vec(2, 'i32'); +const kVec4I = vec(4, 'i32'); +const kVec4U = vec(4, 'u32'); +const kVec2F = vec(2, 'f32'); +const kVec4F = vec(4, 'f32'); +const kVec2H = vec(2, 'f16'); +const kVec4H = vec(4, 'f16'); +// Structures without padding. +const kStructI = structType('SI', [ + ['a', kI32], + ['b', kI32], + ['c', kI32], + ['d', kI32], +]); +const kStructF = structType('SF', [ + ['a', kF32], + ['b', kF32], + ['c', kF32], + ['d', kF32], +]); +const kStructV = structType('SV', [ + ['v', kVec2I], + ['s', kI32], + ['t', kI32], +]); +// Structures with padding. +/** 4 bytes of padding between 'a' and 'b'. */ +const kPaddedI = structType('PI', [ + ['a', kI32], + ['b', kVec2I], +]); +/** 4 bytes of padding between 'a' and 'b'. */ +const kPaddedF = structType('PF', [ + ['a', kF32], + ['b', kVec2F], +]); +/** 4 bytes of trailing padding. */ +const kPaddedTrailingF = structType('PT', [ + ['v', kVec2F], + ['s', kF32], +]); +/** 2 bytes of padding between 'a' and 'b', within the first 4-byte word. */ +const kPaddedH = structType('PH', [ + ['a', kF16], + ['b', kF32], +]); + +/** @returns true if the type 't' has any f16 lanes. */ +export function mixedTypeUsesF16(t: MixedType): boolean { + return t.lanes.some(l => l.elem === 'f16'); +} + +/** @returns the WGSL name of the type 't'. */ +export function mixedTypeName(t: MixedType): string { + return t.name; +} + +/** @returns a WGSL constructor expression for a value of type 't', given an expression per lane. */ +function mixedTypeCtorFromLanes(t: MixedType, lanes: string[]): string { + if (t.struct) { + let next = 0; + const members = t.struct.members.map(m => { + const n = m.type.lanes.length; + const expr = mixedTypeCtorFromLanes(m.type, lanes.slice(next, next + n)); + next += n; + return expr; + }); + return `${t.name}(${members.join(', ')})`; + } + return t.lanes.length === 1 ? lanes[0] : `${t.name}(${lanes.join(', ')})`; +} + +/** + * @returns a WGSL constructor expression for a value of type 't', where lane 'k' has the value + * 'base + k', and 'base' is a WGSL scalar expression. + */ +function mixedTypeCtor(t: MixedType, base: string): string { + const lanes = t.lanes.map(({ elem }, k) => + k === 0 ? `${elem}(${base})` : `${elem}(${base}) + ${elem}(${k})` + ); + return mixedTypeCtorFromLanes(t, lanes); +} + +/** @returns the WGSL declarations of the structures used by the types 'types', including nested. */ +function mixedTypeStructDecls(types: MixedType[]): string { + const decls = new Map(); + const visit = (t: MixedType) => { + if (!t.struct || decls.has(t.name)) { + return; + } + t.struct.members.forEach(m => visit(m.type)); + const members = t.struct.members.map(m => `${m.name} : ${m.type.name}`).join(', '); + decls.set(t.name, `struct ${t.name} { ${members} }`); + }; + types.forEach(visit); + return [...decls.values()].join('\n'); +} + +/** + * @returns WGSL declarations required by the mixed type ops, where 'int' is the type of the loaded + * view, and 'other' is the type of the stored view. This declares the structures used by the types + * and the functions: + * 'mixed_sub(a, b)' - lane-wise 'a - b' + * 'mixed_xor(a, b)' - lane-wise 'a ^ b' + * If either type uses f16, the shader must also 'enable f16;'. + */ +export function mixedTypeDecls(int: MixedType, other: MixedType): string { + const ty = mixedTypeName(int); + const laneWise = (name: string, op: string) => { + if (!int.struct) { + return `fn ${name}(a : ${ty}, b : ${ty}) -> ${ty} { return a ${op} b; }`; + } + const lanes = int.lanes.map(({ access: l }) => ` r${l} = a${l} ${op} b${l};`); + return `fn ${name}(a : ${ty}, b : ${ty}) -> ${ty} {\n var r : ${ty};\n${lanes.join( + '\n' + )}\n return r;\n}`; + }; + return `${mixedTypeStructDecls([int, other])} +${laneWise('mixed_sub', '-')} +${laneWise('mixed_xor', '^')} +`; +} + +/** + * @returns WGSL statements that write each lane of the value 'value' of type 't', as an i32 bit + * pattern, to consecutive elements of the i32 array 'array'. All lanes of 't' must be 32-bit. + */ +export function mixedTypeWriteLanes(t: MixedType, value: string, array: string): string { + return t.lanes + .map(({ access }, k) => `${array}[${k}] = bitcast(${value}${access});`) + .join('\n '); +} + +/** + * The pairs of view types tested. + * 'int' is a type with only i32 or u32 lanes, and at most 4 lanes. It is the only view that is + * loaded from, which keeps the expected results exact: float values are never loaded, so denormal + * flushing and NaN canonicalization cannot affect the results. + * 'other' is a type that is only stored to. + */ +export const kMixedTypePairs: Record = { + // scalar - scalar + i32_f32: { int: kI32, other: kF32 }, + u32_f32: { int: kU32, other: kF32 }, + i32_i32: { int: kI32, other: kI32 }, + // vector - scalar + vec2i_f32: { int: kVec2I, other: kF32 }, + vec4i_f32: { int: kVec4I, other: kF32 }, + vec4u_f32: { int: kVec4U, other: kF32 }, + vec4u_u32: { int: kVec4U, other: kU32 }, + // scalar - vector + i32_vec2f: { int: kI32, other: kVec2F }, + i32_vec4f: { int: kI32, other: kVec4F }, + i32_vec4u: { int: kI32, other: kVec4U }, + i32_vec4i: { int: kI32, other: kVec4I }, + // vector - vector + vec2i_vec2f: { int: kVec2I, other: kVec2F }, + vec4i_vec4f: { int: kVec4I, other: kVec4F }, + vec4u_vec4f: { int: kVec4U, other: kVec4F }, + vec4i_vec2f: { int: kVec4I, other: kVec2F }, + vec2i_vec4f: { int: kVec2I, other: kVec4F }, + vec4i_vec4u: { int: kVec4I, other: kVec4U }, + vec4i_vec4i: { int: kVec4I, other: kVec4I }, + // struct - scalar + structi_f32: { int: kStructI, other: kF32 }, + structv_f32: { int: kStructV, other: kF32 }, + i32_structf: { int: kI32, other: kStructF }, + // struct - vector + structi_vec4f: { int: kStructI, other: kVec4F }, + structi_vec2f: { int: kStructI, other: kVec2F }, + vec4i_structf: { int: kVec4I, other: kStructF }, + // struct - struct + structi_structf: { int: kStructI, other: kStructF }, + structv_structf: { int: kStructV, other: kStructF }, + // f16. The f16 stores only write part of a 4-byte word. + i32_f16: { int: kI32, other: kF16 }, + u32_f16: { int: kU32, other: kF16 }, + i32_vec2h: { int: kI32, other: kVec2H }, + i32_vec4h: { int: kI32, other: kVec4H }, + vec4i_f16: { int: kVec4I, other: kF16 }, + vec2i_vec4h: { int: kVec2I, other: kVec4H }, + vec4i_vec4h: { int: kVec4I, other: kVec4H }, + // padded struct. Stores of padded structures must not write the padding. + paddedi_f32: { int: kPaddedI, other: kF32 }, + paddedi_vec4f: { int: kPaddedI, other: kVec4F }, + i32_paddedf: { int: kI32, other: kPaddedF }, + vec4i_paddedf: { int: kVec4I, other: kPaddedF }, + i32_paddedtf: { int: kI32, other: kPaddedTrailingF }, + vec4i_paddedtf: { int: kVec4I, other: kPaddedTrailingF }, + paddedi_paddedf: { int: kPaddedI, other: kPaddedF }, + paddedi_structf: { int: kPaddedI, other: kStructF }, + structi_paddedf: { int: kStructI, other: kPaddedF }, + // padded struct with f16 + i32_paddedh: { int: kI32, other: kPaddedH }, + vec4i_paddedh: { int: kVec4I, other: kPaddedH }, + paddedi_paddedh: { int: kPaddedI, other: kPaddedH }, +}; + +/** + * How the two views overlap. Views always start on a 4-byte boundary. + * 'same_start' - Both views start at the same byte. + * 'offset' - The views start at different bytes, and a lane of one view overlaps a lane of + * the other. Not all pairs of types can do this while respecting alignment. + * 'padding' - The views overlap, but only lanes of one view overlap padding of the other. + * Only possible for structures with padding. + */ +export const kMixedTypeOverlaps = ['same_start', 'offset', 'padding'] as const; +export type MixedTypeOverlap = (typeof kMixedTypeOverlaps)[number]; + +/** + * @returns the word offsets of the 'int' and 'other' views for the given overlap, or undefined if + * the types cannot be positioned to satisfy the overlap. Word offsets are always at least 4. + */ +export function mixedTypeWords( + int: MixedType, + other: MixedType, + overlap: MixedTypeOverlap +): { intWord: number; otherWord: number } | undefined { + const lanesOverlap = (a: number, b: number) => + int.lanes.some(li => + other.lanes.some(lo => { + const i0 = a * 4 + li.offset; + const o0 = b * 4 + lo.offset; + return i0 < o0 + mixedScalarSize(lo.elem) && o0 < i0 + mixedScalarSize(li.elem); + }) + ); + const viewsOverlap = (a: number, b: number) => + a * 4 < b * 4 + other.size && b * 4 < a * 4 + int.size; + for (let a = 4; a < 12; a++) { + if ((a * 4) % int.align !== 0) continue; + for (let b = 4; b < 12; b++) { + if ((b * 4) % other.align !== 0) continue; + let ok = false; + switch (overlap) { + case 'same_start': + ok = a === b; + break; + case 'offset': + ok = a !== b && lanesOverlap(a, b); + break; + case 'padding': + ok = viewsOverlap(a, b) && !lanesOverlap(a, b); + break; + } + if (ok) { + return { intWord: a, otherWord: b }; + } + } + } + return undefined; +} + +/** A simulated memory of 4-byte words, used to calculate the expected results of mixed type tests. */ +export interface MixedTypeMemory { + ld(ptr: number): number; + st(ptr: number, value: number): void; +} + +/** The runtime arguments passed to a mixed type op. */ +export interface MixedTypeArgs { + /** A value to write through the 'int' view. Also used as a second value for the 'other' view. */ + va: number; + /** A value to write through the 'other' view. */ + vf: number; + /** A loop count. */ + n: number; +} + +/** + * Stores a value of type 't' where lane 'k' is 'base + k', at word 'ptr'. + * Only the bytes of the lanes are written. Padding is left untouched. + */ +function simStore(m: MixedTypeMemory, t: MixedType, ptr: number, base: number) { + t.lanes.forEach(({ elem, offset }, k) => { + const v = base + k; + const word = ptr + Math.floor(offset / 4); + switch (elem) { + case 'i32': + case 'u32': + m.st(word, v | 0); + break; + case 'f32': + m.st(word, f32Bits(v)); + break; + case 'f16': { + const shift = (offset % 4) * 8; + const mask = 0xffff << shift; + m.st(word, (m.ld(word) & ~mask) | ((float32ToFloat16Bits(v) << shift) & mask)); + break; + } + } + }); +} + +/** @returns the lanes of a value of type 't' at word 'ptr', as i32 bit patterns. */ +function simLoad(m: MixedTypeMemory, t: MixedType, ptr: number): number[] { + return t.lanes.map(({ elem, offset }) => { + assert(elem === 'i32' || elem === 'u32', 'only integer views are loaded'); + return m.ld(ptr + offset / 4); + }); +} + +/** + * An operation on a pointer 'pi' to an integer type, and a pointer 'pf' to another type, that may + * refer to overlapping memory. Shaders using these ops must include mixedTypeDecls(int, other). + */ +export interface MixedTypeOp { + /** + * @returns a WGSL function body that returns a value of type 'int', and may use: + * 'pi' - a pointer to 'int' + * 'pf' - a pointer to 'other' + * 'va' - an i32 value + * 'vf' - an f32 value, whose value is exactly representable as an f16 + * 'n' - an i32 loop count + */ + wgsl: (int: MixedType, other: MixedType) => string; + /** + * Simulates 'wgsl', where 'pi' and 'pf' are word offsets into 'm'. + * @returns the lanes of the returned value, as i32 bit patterns. + */ + sim: ( + m: MixedTypeMemory, + int: MixedType, + other: MixedType, + pi: number, + pf: number, + args: MixedTypeArgs + ) => number[]; +} + +/** + * Operations that produce different results if an implementation assumes that differently typed + * pointers do not alias (i.e. applies C/C++ style type-based alias analysis). + */ +export const kMixedTypeOps: Record = { + int_store_other_store_int_load: { + wgsl: (int, other) => ` + *pi = ${mixedTypeCtor(int, 'va')}; + *pf = ${mixedTypeCtor(other, 'vf')}; + return *pi;`, + sim: (m, int, other, pi, pf, { va, vf }) => { + simStore(m, int, pi, va); + simStore(m, other, pf, vf); + return simLoad(m, int, pi); + }, + }, + other_store_int_load: { + wgsl: (int, other) => ` + *pf = ${mixedTypeCtor(other, 'vf')}; + return *pi;`, + sim: (m, int, other, pi, pf, { vf }) => { + simStore(m, other, pf, vf); + return simLoad(m, int, pi); + }, + }, + int_load_other_store_int_load: { + wgsl: (int, other) => ` + let before = *pi; + *pf = ${mixedTypeCtor(other, 'vf')}; + return mixed_sub(*pi, before);`, + sim: (m, int, other, pi, pf, { vf }) => { + const before = simLoad(m, int, pi); + simStore(m, other, pf, vf); + return simLoad(m, int, pi).map((v, k) => (v - before[k]) | 0); + }, + }, + dead_other_store: { + // The first store to '*pf' is only dead if '*pi' does not alias '*pf'. + wgsl: (int, other) => ` + *pf = ${mixedTypeCtor(other, 'vf')}; + let tmp = *pi; + *pf = ${mixedTypeCtor(other, 'va')}; + return tmp;`, + sim: (m, int, other, pi, pf, { va, vf }) => { + simStore(m, other, pf, vf); + const tmp = simLoad(m, int, pi); + simStore(m, other, pf, va); + return tmp; + }, + }, + loop_other_store_int_load: { + // The load of '*pi' must not be hoisted out of the loop. + wgsl: (int, other) => ` + var acc = ${mixedTypeName(int)}(); + for (var i = 0; i < n; i++) { + *pf = ${mixedTypeCtor(other, 'i')}; + acc = mixed_xor(acc, *pi); + } + return acc;`, + sim: (m, int, other, pi, pf, { n }) => { + const acc = new Array(int.lanes.length).fill(0); + for (let i = 0; i < n; i++) { + simStore(m, other, pf, i); + simLoad(m, int, pi).forEach((v, k) => { + acc[k] ^= v; + }); + } + return acc; + }, + }, +}; + +/** The number of 4-byte words in each of the buffers 'A' and 'B' used by runMixedTypeAliasingTest. */ +export const kMixedTypeBufferWords = 32; + +/** + * The runtime parameters passed to the shader of runMixedTypeAliasingTest in 'input.p'. + * p[0]: 'va' - an i32 value to write. + * p[1]: 'vf' - an f32 value to write, passed as an i32 and converted to f32. + * p[2]: 'n' - a loop count. + * p[3]: 'idx' - a runtime value that views may use to form offsets and indices. + */ +export const kMixedTypeParams = [1000, 2000, 4, 1] as const; + +/** The value of 'input.p[3]', which may be used by views to form offsets and indices. */ +export const kMixedTypeIdx = kMixedTypeParams[3]; + +/** + * The kinds of buffer that the views are formed on: + * 'workgroup' - var of type buffer + * 'storage' - var of type buffer + * 'storage_unsized' - var of type buffer + */ +export const kMixedTypeBuffers = ['workgroup', 'storage', 'storage_unsized'] as const; +export type MixedTypeBuffer = (typeof kMixedTypeBuffers)[number]; + +interface MixedTypeAliasingParams { + /** The kind of buffer the views are formed on. */ + buffer: MixedTypeBuffer; + /** The type of the view that is loaded from. */ + int: MixedType; + /** The type of the view that is only stored to. */ + other: MixedType; + /** The word offset of the 'int' view. */ + intWord: number; + /** The word offset of the 'other' view. */ + otherWord: number; + /** + * @returns a WGSL expression for a pointer to a value of type 'type' at word offset 'word' of the + * buffer variable named 'buffer'. May use 'input.p[3]' as a runtime value. + */ + view: (type: MixedType, buffer: string, word: number) => string; + /** + * If true, the 'other' view is formed in the same buffer as the 'int' view. + * Otherwise it is formed in a different buffer. + */ + aliased: boolean; + /** The operation to perform. */ + op: MixedTypeOp; + /** + * If true, the views are formed in the entry point and passed to 'f' as pointer parameters. + * Otherwise the views are formed within 'f'. + */ + pointerParams?: boolean; + /** WGSL directives placed at the start of the shader. */ + directives?: string; +} + +/** + * Runs a test where two differently typed pointers are formed from buffer views and used within a + * single function, or passed to a function as pointer parameters if 'pointerParams' is true. + * If 'aliased' is true, then both views refer to the same buffer, otherwise the 'other' view refers + * to a different buffer. + * + * The shader initializes buffers 'A' and 'B' from 'input', calls the op, and writes the lanes of + * the result of the op, followed by the contents of 'A' and 'B', to 'output'. + */ +export function runMixedTypeAliasingTest(t: GPUTest, params: MixedTypeAliasingParams) { + t.skipIfLanguageFeatureNotSupported('buffer_view'); + const usesF16 = mixedTypeUsesF16(params.int) || mixedTypeUsesF16(params.other); + if (usesF16) { + t.skipIfDeviceDoesNotHaveFeature('shader-f16'); + } + + const N = kMixedTypeBufferWords; + let decls = ''; + switch (params.buffer) { + case 'workgroup': + decls = `var A : buffer<${N * 4}>;\nvar B : buffer<${N * 4}>;`; + break; + case 'storage': + decls = `@group(0) @binding(2) var A : buffer<${N * 4}>; +@group(0) @binding(3) var B : buffer<${N * 4}>;`; + break; + case 'storage_unsized': + decls = `@group(0) @binding(2) var A : buffer; +@group(0) @binding(3) var B : buffer;`; + break; + } + + const intTy = mixedTypeName(params.int); + const otherTy = mixedTypeName(params.other); + const writeResult = mixedTypeWriteLanes(params.int, 'r', 'output.r'); + const viewI = params.view(params.int, 'A', params.intWord); + const viewF = params.view(params.other, params.aliased ? 'A' : 'B', params.otherWord); + // Only storage pointers may specify an access mode. + const ptr = (ty: string) => + params.buffer === 'workgroup' ? `ptr` : `ptr`; + let fn = ''; + let call = ''; + if (params.pointerParams) { + fn = `fn f(pi : ${ptr(intTy)}, pf : ${ptr(otherTy)}, + va : i32, vf : f32, n : i32) -> ${intTy} { + ${params.op.wgsl(params.int, params.other)} +}`; + call = `f(${viewI}, ${viewF}, input.p[0], f32(input.p[1]), input.p[2])`; + } else { + fn = `fn f(va : i32, vf : f32, n : i32) -> ${intTy} { + let pi = ${viewI}; + let pf = ${viewF}; + ${params.op.wgsl(params.int, params.other)} +}`; + call = `f(input.p[0], f32(input.p[1]), input.p[2])`; + } + + const wgsl = `${params.directives ?? ''} +${usesF16 ? 'enable f16;' : ''} +struct In { + a : array, + b : array, + p : array, +} + +struct Out { + r : array, + a : array, + b : array, +} + +@group(0) @binding(0) var input : In; +@group(0) @binding(1) var output : Out; + +${decls} + +${mixedTypeDecls(params.int, params.other)} + +${fn} + +@compute @workgroup_size(1) +fn main() { + *bufferView>(&A, 0) = input.a; + *bufferView>(&B, 0) = input.b; + let r = ${call}; + ${writeResult} + output.a = *bufferView>(&A, 0); + output.b = *bufferView>(&B, 0); +} +`; + + // Simulate the op. 'A' is at word 0, and 'B' is at word N. + const initA = Array.from({ length: N }, (_, i) => 10 + i); + const initB = Array.from({ length: N }, (_, i) => 100 + i); + const mem = [...initA, ...initB]; + const memory: MixedTypeMemory = { + ld: ptr => mem[ptr], + st: (ptr, value) => { + mem[ptr] = value | 0; + }, + }; + const [va, vf, n] = kMixedTypeParams; + const pi = params.intWord; + const pf = (params.aliased ? 0 : N) + params.otherWord; + const r = params.op.sim(memory, params.int, params.other, pi, pf, { va, vf, n }); + const rOut = [0, 0, 0, 0]; + r.forEach((v, k) => { + rOut[k] = v; + }); + const expected = new Int32Array([...rOut, ...mem]); + + const pipeline = t.device.createComputePipeline({ + layout: 'auto', + compute: { module: t.device.createShaderModule({ code: wgsl }) }, + }); + + const inputBuffer = t.makeBufferWithContents( + new Int32Array([...initA, ...initB, ...kMixedTypeParams]), + GPUBufferUsage.STORAGE + ); + const outputBuffer = t.createBufferTracked({ + size: expected.byteLength, + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, + }); + const entries: GPUBindGroupEntry[] = [ + { binding: 0, resource: { buffer: inputBuffer } }, + { binding: 1, resource: { buffer: outputBuffer } }, + ]; + if (params.buffer !== 'workgroup') { + for (const binding of [2, 3]) { + const buffer = t.createBufferTracked({ size: N * 4, usage: GPUBufferUsage.STORAGE }); + entries.push({ binding, resource: { buffer } }); + } + } + const bg = t.device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries }); + + const encoder = t.device.createCommandEncoder(); + const pass = encoder.beginComputePass(); + pass.setPipeline(pipeline); + pass.setBindGroup(0, bg); + pass.dispatchWorkgroups(1); + pass.end(); + t.queue.submit([encoder.finish()]); + + t.expectGPUBufferValuesEqual(outputBuffer, expected); +} diff --git a/src/webgpu/shader/execution/expression/call/user/ptr_aliasing.spec.ts b/src/webgpu/shader/execution/expression/call/user/ptr_aliasing.spec.ts new file mode 100644 index 000000000000..6f949460f483 --- /dev/null +++ b/src/webgpu/shader/execution/expression/call/user/ptr_aliasing.spec.ts @@ -0,0 +1,1178 @@ +export const description = ` +Execution tests for aliased pointer parameters. + +With the 'unrestricted_aliasing' language feature, pointer arguments passed to a user-declared +function may alias each other (or a module-scope variable accessed by the callee), even when one of +the accesses is a write. These tests check that implementations honour the WGSL memory semantics in +that case, i.e. that no code generation or downstream compiler step assumes that pointer parameters +are non-aliasing. + +Every test is run with 'aliased' set to true and false: + * aliased=false: The pointers refer to distinct memory locations with distinct root identifiers. + These cases are valid without the 'unrestricted_aliasing' language feature and act as a control + for the expectations. + * aliased=true: The pointers refer to overlapping memory locations. These cases require the + 'unrestricted_aliasing' language feature. + +Cases that form pointers using runtime indices into the same root variable always require the +'unrestricted_aliasing' language feature, as they cannot be statically proven not to alias. +`; + +import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; +import { keysOf } from '../../../../../../common/util/data_tables.js'; +import { AllFeaturesMaxLimitsGPUTest, GPUTest } from '../../../../../gpu_test.js'; +import { + kMixedTypeBuffers, + kMixedTypeIdx, + kMixedTypeOps, + kMixedTypeOverlaps, + kMixedTypePairs, + MixedType, + mixedTypeName, + mixedTypeWords, + runMixedTypeAliasingTest, +} from '../mixed_type_aliasing_utils.js'; + +export const g = makeTestGroup(AllFeaturesMaxLimitsGPUTest); + +type AddressSpace = 'function' | 'private' | 'workgroup' | 'storage'; + +const kAddressSpaces = ['function', 'private', 'workgroup', 'storage'] as const; +const kModuleScopeAddressSpaces = ['private', 'workgroup', 'storage'] as const; +const kAtomicAddressSpaces = ['workgroup', 'storage'] as const; + +/** + * WGSL type declarations shared by all shaders. + * + * Data is made entirely of i32 so it has no padding, and its memory layout in i32 units is: + * x: 0, y: 1, arr: 2..5, s: 6..9, big: 10..41 + */ +const kTypeDecls = ` +struct S { + a : i32, + b : i32, + c : i32, + d : i32, +} + +struct Data { + x : i32, + y : i32, + arr : array, + s : S, + big : array, +} + +struct AtomicData { + x : atomic, + y : atomic, + arr : array, 4>, +} + +struct In { + a : Data, + b : Data, + p : array, +} + +struct Out { + r : array, + a : Data, + b : Data, +} +`; + +/** The number of i32s in 'Data'. */ +const kDataSize = 42; +/** The offset of 'Data.big' in i32 units. */ +const kBig = 10; +/** The number of elements in 'Data.big'. */ +const kBigSize = 32; +/** The offset of variable 'B' in the simulated memory. Variable 'A' is at offset 0. */ +const kB = kDataSize; +/** The initial contents of variable 'A'. */ +const kInitA = Array.from({ length: kDataSize }, (_, i) => 10 + i); +/** The initial contents of variable 'B'. */ +const kInitB = Array.from({ length: kDataSize }, (_, i) => 100 + i); +/** + * The runtime parameters passed to the shader in 'input.p'. + * Shaders use these to prevent the compiler from constant folding the tests away. + * p[0]: 'va' - a value to write. + * p[1]: 'vb' - another value to write. + * p[2]: 'n' - a loop count. + * p[3]: 'idx' - a runtime index. + * p[4]: 'm' - a long loop count, used to iterate over 'Data.big'. Must be less than kBigSize. + * p[5..7]: unused. + */ +const kParams = [1000, 2000, 4, 1, kBigSize - 1, 0, 0, 0]; +const [kVA, kVB, kN, kIdx, kM] = kParams; + +/** A simulation of the memory of variables 'A' and 'B', and the output results. */ +class Memory { + readonly mem: number[] = [...kInitA, ...kInitB]; + readonly r: number[] = [0, 0, 0, 0]; + ld(ptr: number): number { + return this.mem[ptr]; + } + st(ptr: number, value: number) { + this.mem[ptr] = value | 0; + } + /** @returns the expected contents of the 'Out' buffer. */ + expected(): Int32Array { + return new Int32Array([...this.r, ...this.mem]); + } +} + +/** @returns the WGSL pointer type for a pointer to 'type' in the given address space. */ +function ptr(space: AddressSpace, type: string) { + switch (space) { + case 'storage': + return `ptr`; + default: + return `ptr<${space}, ${type}>`; + } +} + +interface ShaderParams { + /** The address space of variables 'A' and 'B'. */ + space: AddressSpace; + /** Whether the shader requires the 'unrestricted_aliasing' language feature. */ + requiresAliasing: boolean; + /** Module-scope helper functions. */ + helpers: string; + /** The body of the entry point. Runs after 'A' and 'B' are initialized. */ + body: string; + /** The expected contents of the output buffer. */ + expected: Int32Array; + /** + * If true, then 'A' and 'B' are of type 'AtomicData', and 'body' is responsible for initializing + * them and for writing them to 'output'. + */ + atomic?: boolean; + /** + * If true, then 'A' and 'B' are untyped 'buffer' variables with the same size as 'Data'. + * They are initialized from, and written to, 'output' through 'bufferView'. + * Requires the 'buffer_view' language feature, and 'space' must be 'workgroup' or 'storage'. + */ + bufferView?: boolean; +} + +/** + * Builds and runs a compute shader that declares two variables 'A' and 'B' in the given address + * space, initializes them from 'input', runs 'body', then writes 'A' and 'B' to 'output'. + */ +function run(t: GPUTest, params: ShaderParams) { + t.skipIfLanguageFeatureNotSupported('unrestricted_pointer_parameters'); + //if (params.requiresAliasing) { + // t.skipIfLanguageFeatureNotSupported('unrestricted_aliasing'); + //} + if (params.bufferView) { + t.skipIfLanguageFeatureNotSupported('buffer_view'); + } + + const varType = params.atomic + ? 'AtomicData' + : params.bufferView + ? `buffer<${kDataSize * 4}>` + : 'Data'; + let moduleVars = ''; + let functionVars = ''; + switch (params.space) { + case 'function': + functionVars = `var A : ${varType};\n var B : ${varType};`; + break; + case 'private': + case 'workgroup': + moduleVars = `var<${params.space}> A : ${varType};\nvar<${params.space}> B : ${varType};`; + break; + case 'storage': + moduleVars = ` +@group(0) @binding(2) var A : ${varType}; +@group(0) @binding(3) var B : ${varType};`; + break; + } + + let init = `A = input.a;\n B = input.b;`; + let dump = `output.a = A;\n output.b = B;`; + if (params.atomic) { + init = ''; + dump = ''; + } else if (params.bufferView) { + init = `*bufferView(&A, 0u) = input.a;\n *bufferView(&B, 0u) = input.b;`; + dump = `output.a = *bufferView(&A, 0u);\n output.b = *bufferView(&B, 0u);`; + } + + const code = ` +${params.requiresAliasing ? '// requires unrestricted_aliasing;' : ''} + +${kTypeDecls} + +@group(0) @binding(0) var input : In; +@group(0) @binding(1) var output : Out; + +${moduleVars} + +${params.helpers} + +@compute @workgroup_size(1) +fn main() { + ${functionVars} + ${init} + ${params.body} + ${dump} +} +`; + + const pipeline = t.device.createComputePipeline({ + layout: 'auto', + compute: { + module: t.device.createShaderModule({ code }), + entryPoint: 'main', + }, + }); + + const inputBuffer = t.makeBufferWithContents( + new Int32Array([...kInitA, ...kInitB, ...kParams]), + GPUBufferUsage.STORAGE + ); + const outputBuffer = t.createBufferTracked({ + size: params.expected.byteLength, + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, + }); + + const entries: GPUBindGroupEntry[] = [ + { binding: 0, resource: { buffer: inputBuffer } }, + { binding: 1, resource: { buffer: outputBuffer } }, + ]; + if (params.space === 'storage') { + for (const binding of [2, 3]) { + const buffer = t.createBufferTracked({ + size: kDataSize * 4, + usage: GPUBufferUsage.STORAGE, + }); + entries.push({ binding, resource: { buffer } }); + } + } + + const bindGroup = t.device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries, + }); + + const encoder = t.device.createCommandEncoder(); + const pass = encoder.beginComputePass(); + pass.setPipeline(pipeline); + pass.setBindGroup(0, bindGroup); + pass.dispatchWorkgroups(1); + pass.end(); + t.queue.submit([encoder.finish()]); + + t.expectGPUBufferValuesEqual(outputBuffer, params.expected); +} + +/** Scalar i32 memory locations within 'Data' that pointers can be formed to. */ +const kScalarTargets = { + scalar: { wgsl: 'x', offset: 0 }, + array_element: { wgsl: 'arr[1]', offset: 3 }, + struct_member: { wgsl: 's.b', offset: 7 }, +}; + +interface TwoPointerOp { + /** The body of 'fn f(pa : ptr, pb : ptr, va : i32, vb : i32, n : i32) -> i32'. */ + wgsl: string; + /** Simulates 'f'. */ + sim: (m: Memory, pa: number, pb: number) => number; +} + +const kTwoPointerOps: Record = { + store_store_load: { + wgsl: ` + *pa = va; + *pb = vb; + return *pa;`, + sim: (m, pa, pb) => { + m.st(pa, kVA); + m.st(pb, kVB); + return m.ld(pa); + }, + }, + store_load: { + wgsl: ` + *pa = va; + return *pb;`, + sim: (m, pa, pb) => { + m.st(pa, kVA); + return m.ld(pb); + }, + }, + load_store_load: { + wgsl: ` + let before = *pa; + *pb = vb; + return *pa - before;`, + sim: (m, pa, pb) => { + const before = m.ld(pa); + m.st(pb, kVB); + return m.ld(pa) - before; + }, + }, + compound_assign: { + wgsl: ` + *pa += *pb; + *pa += *pb; + return *pa;`, + sim: (m, pa, pb) => { + m.st(pa, m.ld(pa) + m.ld(pb)); + m.st(pa, m.ld(pa) + m.ld(pb)); + return m.ld(pa); + }, + }, + increment: { + wgsl: ` + (*pa)++; + (*pb)++; + return *pa;`, + sim: (m, pa, pb) => { + m.st(pa, m.ld(pa) + 1); + m.st(pb, m.ld(pb) + 1); + return m.ld(pa); + }, + }, + swap: { + wgsl: ` + let tmp = *pa; + *pa = *pb; + *pb = tmp; + return *pa;`, + sim: (m, pa, pb) => { + const tmp = m.ld(pa); + m.st(pa, m.ld(pb)); + m.st(pb, tmp); + return m.ld(pa); + }, + }, + loop_accumulate: { + // The load of '*pb' must not be hoisted out of the loop. + wgsl: ` + for (var i = 0; i < n; i++) { + *pa += *pb; + } + return *pa;`, + sim: (m, pa, pb) => { + for (let i = 0; i < kN; i++) { + m.st(pa, m.ld(pa) + m.ld(pb)); + } + return m.ld(pa); + }, + }, + loop_store: { + // The load of '*pb' must not be hoisted out of the loop. + wgsl: ` + var sum = 0; + for (var i = 0; i < n; i++) { + *pa = i; + sum += *pb; + } + return sum;`, + sim: (m, pa, pb) => { + let sum = 0; + for (let i = 0; i < kN; i++) { + m.st(pa, i); + sum += m.ld(pb); + } + return sum; + }, + }, + nested_calls: { + wgsl: ` + *pa = va; + store_i32(pb, vb); + return load_i32(pa);`, + sim: (m, pa, pb) => { + m.st(pa, kVA); + m.st(pb, kVB); + return m.ld(pa); + }, + }, + nested_calls_swapped: { + // Pass the pointers through to another function in the opposite order. + wgsl: ` + return store_store_load(pb, pa, vb, va);`, + sim: (m, pa, pb) => { + m.st(pb, kVB); + m.st(pa, kVA); + return m.ld(pb); + }, + }, + dead_store: { + // The first store to '*pa' is only dead if '*pb' does not alias '*pa'. + wgsl: ` + *pa = va; + let tmp = *pb; + *pa = vb; + return tmp;`, + sim: (m, pa, pb) => { + m.st(pa, kVA); + const tmp = m.ld(pb); + m.st(pa, kVB); + return tmp; + }, + }, + dead_store_nested_call: { + // As 'dead_store', but the intervening read happens in another function. + wgsl: ` + *pa = va; + let tmp = load_i32(pb); + *pa = vb; + return tmp;`, + sim: (m, pa, pb) => { + m.st(pa, kVA); + const tmp = m.ld(pb); + m.st(pa, kVB); + return tmp; + }, + }, + dead_store_loop: { + // Each store to '*pa' is overwritten by the next iteration, so only the last is live if '*pb' + // does not alias '*pa'. + wgsl: ` + var sum = 0; + for (var i = 0; i < n; i++) { + *pa = va + i; + sum += *pb; + *pa = vb; + } + return sum;`, + sim: (m, pa, pb) => { + let sum = 0; + for (let i = 0; i < kN; i++) { + m.st(pa, kVA + i); + sum += m.ld(pb); + m.st(pa, kVB); + } + return sum; + }, + }, +}; + +g.test('two_pointers') + .desc( + `Test that a function with two i32 pointer parameters behaves correctly when both pointers refer +to the same memory location.` + ) + .params(u => + u + .combine('address_space', kAddressSpaces) + .combine('target', keysOf(kScalarTargets)) + .combine('aliased', [true, false]) + .beginSubcases() + .combine('op', keysOf(kTwoPointerOps)) + ) + .fn(t => { + const space = t.params.address_space; + const target = kScalarTargets[t.params.target]; + const op = kTwoPointerOps[t.params.op]; + const pi32 = ptr(space, 'i32'); + + const pa = 0 + target.offset; + const pb = (t.params.aliased ? 0 : kB) + target.offset; + const m = new Memory(); + m.r[0] = op.sim(m, pa, pb); + + run(t, { + space, + requiresAliasing: t.params.aliased, + helpers: ` +fn store_i32(p : ${pi32}, v : i32) { + *p = v; +} + +fn load_i32(p : ${pi32}) -> i32 { + return *p; +} + +fn store_store_load(pa : ${pi32}, pb : ${pi32}, va : i32, vb : i32) -> i32 { + *pa = va; + *pb = vb; + return *pa; +} + +fn f(pa : ${pi32}, pb : ${pi32}, va : i32, vb : i32, n : i32) -> i32 { + ${op.wgsl} +}`, + body: ` + output.r[0] = f(&A.${target.wgsl}, &${t.params.aliased ? 'A' : 'B'}.${target.wgsl}, + input.p[0], input.p[1], input.p[2]);`, + expected: m.expected(), + }); + }); + +g.test('two_pointers_dynamic_index') + .desc( + `Test that a function with two i32 pointer parameters behaves correctly when the pointers are +formed from runtime indices into the same array, and those indices may or may not be equal. + +These cases always require the 'unrestricted_aliasing' language feature.` + ) + .params(u => + u + .combine('address_space', kAddressSpaces) + .combine('aliased', [true, false]) + .beginSubcases() + .combine('op', keysOf(kTwoPointerOps)) + ) + .fn(t => { + const space = t.params.address_space; + const op = kTwoPointerOps[t.params.op]; + const pi32 = ptr(space, 'i32'); + + const kArr = 2; + const pa = kArr + kIdx; + const pb = kArr + kIdx + (t.params.aliased ? 0 : 1); + const m = new Memory(); + m.r[0] = op.sim(m, pa, pb); + + run(t, { + space, + requiresAliasing: true, + helpers: ` +fn store_i32(p : ${pi32}, v : i32) { + *p = v; +} + +fn load_i32(p : ${pi32}) -> i32 { + return *p; +} + +fn store_store_load(pa : ${pi32}, pb : ${pi32}, va : i32, vb : i32) -> i32 { + *pa = va; + *pb = vb; + return *pa; +} + +fn f(pa : ${pi32}, pb : ${pi32}, va : i32, vb : i32, n : i32) -> i32 { + ${op.wgsl} +}`, + body: ` + let i = input.p[3]; + let j = input.p[3] + ${t.params.aliased ? 0 : 1}; + output.r[0] = f(&A.arr[i], &A.arr[j], input.p[0], input.p[1], input.p[2]);`, + expected: m.expected(), + }); + }); + +/** Composite types within 'Data' that pointers can be formed to. */ +const kComposites = { + array: { + type: 'array', + member: 'arr', + offset: 2, + access: (i: number | string) => `[${i}]`, + ctor: (args: string) => `array(${args})`, + }, + struct: { + type: 'S', + member: 's', + offset: 6, + access: (i: number | string) => `.${'abcd'[Number(i)]}`, + ctor: (args: string) => `S(${args})`, + }, +}; +type Composite = (typeof kComposites)[keyof typeof kComposites]; + +interface CompositeAndElementOp { + /** + * The body of 'fn f(pc : ptr, pe : ptr, va : i32, vb : i32, idx : i32) -> i32', where + * 'pc' points to a composite and 'pe' points to its element 1 (or the equivalent element of a + * different variable). + */ + wgsl: (c: Composite) => string; + /** Simulates 'f'. */ + sim: (m: Memory, pc: number, pe: number) => number; +} + +const kCompositeAndElementOps: Record = { + element_then_whole: { + wgsl: c => ` + *pe = va; + *pc = ${c.ctor('vb, vb + 1, vb + 2, vb + 3')}; + return *pe;`, + sim: (m, pc, pe) => { + m.st(pe, kVA); + for (let i = 0; i < 4; i++) { + m.st(pc + i, kVB + i); + } + return m.ld(pe); + }, + }, + whole_then_element: { + wgsl: c => ` + *pc = ${c.ctor('vb, vb + 1, vb + 2, vb + 3')}; + *pe = va; + return (*pc)${c.access(1)};`, + sim: (m, pc, pe) => { + for (let i = 0; i < 4; i++) { + m.st(pc + i, kVB + i); + } + m.st(pe, kVA); + return m.ld(pc + 1); + }, + }, + element_then_load_whole: { + wgsl: c => ` + *pe = va; + let v = *pc; + return v${c.access(1)};`, + sim: (m, pc, pe) => { + m.st(pe, kVA); + return m.ld(pc + 1); + }, + }, + sub_element_write: { + wgsl: c => ` + (*pc)${c === kComposites.array ? '[idx]' : c.access(1)} = va; + return *pe;`, + sim: (m, pc, pe) => { + m.st(pc + kIdx, kVA); + return m.ld(pe); + }, + }, + load_whole_modify_store_whole: { + // The copy 'v' must be taken before the store to '*pe'. + wgsl: c => ` + var v = *pc; + *pe = va; + v${c.access(0)} = v${c.access(1)}; + *pc = v; + return *pe;`, + sim: (m, pc, pe) => { + const v = [0, 1, 2, 3].map(i => m.ld(pc + i)); + m.st(pe, kVA); + v[0] = v[1]; + for (let i = 0; i < 4; i++) { + m.st(pc + i, v[i]); + } + return m.ld(pe); + }, + }, +}; + +g.test('composite_and_element') + .desc( + `Test that a function taking a pointer to a composite and a pointer to an i32 behaves correctly +when the i32 pointer points to an element of the composite.` + ) + .params(u => + u + .combine('address_space', kAddressSpaces) + .combine('composite', keysOf(kComposites)) + .combine('aliased', [true, false]) + .beginSubcases() + .combine('op', keysOf(kCompositeAndElementOps)) + ) + .fn(t => { + const space = t.params.address_space; + const c = kComposites[t.params.composite]; + const op = kCompositeAndElementOps[t.params.op]; + + const pc = c.offset; + const pe = (t.params.aliased ? 0 : kB) + c.offset + 1; + const m = new Memory(); + m.r[0] = op.sim(m, pc, pe); + + run(t, { + space, + requiresAliasing: t.params.aliased, + helpers: ` +fn f(pc : ${ptr(space, c.type)}, pe : ${ptr(space, 'i32')}, va : i32, vb : i32, idx : i32) -> i32 { + ${op.wgsl(c)} +}`, + body: ` + output.r[0] = f(&A.${c.member}, &${t.params.aliased ? 'A' : 'B'}.${c.member}${c.access(1)}, + input.p[0], input.p[1], input.p[3]);`, + expected: m.expected(), + }); + }); + +interface CompositeCopyOp { + /** The body of 'fn f(pd : ptr, ps : ptr) -> i32'. */ + wgsl: (c: Composite) => string; + /** Simulates 'f'. */ + sim: (m: Memory, pd: number, ps: number) => number; +} + +const kCompositeCopyOps: Record = { + reverse_construct: { + // All loads of '*ps' happen before the store to '*pd'. + wgsl: c => ` + *pd = ${c.ctor([3, 2, 1, 0].map(i => `(*ps)${c.access(i)}`).join(', '))}; + return (*pd)${c.access(0)};`, + sim: (m, pd, ps) => { + const v = [3, 2, 1, 0].map(i => m.ld(ps + i)); + for (let i = 0; i < 4; i++) { + m.st(pd + i, v[i]); + } + return m.ld(pd); + }, + }, + memberwise_reverse: { + // Loads and stores are interleaved. + wgsl: c => + [0, 1, 2, 3].map(i => `\n (*pd)${c.access(i)} = (*ps)${c.access(3 - i)};`).join('') + + `\n return (*pd)${c.access(0)};`, + sim: (m, pd, ps) => { + for (let i = 0; i < 4; i++) { + m.st(pd + i, m.ld(ps + 3 - i)); + } + return m.ld(pd); + }, + }, + copy_then_modify: { + wgsl: c => ` + *pd = *ps; + (*pd)${c.access(0)} += 1; + return (*ps)${c.access(0)};`, + sim: (m, pd, ps) => { + for (let i = 0; i < 4; i++) { + m.st(pd + i, m.ld(ps + i)); + } + m.st(pd, m.ld(pd) + 1); + return m.ld(ps); + }, + }, +}; + +g.test('composite_copy') + .desc( + `Test that a function that copies between two composite pointers behaves correctly when both +pointers refer to the same composite.` + ) + .params(u => + u + .combine('address_space', kAddressSpaces) + .combine('composite', keysOf(kComposites)) + .combine('aliased', [true, false]) + .beginSubcases() + .combine('op', keysOf(kCompositeCopyOps)) + ) + .fn(t => { + const space = t.params.address_space; + const c = kComposites[t.params.composite]; + const op = kCompositeCopyOps[t.params.op]; + const pc = ptr(space, c.type); + + const pd = c.offset; + const ps = (t.params.aliased ? 0 : kB) + c.offset; + const m = new Memory(); + m.r[0] = op.sim(m, pd, ps); + + run(t, { + space, + requiresAliasing: t.params.aliased, + helpers: ` +fn f(pd : ${pc}, ps : ${pc}) -> i32 { + ${op.wgsl(c)} +}`, + body: ` + output.r[0] = f(&A.${c.member}, &${t.params.aliased ? 'A' : 'B'}.${c.member});`, + expected: m.expected(), + }); + }); + +interface ModuleScopeOp { + /** + * The body of 'fn f(p : ptr, va : i32, vb : i32, n : i32) -> i32', where 'global' is an + * expression that directly accesses the module-scope variable 'A'. + */ + wgsl: (global: string) => string; + /** Simulates 'f', where 'global' is the offset of the memory accessed by the 'global' expression. */ + sim: (m: Memory, p: number, global: number) => number; +} + +const kModuleScopeOps: Record = { + ptr_store_global_store_ptr_load: { + wgsl: g => ` + *p = va; + ${g} = vb; + return *p;`, + sim: (m, p, g) => { + m.st(p, kVA); + m.st(g, kVB); + return m.ld(p); + }, + }, + global_store_ptr_store_global_load: { + wgsl: g => ` + ${g} = va; + *p = vb; + return ${g};`, + sim: (m, p, g) => { + m.st(g, kVA); + m.st(p, kVB); + return m.ld(g); + }, + }, + ptr_store_global_load: { + wgsl: g => ` + *p = va; + return ${g};`, + sim: (m, p, g) => { + m.st(p, kVA); + return m.ld(g); + }, + }, + loop_accumulate: { + wgsl: g => ` + for (var i = 0; i < n; i++) { + ${g} += *p; + } + return ${g};`, + sim: (m, p, g) => { + for (let i = 0; i < kN; i++) { + m.st(g, m.ld(g) + m.ld(p)); + } + return m.ld(g); + }, + }, + whole_global_store: { + wgsl: _ => ` + *p = va; + A = input.b; + return *p;`, + sim: (m, p, _) => { + m.st(p, kVA); + for (let i = 0; i < kDataSize; i++) { + m.st(i, kInitB[i]); + } + return m.ld(p); + }, + }, +}; + +g.test('one_pointer_one_module_scope') + .desc( + `Test that a function with a pointer parameter behaves correctly when the pointer refers to a +module-scope variable that is also directly accessed by the function.` + ) + .params(u => + u + .combine('address_space', kModuleScopeAddressSpaces) + .combine('target', keysOf(kScalarTargets)) + .combine('aliased', [true, false]) + .beginSubcases() + .combine('op', keysOf(kModuleScopeOps)) + ) + .fn(t => { + const space = t.params.address_space; + const target = kScalarTargets[t.params.target]; + const op = kModuleScopeOps[t.params.op]; + + const p = (t.params.aliased ? 0 : kB) + target.offset; + const m = new Memory(); + m.r[0] = op.sim(m, p, target.offset); + + run(t, { + space, + requiresAliasing: t.params.aliased, + helpers: ` +fn f(p : ${ptr(space, 'i32')}, va : i32, vb : i32, n : i32) -> i32 { + ${op.wgsl(`A.${target.wgsl}`)} +}`, + body: ` + output.r[0] = f(&${t.params.aliased ? 'A' : 'B'}.${target.wgsl}, + input.p[0], input.p[1], input.p[2]);`, + expected: m.expected(), + }); + }); + +interface AtomicOp { + /** The body of 'fn f(pa : ptr>, pb : ptr>, va : i32, vb : i32) -> i32'. */ + wgsl: string; + /** Simulates 'f'. */ + sim: (m: Memory, pa: number, pb: number) => number; +} + +const kAtomicOps: Record = { + add_add_load: { + wgsl: ` + atomicAdd(pa, va); + atomicAdd(pb, vb); + return atomicLoad(pa);`, + sim: (m, pa, pb) => { + m.st(pa, m.ld(pa) + kVA); + m.st(pb, m.ld(pb) + kVB); + return m.ld(pa); + }, + }, + store_exchange: { + wgsl: ` + atomicStore(pa, va); + return atomicExchange(pb, vb);`, + sim: (m, pa, pb) => { + m.st(pa, kVA); + const old = m.ld(pb); + m.st(pb, kVB); + return old; + }, + }, + exchange_load: { + wgsl: ` + let old = atomicExchange(pa, va); + return old + atomicLoad(pb);`, + sim: (m, pa, pb) => { + const old = m.ld(pa); + m.st(pa, kVA); + return old + m.ld(pb); + }, + }, +}; + +/** Atomic i32 memory locations within 'AtomicData' that pointers can be formed to. */ +const kAtomicTargets = { + scalar: { wgsl: 'x', offset: 0 }, + array_element: { wgsl: 'arr[1]', offset: 3 }, +}; + +g.test('two_atomic_pointers') + .desc( + `Test that a function with two atomic pointer parameters behaves correctly when both pointers +refer to the same atomic.` + ) + .params(u => + u + .combine('address_space', kAtomicAddressSpaces) + .combine('target', keysOf(kAtomicTargets)) + .combine('aliased', [true, false]) + .beginSubcases() + .combine('op', keysOf(kAtomicOps)) + ) + .fn(t => { + const space = t.params.address_space; + const target = kAtomicTargets[t.params.target]; + const op = kAtomicOps[t.params.op]; + const patomic = ptr(space, 'atomic'); + + const pa = target.offset; + const pb = (t.params.aliased ? 0 : kB) + target.offset; + const m = new Memory(); + m.r[0] = op.sim(m, pa, pb); + // 's' is not present in 'AtomicData', so is not written to the output. + for (const base of [0, kB]) { + for (let i = 6; i < kDataSize; i++) { + m.mem[base + i] = 0; + } + } + + run(t, { + space, + requiresAliasing: t.params.aliased, + atomic: true, + helpers: ` +fn f(pa : ${patomic}, pb : ${patomic}, va : i32, vb : i32) -> i32 { + ${op.wgsl} +}`, + body: ` + atomicStore(&A.x, input.a.x); + atomicStore(&A.y, input.a.y); + atomicStore(&B.x, input.b.x); + atomicStore(&B.y, input.b.y); + for (var i = 0; i < 4; i++) { + atomicStore(&A.arr[i], input.a.arr[i]); + atomicStore(&B.arr[i], input.b.arr[i]); + } + + output.r[0] = f(&A.${target.wgsl}, &${t.params.aliased ? 'A' : 'B'}.${target.wgsl}, + input.p[0], input.p[1]); + + output.a.x = atomicLoad(&A.x); + output.a.y = atomicLoad(&A.y); + output.b.x = atomicLoad(&B.x); + output.b.y = atomicLoad(&B.y); + for (var i = 0; i < 4; i++) { + output.a.arr[i] = atomicLoad(&A.arr[i]); + output.b.arr[i] = atomicLoad(&B.arr[i]); + }`, + expected: m.expected(), + }); + }); + +interface LoopCarriedOp { + /** + * The body of 'fn f(pd : ptr>, ps : ptr>, m : i32) -> i32'. + * When 'pd' and 'ps' alias, each iteration depends on a value written by an earlier iteration, + * so the loop must not be vectorized as if the arrays were disjoint. + */ + wgsl: string; + /** Simulates 'f'. */ + sim: (m: Memory, pd: number, ps: number) => number; +} + +const kLoopCarriedOps: Record = { + shift_up: { + // Aliased: ps[0] is propagated to every element. + wgsl: ` + for (var i = 0; i < m; i++) { + (*pd)[i + 1] = (*ps)[i]; + } + return (*pd)[m];`, + sim: (mem, pd, ps) => { + for (let i = 0; i < kM; i++) { + mem.st(pd + i + 1, mem.ld(ps + i)); + } + return mem.ld(pd + kM); + }, + }, + shift_up_by_two: { + // Aliased: a dependence distance of 2, which is less than typical vector widths. + wgsl: ` + for (var i = 0; i < m - 1; i++) { + (*pd)[i + 2] = (*ps)[i] * 2; + } + return (*pd)[m];`, + sim: (mem, pd, ps) => { + for (let i = 0; i < kM - 1; i++) { + mem.st(pd + i + 2, mem.ld(ps + i) * 2); + } + return mem.ld(pd + kM); + }, + }, + running_sum: { + // Aliased: a prefix-sum style recurrence. + wgsl: ` + for (var i = 0; i < m; i++) { + (*pd)[i + 1] = (*ps)[i] + (*ps)[i + 1]; + } + return (*pd)[m];`, + sim: (mem, pd, ps) => { + for (let i = 0; i < kM; i++) { + mem.st(pd + i + 1, mem.ld(ps + i) + mem.ld(ps + i + 1)); + } + return mem.ld(pd + kM); + }, + }, + reverse: { + // Aliased: the second half reads values already written by the first half, producing a + // palindrome rather than a reversal. + wgsl: ` + for (var i = 0; i <= m; i++) { + (*pd)[i] = (*ps)[m - i]; + } + return (*pd)[0];`, + sim: (mem, pd, ps) => { + for (let i = 0; i <= kM; i++) { + mem.st(pd + i, mem.ld(ps + kM - i)); + } + return mem.ld(pd); + }, + }, +}; + +g.test('loop_carried_dependence') + .desc( + `Test that loops over two array pointers behave correctly when both pointers refer to the same +array, creating a loop-carried dependence. The loop count is a runtime value large enough to make +vectorization attractive.` + ) + .params(u => + u + .combine('address_space', kAddressSpaces) + .combine('aliased', [true, false]) + .beginSubcases() + .combine('op', keysOf(kLoopCarriedOps)) + ) + .fn(t => { + const space = t.params.address_space; + const op = kLoopCarriedOps[t.params.op]; + const parr = ptr(space, `array`); + + const pd = kBig; + const ps = (t.params.aliased ? 0 : kB) + kBig; + const m = new Memory(); + m.r[0] = op.sim(m, pd, ps); + + run(t, { + space, + requiresAliasing: t.params.aliased, + helpers: ` +fn f(pd : ${parr}, ps : ${parr}, m : i32) -> i32 { + ${op.wgsl} +}`, + body: ` + output.r[0] = f(&A.big, &${t.params.aliased ? 'A' : 'B'}.big, input.p[4]);`, + expected: m.expected(), + }); + }); + +g.test('buffer_view_mixed_types') + .desc( + `Test that a function taking two differently typed pointer parameters behaves correctly when +the pointers are views of overlapping bytes of a buffer, formed with bufferView or bufferArrayView. + +The views are formed in the entry point and passed to the function. This is the pointer parameter +counterpart of the bufferView and bufferArrayView 'mixed_types_aliasing' tests. + +Only the integer view is loaded from, so that float denormal flushing and NaN canonicalization +cannot affect the results. + + * 'pair' selects the types of the two views, covering combinations of scalars, vectors, and + structures (with and without padding), including f16 types. + * 'overlap' selects whether the views start at the same byte, partially overlap, or only overlap + padding. + * 'view' selects how the views are formed: + - 'constant_offset': bufferView with a constant byte offset. + - 'dynamic_offset': bufferView with a runtime byte offset. + - 'dynamic_index': an element of bufferArrayView, with a runtime element index.` + ) + .params(u => + u + .combine('buffer', kMixedTypeBuffers) + .combine('pair', keysOf(kMixedTypePairs)) + .combine('overlap', kMixedTypeOverlaps) + .filter(t => { + const pair = kMixedTypePairs[t.pair]; + return mixedTypeWords(pair.int, pair.other, t.overlap) !== undefined; + }) + .combine('view', ['constant_offset', 'dynamic_offset', 'dynamic_index'] as const) + .combine('aliased', [true, false]) + .beginSubcases() + .combine('op', keysOf(kMixedTypeOps)) + ) + .fn(t => { + //if (t.params.aliased) { + // t.skipIfLanguageFeatureNotSupported('unrestricted_aliasing'); + //} + const pair = kMixedTypePairs[t.params.pair]; + const words = mixedTypeWords(pair.int, pair.other, t.params.overlap)!; + // The dynamic offset adds 'input.p[3] * 16' bytes, which preserves the alignment of all types. + const dynamicWords = 4 * kMixedTypeIdx; + const view = (type: MixedType, buffer: string, word: number) => { + const ty = mixedTypeName(type); + switch (t.params.view) { + case 'constant_offset': + return `bufferView<${ty}>(&${buffer}, ${word * 4}u)`; + case 'dynamic_offset': + return `bufferView<${ty}>(&${buffer}, ${ + (word - dynamicWords) * 4 + }u + u32(input.p[3]) * 16u)`; + case 'dynamic_index': { + // The array starts one element before the target. All of the tested types have an + // element stride equal to their size. Each view spans 32 bytes, which keeps all views + // within the buffer. + const base = word * 4 - type.size * kMixedTypeIdx; + return `&(*bufferArrayView>(&${buffer}, ${base}u, 32u))[input.p[3]]`; + } + } + }; + runMixedTypeAliasingTest(t, { + buffer: t.params.buffer, + int: pair.int, + other: pair.other, + intWord: words.intWord, + otherWord: words.otherWord, + view, + aliased: t.params.aliased, + op: kMixedTypeOps[t.params.op], + pointerParams: true, + directives: t.params.aliased ? '// requires unrestricted_aliasing;' : '', + }); + }); diff --git a/src/webgpu/shader/validation/functions/alias_analysis.spec.ts b/src/webgpu/shader/validation/functions/alias_analysis.spec.ts index b428e4ee9584..c53c74fa6c01 100644 --- a/src/webgpu/shader/validation/functions/alias_analysis.spec.ts +++ b/src/webgpu/shader/validation/functions/alias_analysis.spec.ts @@ -32,10 +32,24 @@ const kUses: Record = { type UseName = keyof typeof kUses; -function shouldPass(aliased: boolean, ...uses: UseName[]): boolean { +/** + * @returns true if aliased pointer arguments with a write access are a shader-creation error. + * This is the case unless the 'unrestricted_aliasing' language feature is supported. + */ +function aliasingRestricted(t: ShaderValidationTest): boolean { + return !t.hasLanguageFeature('unrestricted_aliasing'); +} + +function shouldPass(t: ShaderValidationTest, aliased: boolean, ...uses: UseName[]): boolean { // Expect fail if the pointers are aliased and at least one of the accesses is a write. // If either of the accesses is a "no access" then expect pass. - return !aliased || !uses.some(u => kUses[u].is_write) || uses.includes('no_access'); + // If the 'unrestricted_aliasing' language feature is supported then always expect pass. + return ( + !aliasingRestricted(t) || + !aliased || + !uses.some(u => kUses[u].is_write) || + uses.includes('no_access') + ); } type AddressSpace = 'private' | 'function' | 'storage' | 'uniform' | 'workgroup'; @@ -131,7 +145,7 @@ fn caller() { callee(&x, ${t.params.aliased ? `&x` : `&y`}); } `; - t.expectCompileResult(shouldPass(t.params.aliased, t.params.a_use, t.params.b_use), code); + t.expectCompileResult(shouldPass(t, t.params.aliased, t.params.a_use, t.params.b_use), code); }); g.test('two_pointers_to_array_elements') @@ -165,7 +179,7 @@ fn caller() { callee(&x[${t.params.index}], ${t.params.aliased ? `&x[0]` : `&y[0]`}); } `; - t.expectCompileResult(shouldPass(t.params.aliased, t.params.a_use, t.params.b_use), code); + t.expectCompileResult(shouldPass(t, t.params.aliased, t.params.a_use, t.params.b_use), code); }); g.test('two_pointers_to_array_elements_indirect') @@ -207,7 +221,7 @@ fn caller() { index(&x, ${t.params.aliased ? `&x` : `&y`}); } `; - t.expectCompileResult(shouldPass(t.params.aliased, t.params.a_use, t.params.b_use), code); + t.expectCompileResult(shouldPass(t, t.params.aliased, t.params.a_use, t.params.b_use), code); }); g.test('two_pointers_to_struct_members') @@ -246,7 +260,7 @@ fn caller() { callee(&x.${t.params.member}, ${t.params.aliased ? `&x.a` : `&y.a`}); } `; - t.expectCompileResult(shouldPass(t.params.aliased, t.params.a_use, t.params.b_use), code); + t.expectCompileResult(shouldPass(t, t.params.aliased, t.params.a_use, t.params.b_use), code); }); g.test('two_pointers_to_struct_members_indirect') @@ -293,7 +307,7 @@ fn caller() { access(&x, ${t.params.aliased ? `&x` : `&y`}); } `; - t.expectCompileResult(shouldPass(t.params.aliased, t.params.a_use, t.params.b_use), code); + t.expectCompileResult(shouldPass(t, t.params.aliased, t.params.a_use, t.params.b_use), code); }); g.test('one_pointer_one_module_scope') @@ -325,7 +339,7 @@ fn caller() { callee(${t.params.aliased ? `&x` : `&y`}); } `; - t.expectCompileResult(shouldPass(t.params.aliased, t.params.a_use, t.params.b_use), code); + t.expectCompileResult(shouldPass(t, t.params.aliased, t.params.a_use, t.params.b_use), code); }); g.test('subcalls') @@ -371,7 +385,7 @@ fn caller() { callee(&x, ${t.params.aliased ? `&x` : `&y`}); } `; - t.expectCompileResult(shouldPass(t.params.aliased, t.params.a_use, t.params.b_use), code); + t.expectCompileResult(shouldPass(t, t.params.aliased, t.params.a_use, t.params.b_use), code); }); g.test('member_accessors') @@ -406,7 +420,7 @@ fn caller() { callee(&x, ${t.params.aliased ? `&x` : `&y`}); } `; - t.expectCompileResult(shouldPass(t.params.aliased, t.params.a_use, t.params.b_use), code); + t.expectCompileResult(shouldPass(t, t.params.aliased, t.params.a_use, t.params.b_use), code); }); g.test('swizzles') @@ -442,7 +456,7 @@ fn caller() { callee(&x, ${t.params.aliased ? `&x` : `&y`}); } `; - t.expectCompileResult(shouldPass(t.params.aliased, t.params.a_use, 'let_init'), code); + t.expectCompileResult(shouldPass(t, t.params.aliased, t.params.a_use, 'let_init'), code); }); g.test('same_pointer_read_and_write') @@ -572,7 +586,9 @@ fn caller() { } `; const shouldFail = - t.params.aliased && (isWrite(t.params.builtin_a) || isWrite(t.params.builtin_b)); + aliasingRestricted(t) && + t.params.aliased && + (isWrite(t.params.builtin_a) || isWrite(t.params.builtin_b)); t.expectCompileResult(!shouldFail, code); }); @@ -606,7 +622,9 @@ fn caller() { } `; const shouldFail = - t.params.aliased && (isWrite(t.params.builtin_a) || isWrite(t.params.builtin_b)); + aliasingRestricted(t) && + t.params.aliased && + (isWrite(t.params.builtin_a) || isWrite(t.params.builtin_b)); t.expectCompileResult(!shouldFail, code); }); @@ -645,7 +663,9 @@ fn caller() { } `; const shouldFail = - t.params.aliased && (isWrite(t.params.builtin_a) || isWrite(t.params.builtin_b)); + aliasingRestricted(t) && + t.params.aliased && + (isWrite(t.params.builtin_a) || isWrite(t.params.builtin_b)); t.expectCompileResult(!shouldFail, code); }); @@ -678,7 +698,9 @@ fn caller() { } `; const shouldFail = - t.params.aliased && (isWrite(t.params.builtin_a) || isWrite(t.params.builtin_b)); + aliasingRestricted(t) && + t.params.aliased && + (isWrite(t.params.builtin_a) || isWrite(t.params.builtin_b)); t.expectCompileResult(!shouldFail, code); }); @@ -718,6 +740,40 @@ fn caller() { callee(&x, &${t.params.aliased ? 'x' : 'y'}); } `; - const shouldFail = t.params.aliased && t.params.use === 'store'; + const shouldFail = aliasingRestricted(t) && t.params.aliased && t.params.use === 'store'; t.expectCompileResult(!shouldFail, code); }); + +g.test('requires_unrestricted_aliasing') + .desc( + `Test that aliased pointer arguments with write accesses are valid with a +'requires unrestricted_aliasing' directive, iff the language feature is supported.` + ) + .params(u => u.combine('address_space', kWritableAddressSpaces).combine('aliased', [true, false])) + .fn(t => { + if (requiresUnrestrictedPointerParameters(t.params.address_space)) { + t.skipIfLanguageFeatureNotSupported('unrestricted_pointer_parameters'); + } + + const code = ` +requires unrestricted_aliasing; + +${maybeDeclareModuleScopeVar('x', t.params.address_space, 'i32')} +${maybeDeclareModuleScopeVar('y', t.params.address_space, 'i32')} + +fn callee(pa : ${ptr(t.params.address_space, 'i32')}, + pb : ${ptr(t.params.address_space, 'i32')}) -> i32 { + *pa = 1; + *pb = 2; + return *pa; +} + +fn caller() { + ${maybeDeclareFunctionScopeVar('x', t.params.address_space, 'i32')} + ${maybeDeclareFunctionScopeVar('y', t.params.address_space, 'i32')} + callee(&x, ${t.params.aliased ? `&x` : `&y`}); +} +`; + // The 'requires' directive is an error if the feature is not supported, regardless of aliasing. + t.expectCompileResult(t.hasLanguageFeature('unrestricted_aliasing'), code); + });