diff options
Diffstat (limited to 'dom/webgpu/tests/cts/checkout/src/webgpu/shader/execution/memory_model/memory_model_setup.ts')
| -rw-r--r-- | dom/webgpu/tests/cts/checkout/src/webgpu/shader/execution/memory_model/memory_model_setup.ts | 341 |
1 files changed, 271 insertions, 70 deletions
diff --git a/dom/webgpu/tests/cts/checkout/src/webgpu/shader/execution/memory_model/memory_model_setup.ts b/dom/webgpu/tests/cts/checkout/src/webgpu/shader/execution/memory_model/memory_model_setup.ts index f8e5b9034c..8dee32b72d 100644 --- a/dom/webgpu/tests/cts/checkout/src/webgpu/shader/execution/memory_model/memory_model_setup.ts +++ b/dom/webgpu/tests/cts/checkout/src/webgpu/shader/execution/memory_model/memory_model_setup.ts @@ -1,5 +1,7 @@ -import { GPUTest } from '../../../gpu_test'; +import { GPUTest } from '../../../gpu_test.js'; import { checkElementsPassPredicate } from '../../../util/check_contents.js'; +import { align } from '../../../util/math.js'; +import { PRNG } from '../../../util/prng.js'; /* All buffer sizes are counted in units of 4-byte words. */ @@ -15,6 +17,9 @@ import { checkElementsPassPredicate } from '../../../util/check_contents.js'; export type AccessValueType = 'f16' | 'u32'; export const kAccessValueTypes = ['f16', 'u32'] as const; +/** The width used for textures (default compat limit in WebGPU). */ +const kWidth = 4096; + /* Parameter values are set heuristically, typically by a time-intensive search. */ export type MemoryModelTestParams = { /* Number of invocations per workgroup. The workgroups are 1-dimensional. */ @@ -76,12 +81,33 @@ const numReadOutputs = 2; type BufferWithSource = { /** Buffer used by shader code. */ deviceBuf: GPUBuffer; - /** Buffer populated from the host size, data is copied to device buffer for use by shader. */ + /** Buffer populated from the host side, data is copied to device buffer for use by shader. */ srcBuf: GPUBuffer; /** Size in bytes of the buffer. */ size: number; }; +/** Represents a device texture and a utility buffer for resetting memory and copying parameters. */ +type TextureWithSource = { + /** Texture used by shader code. */ + deviceTex: GPUTexture; + /** Buffer populated from the host side, data is copied to device buffer for use by shader. */ + srcBuf: GPUBuffer; + /** Size in bytes of the buffer. */ + size: number; +}; + +type SubBufferWithSource = { + /** Buffer used by shader code. This buffer is shared for multiple used */ + deviceBuf: GPUBuffer; + /** Buffer populated from the host side, data is copied to device buffer for use by shader. */ + srcBuf: GPUBuffer; + /** Size in bytes of this portion of the buffer. */ + size: number; + /** Offset in bytes of this portion of the buffer */ + offset: number; +}; + /** Specifies the buffers used during a memory model test. */ type MemoryModelBuffers = { /** This is the memory region that testing threads read from and write to. */ @@ -102,6 +128,10 @@ type MemoryModelBuffers = { stressParams: BufferWithSource; }; +type MemoryModelTextures = { + testLocations: TextureWithSource; +}; + /** The number of stress params to add to the stress params buffer. */ const numStressParams = 12; const barrierParamIndex = 0; @@ -128,11 +158,11 @@ const bytesPerWord = 4; * - enable directives, if necessary * - the type alias for AccessValueType */ -function shaderPreamble(accessValueType: AccessValueType): string { +function shaderPreamble(accessValueType: AccessValueType, constants: string): string { if (accessValueType === 'f16') { - return 'enable f16;\nalias AccessValueTy = f16;\n'; + return `enable f16;\nalias AccessValueTy = f16;\n${constants}\n`; } - return `alias AccessValueTy = ${accessValueType};\n`; + return `alias AccessValueTy = ${accessValueType};\n${constants}\n`; } /** @@ -175,10 +205,14 @@ export class MemoryModelTester { protected test: GPUTest; protected params: MemoryModelTestParams; protected buffers: MemoryModelBuffers; + protected textures: MemoryModelTextures | undefined; protected testPipeline: GPUComputePipeline; protected testBindGroup: GPUBindGroup; + protected textureBindGroup: GPUBindGroup | undefined; protected resultPipeline: GPUComputePipeline; protected resultBindGroup: GPUBindGroup; + protected prng: PRNG; + protected useTexture: boolean; /** Sets up a memory model test by initializing buffers and pipeline layouts. */ constructor( @@ -186,24 +220,36 @@ export class MemoryModelTester { params: MemoryModelTestParams, testShader: string, resultShader: string, - accessValueType: AccessValueType = 'u32' + accessValueType: AccessValueType = 'u32', + useTexture: boolean = false ) { + this.prng = new PRNG(1); this.test = t; this.params = params; - - testShader = shaderPreamble(accessValueType) + testShader; - resultShader = shaderPreamble(accessValueType) + resultShader; + this.useTexture = useTexture; + + const workgroupXSize = Math.min(params.workgroupSize, t.device.limits.maxComputeWorkgroupSizeX); + const constants = ` + const kNumBarriers = 1u; // MAINTENANCE_TODO: make barrier not an array + const kMaxWorkgroups = ${params.maxWorkgroups}u; + const kScratchMemorySize = ${params.scratchMemorySize}u; + const kWorkgroupXSize = ${workgroupXSize}u; + `; + testShader = shaderPreamble(accessValueType, constants) + testShader; + resultShader = shaderPreamble(accessValueType, constants) + resultShader; // set up buffers - const testingThreads = this.params.workgroupSize * this.params.testingWorkgroups; + const testingThreads = workgroupXSize * this.params.testingWorkgroups; const testLocationsSize = testingThreads * numMemLocations * this.params.memStride * bytesPerWord; const testLocationsBuffer: BufferWithSource = { deviceBuf: this.test.device.createBuffer({ + label: 'testLocationsBuffer', size: testLocationsSize, usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE, }), srcBuf: this.test.device.createBuffer({ + label: 'testLocationsSrcBuf', size: testLocationsSize, usage: GPUBufferUsage.COPY_SRC, }), @@ -213,10 +259,12 @@ export class MemoryModelTester { const readResultsSize = testingThreads * numReadOutputs * bytesPerWord; const readResultsBuffer: BufferWithSource = { deviceBuf: this.test.device.createBuffer({ + label: 'readResultsBuffer', size: readResultsSize, usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE, }), srcBuf: this.test.device.createBuffer({ + label: 'readResultsSrcBuf', size: readResultsSize, usage: GPUBufferUsage.COPY_SRC, }), @@ -226,10 +274,12 @@ export class MemoryModelTester { const testResultsSize = this.params.numBehaviors * bytesPerWord; const testResultsBuffer: BufferWithSource = { deviceBuf: this.test.device.createBuffer({ + label: 'testResultsBuffer', size: testResultsSize, usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, }), srcBuf: this.test.device.createBuffer({ + label: 'testResultsSrcBuffer', size: testResultsSize, usage: GPUBufferUsage.COPY_SRC, }), @@ -249,52 +299,87 @@ export class MemoryModelTester { size: shuffledWorkgroupsSize, }; - const barrierSize = bytesPerWord; - const barrierBuffer: BufferWithSource = { - deviceBuf: this.test.device.createBuffer({ - size: barrierSize, - usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE, - }), + if (this.useTexture) { + const numTexels = testLocationsSize / bytesPerWord; + const width = kWidth; + const height = numTexels / width; + const textureSize: GPUExtent3D = { width, height }; + const textureLocations: TextureWithSource = { + deviceTex: this.test.device.createTexture({ + format: 'r32uint', + dimension: '2d', + size: textureSize, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.STORAGE_BINDING, + }), + srcBuf: testLocationsBuffer.srcBuf, + size: testLocationsSize, + }; + this.textures = { + testLocations: textureLocations, + }; + } + + // Combine 3 arrays into 1 buffer as we need to keep the number of storage buffers to 4 for compat. + const falseSharingAvoidanceQuantum = 4096; + const barrierSize = align(bytesPerWord, falseSharingAvoidanceQuantum); + const scratchpadSize = align( + this.params.scratchMemorySize * bytesPerWord, + falseSharingAvoidanceQuantum + ); + const scratchMemoryLocationsSize = align( + this.params.maxWorkgroups * bytesPerWord, + falseSharingAvoidanceQuantum + ); + const comboSize = barrierSize + scratchpadSize + scratchMemoryLocationsSize; + + const comboBuffer = this.test.device.createBuffer({ + label: 'comboBuffer', + size: comboSize, + usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE, + }); + + const barrierBuffer: SubBufferWithSource = { + deviceBuf: comboBuffer, srcBuf: this.test.device.createBuffer({ + label: 'barrierSrcBuf', size: barrierSize, usage: GPUBufferUsage.COPY_SRC, }), size: barrierSize, + offset: 0, }; - const scratchpadSize = this.params.scratchMemorySize * bytesPerWord; - const scratchpadBuffer: BufferWithSource = { - deviceBuf: this.test.device.createBuffer({ - size: scratchpadSize, - usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE, - }), + const scratchpadBuffer: SubBufferWithSource = { + deviceBuf: comboBuffer, srcBuf: this.test.device.createBuffer({ + label: 'scratchpadSrcBuf', size: scratchpadSize, usage: GPUBufferUsage.COPY_SRC, }), size: scratchpadSize, + offset: barrierSize, }; - const scratchMemoryLocationsSize = this.params.maxWorkgroups * bytesPerWord; - const scratchMemoryLocationsBuffer: BufferWithSource = { - deviceBuf: this.test.device.createBuffer({ - size: scratchMemoryLocationsSize, - usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE, - }), + const scratchMemoryLocationsBuffer: SubBufferWithSource = { + deviceBuf: comboBuffer, srcBuf: this.test.device.createBuffer({ + label: 'scratchMemoryLocationsSrcBuf', size: scratchMemoryLocationsSize, usage: GPUBufferUsage.COPY_SRC | GPUBufferUsage.MAP_WRITE, }), size: scratchMemoryLocationsSize, + offset: barrierSize + scratchpadSize, }; const stressParamsSize = numStressParams * bytesPerWord; const stressParamsBuffer: BufferWithSource = { deviceBuf: this.test.device.createBuffer({ + label: 'stressParamsBuffer', size: stressParamsSize, usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.UNIFORM, }), srcBuf: this.test.device.createBuffer({ + label: 'stressParamsSrcBuf', size: stressParamsSize, usage: GPUBufferUsage.COPY_SRC | GPUBufferUsage.MAP_WRITE, }), @@ -314,19 +399,50 @@ export class MemoryModelTester { // set up pipeline layouts const testLayout = this.test.device.createBindGroupLayout({ + label: 'testLayout', entries: [ { binding: 0, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } }, { binding: 1, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } }, { binding: 2, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'read-only-storage' } }, { binding: 3, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } }, - { binding: 4, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } }, - { binding: 5, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } }, - { binding: 6, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'uniform' } }, + { binding: 4, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'uniform' } }, ], }); + + let layouts: GPUBindGroupLayout[] = [testLayout]; + if (this.useTexture) { + const textureLayout = this.test.device.createBindGroupLayout({ + label: 'textureLayout', + entries: [ + { + binding: 0, + visibility: GPUShaderStage.COMPUTE, + storageTexture: { + access: 'read-write', + format: 'r32uint', + viewDimension: '2d', + }, + }, + ], + }); + layouts = [testLayout, textureLayout]; + + const texLocations = (this.textures as MemoryModelTextures).testLocations.deviceTex; + this.textureBindGroup = this.test.device.createBindGroup({ + label: 'textureBindGroup', + entries: [ + { + binding: 0, + resource: texLocations.createView(), + }, + ], + layout: textureLayout, + }); + } this.testPipeline = this.test.device.createComputePipeline({ + label: 'testPipeline', layout: this.test.device.createPipelineLayout({ - bindGroupLayouts: [testLayout], + bindGroupLayouts: layouts, }), compute: { module: this.test.device.createShaderModule({ @@ -336,19 +452,19 @@ export class MemoryModelTester { }, }); this.testBindGroup = this.test.device.createBindGroup({ + label: 'testBindGroup', entries: [ { binding: 0, resource: { buffer: this.buffers.testLocations.deviceBuf } }, { binding: 1, resource: { buffer: this.buffers.readResults.deviceBuf } }, { binding: 2, resource: { buffer: this.buffers.shuffledWorkgroups.deviceBuf } }, - { binding: 3, resource: { buffer: this.buffers.barrier.deviceBuf } }, - { binding: 4, resource: { buffer: this.buffers.scratchpad.deviceBuf } }, - { binding: 5, resource: { buffer: this.buffers.scratchMemoryLocations.deviceBuf } }, - { binding: 6, resource: { buffer: this.buffers.stressParams.deviceBuf } }, + { binding: 3, resource: { buffer: comboBuffer } }, + { binding: 4, resource: { buffer: this.buffers.stressParams.deviceBuf } }, ], layout: testLayout, }); const resultLayout = this.test.device.createBindGroupLayout({ + label: 'resultLayout', entries: [ { binding: 0, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } }, { binding: 1, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } }, @@ -357,6 +473,7 @@ export class MemoryModelTester { ], }); this.resultPipeline = this.test.device.createComputePipeline({ + label: 'resultPipeline', layout: this.test.device.createPipelineLayout({ bindGroupLayouts: [resultLayout], }), @@ -368,6 +485,7 @@ export class MemoryModelTester { }, }); this.resultBindGroup = this.test.device.createBindGroup({ + label: 'resultBindGroup', entries: [ { binding: 0, resource: { buffer: this.buffers.testLocations.deviceBuf } }, { binding: 1, resource: { buffer: this.buffers.readResults.deviceBuf } }, @@ -402,10 +520,16 @@ export class MemoryModelTester { this.copyBufferToBuffer(encoder, this.buffers.scratchpad); this.copyBufferToBuffer(encoder, this.buffers.scratchMemoryLocations); this.copyBufferToBuffer(encoder, this.buffers.stressParams); + if (this.useTexture) { + this.copyBufferToTexture(encoder, (this.textures as MemoryModelTextures).testLocations); + } const testPass = encoder.beginComputePass(); testPass.setPipeline(this.testPipeline); testPass.setBindGroup(0, this.testBindGroup); + if (this.useTexture) { + testPass.setBindGroup(1, this.textureBindGroup as GPUBindGroup); + } testPass.dispatchWorkgroups(numWorkgroups); testPass.end(); @@ -443,8 +567,8 @@ export class MemoryModelTester { * If the weak index's value is not 0, it means the test has observed a behavior disallowed by the memory model and * is considered a test failure. */ - protected checkResult(weakIndex: number): (i: number, v: number) => boolean { - return function (i: number, v: number): boolean { + protected checkResult(weakIndex: number): (i: number, v: number | bigint) => boolean { + return function (i: number, v: number | bigint): boolean { if (i === weakIndex && v > 0) { return false; } @@ -453,7 +577,7 @@ export class MemoryModelTester { } /** Returns a printer function that visualizes the results of checking the test results. */ - protected resultPrinter(weakIndex: number): (i: number) => string | number { + protected resultPrinter(weakIndex: number): (i: number) => string | number | bigint { return function (i: number): string | number { if (i === weakIndex) { return 0; @@ -464,16 +588,42 @@ export class MemoryModelTester { } /** Utility method that simplifies copying source buffers to device buffers. */ - protected copyBufferToBuffer(encoder: GPUCommandEncoder, buffer: BufferWithSource): void { - encoder.copyBufferToBuffer(buffer.srcBuf, 0, buffer.deviceBuf, 0, buffer.size); + protected copyBufferToBuffer( + encoder: GPUCommandEncoder, + buffer: BufferWithSource | SubBufferWithSource + ): void { + encoder.copyBufferToBuffer( + buffer.srcBuf, + 0, + buffer.deviceBuf, + (buffer as SubBufferWithSource).offset || 0, + buffer.size + ); } - /** Returns a random integer between 0 and the max. */ + /** Utility method that simplifies copying source buffers to device textures. */ + protected copyBufferToTexture(encoder: GPUCommandEncoder, texture: TextureWithSource): void { + const bytesPerWord = 4; // always uses r32uint format. + const numTexels = texture.size / bytesPerWord; + const size: GPUExtent3D = { width: kWidth, height: numTexels / kWidth }; + encoder.copyBufferToTexture( + { + buffer: texture.srcBuf, + offset: 0, + bytesPerRow: kWidth * bytesPerWord, + rowsPerImage: size.height, + }, + { texture: texture.deviceTex }, + size + ); + } + + /** Returns a random integer in the range [0, max). */ protected getRandomInt(max: number): number { - return Math.floor(Math.random() * max); + return this.prng.randomU32() % max; } - /** Returns a random number in between the min and max values. */ + /** Returns a random number in the range [min, max). */ protected getRandomInRange(min: number, max: number): number { if (min === max) { return min; @@ -626,7 +776,19 @@ const shaderMemStructures = ` }; struct IndexMemory { - value: array<u32> + value: array<u32>, + }; + + struct AtomicMemoryBarrier { + value: array<atomic<u32>, kNumBarriers> + }; + + struct IndexMemoryScratchpad { + value: array<u32, kMaxWorkgroups>, + }; + + struct IndexMemoryScratchLocations { + value: array<u32, kScratchMemorySize>, }; struct ReadResult { @@ -635,7 +797,14 @@ const shaderMemStructures = ` }; struct ReadResults { - value: array<ReadResult> + value: array<ReadResult>, + }; + + // These arrays are combine into 1 buffer because compat mode only supports 4 storage buffers by default. + struct CombinedData { + barrier: AtomicMemoryBarrier, + scratchpad: IndexMemoryScratchpad, + scratch_locations: IndexMemoryScratchLocations, }; struct StressParamsMemory { @@ -687,10 +856,8 @@ const twoBehaviorTestResultStructure = ` const commonTestShaderBindings = ` @group(0) @binding(1) var<storage, read_write> results : ReadResults; @group(0) @binding(2) var<storage, read> shuffled_workgroups : IndexMemory; - @group(0) @binding(3) var<storage, read_write> barrier : AtomicMemory; - @group(0) @binding(4) var<storage, read_write> scratchpad : IndexMemory; - @group(0) @binding(5) var<storage, read_write> scratch_locations : IndexMemory; - @group(0) @binding(6) var<uniform> stress_params : StressParamsMemory; + @group(0) @binding(3) var<storage, read_write> combo : CombinedData; + @group(0) @binding(4) var<uniform> stress_params : StressParamsMemory; `; /** The combined bindings for a test on atomic memory. */ @@ -709,6 +876,11 @@ const nonAtomicTestShaderBindings = [ commonTestShaderBindings, ].join('\n'); +/** The extra binding for texture non-atomic texture tests. */ +const textureBindings = ` +@group(1) @binding(0) var texture_locations : texture_storage_2d<r32uint, read_write>; +`; + /** Bindings used in the result aggregation phase of the test. */ const resultShaderBindings = ` @group(0) @binding(0) var<storage, read_write> test_locations : Memory; @@ -750,6 +922,16 @@ const memoryLocationFunctions = ` } `; +/** + * Function to convert an index into an equivalent 2D coordinate for the texture. + */ +const textureFunctions = ` + const kWidth = ${kWidth}; + fn indexToCoord(idx : u32) -> vec2u { + return vec2u(idx % kWidth, idx / kWidth); + } +`; + /** Functions that help add stress to the test. */ const testShaderFunctions = ` //Force the invocations in the workgroup to wait for each other, but without the general memory ordering @@ -758,12 +940,12 @@ const testShaderFunctions = ` // the barrier but does not overly reduce testing throughput. fn spin(limit: u32) { var i : u32 = 0u; - var bar_val : u32 = atomicAdd(&barrier.value[0], 1u); + var bar_val : u32 = atomicAdd(&combo.barrier.value[0], 1u); loop { if (i == 1024u || bar_val >= limit) { break; } - bar_val = atomicAdd(&barrier.value[0], 0u); + bar_val = atomicAdd(&combo.barrier.value[0], 0u); i = i + 1u; } } @@ -773,44 +955,44 @@ const testShaderFunctions = ` // the compiler optimizing out unused loads, where 100,000 is larger than the maximum number of stress iterations used // in any test. fn do_stress(iterations: u32, pattern: u32, workgroup_id: u32) { - let addr = scratch_locations.value[workgroup_id]; + let addr = combo.scratch_locations.value[workgroup_id]; switch(pattern) { case 0u: { for(var i: u32 = 0u; i < iterations; i = i + 1u) { - scratchpad.value[addr] = i; - scratchpad.value[addr] = i + 1u; + combo.scratchpad.value[addr] = i; + combo.scratchpad.value[addr] = i + 1u; } } case 1u: { for(var i: u32 = 0u; i < iterations; i = i + 1u) { - scratchpad.value[addr] = i; - let tmp1: u32 = scratchpad.value[addr]; + combo.scratchpad.value[addr] = i; + let tmp1: u32 = combo.scratchpad.value[addr]; if (tmp1 > 100000u) { - scratchpad.value[addr] = i; + combo.scratchpad.value[addr] = i; break; } } } case 2u: { for(var i: u32 = 0u; i < iterations; i = i + 1u) { - let tmp1: u32 = scratchpad.value[addr]; + let tmp1: u32 = combo.scratchpad.value[addr]; if (tmp1 > 100000u) { - scratchpad.value[addr] = i; + combo.scratchpad.value[addr] = i; break; } - scratchpad.value[addr] = i; + combo.scratchpad.value[addr] = i; } } case 3u: { for(var i: u32 = 0u; i < iterations; i = i + 1u) { - let tmp1: u32 = scratchpad.value[addr]; + let tmp1: u32 = combo.scratchpad.value[addr]; if (tmp1 > 100000u) { - scratchpad.value[addr] = i; + combo.scratchpad.value[addr] = i; break; } - let tmp2: u32 = scratchpad.value[addr]; + let tmp2: u32 = combo.scratchpad.value[addr]; if (tmp2 > 100000u) { - scratchpad.value[addr] = i; + combo.scratchpad.value[addr] = i; break; } } @@ -827,7 +1009,7 @@ const testShaderFunctions = ` */ const shaderEntryPoint = ` // Change to pipeline overridable constant when possible. - const workgroupXSize = 256u; + const workgroupXSize = kWorkgroupXSize; @compute @workgroup_size(workgroupXSize) fn main( @builtin(local_invocation_id) local_invocation_id : vec3<u32>, @builtin(workgroup_id) workgroup_id : vec3<u32>) { @@ -980,6 +1162,18 @@ const storageMemoryNonAtomicTestShaderCode = [ testShaderCommonHeader, ].join('\n'); +/** The common shader code for the test shaders that perform non-atomic texture memory litmus tests. */ +const textureMemoryNonAtomicTestShaderCode = [ + shaderMemStructures, + nonAtomicTestShaderBindings, + textureBindings, + memoryLocationFunctions, + textureFunctions, + testShaderFunctions, + shaderEntryPoint, + testShaderCommonHeader, +].join('\n'); + /** The common shader code for test shaders that perform atomic workgroup class memory litmus tests. */ const workgroupMemoryAtomicTestShaderCode = [ shaderMemStructures, @@ -1023,6 +1217,8 @@ export enum MemoryType { AtomicWorkgroupClass = 'atomic_workgroup', /** Non-atomic memory in the workgroup address space. */ NonAtomicWorkgroupClass = 'non_atomic_workgroup', + /** Non-atomic memory in a texture. */ + NonAtomicTextureClass = 'non_atomic_texture', } /** @@ -1052,21 +1248,26 @@ export function buildTestShader( testType: TestType ): string { let memoryTypeCode; - let isStorageAS = false; + let isGlobalSpace = false; switch (memoryType) { case MemoryType.AtomicStorageClass: memoryTypeCode = storageMemoryAtomicTestShaderCode; - isStorageAS = true; + isGlobalSpace = true; break; case MemoryType.NonAtomicStorageClass: memoryTypeCode = storageMemoryNonAtomicTestShaderCode; - isStorageAS = true; + isGlobalSpace = true; break; case MemoryType.AtomicWorkgroupClass: memoryTypeCode = workgroupMemoryAtomicTestShaderCode; break; case MemoryType.NonAtomicWorkgroupClass: memoryTypeCode = workgroupMemoryNonAtomicTestShaderCode; + break; + case MemoryType.NonAtomicTextureClass: + memoryTypeCode = textureMemoryNonAtomicTestShaderCode; + isGlobalSpace = true; + break; } let testTypeCode; switch (testType) { @@ -1074,7 +1275,7 @@ export function buildTestShader( testTypeCode = interWorkgroupTestShaderCode; break; case TestType.IntraWorkgroup: - if (isStorageAS) { + if (isGlobalSpace) { testTypeCode = storageIntraWorkgroupTestShaderCode; } else { testTypeCode = intraWorkgroupTestShaderCode; |