-Title: WebGPU Explainer -Shortname: webgpu-explainer -Level: 1 -Status: LS -Group: webgpu -URL: https://gpuweb.github.io/gpuweb/explainer/ -Issue Tracking: gpuweb/gpuweb#1321 https://github.com/gpuweb/gpuweb/issues/1321 -No Editor: true -No Abstract: true -Markup Shorthands: markdown yes -Markup Shorthands: dfn yes -Markup Shorthands: idl yes -Markup Shorthands: css no -Assume Explicit For: yes -Boilerplate: repository-issue-tracking no -- -Issue(tabatkins/bikeshed#2006): Set up cross-linking into the WebGPU and WGSL specs. - - -# Motivation # {#motivation} - -See [Introduction](https://gpuweb.github.io/gpuweb/#introduction). - - -# Security/Privacy # {#security} - -See [Malicious use considerations](https://gpuweb.github.io/gpuweb/#malicious-use). - - -# Additional Background # {#background} - - -## Sandboxed GPU Processes in Web Browsers ## {#gpu-process} - -A major design constraint for WebGPU is that it must be implementable and efficient in browsers that use a GPU-process architecture. -GPU drivers need access to additional kernel syscalls than what's otherwise used for Web content, and many GPU drivers are prone to hangs or crashes. -To improve stability and sandboxing, browsers use a special process that contains the GPU driver and talks with the rest of the browser through asynchronous IPC. -GPU processes are (or will be) used in Chromium, Gecko, and WebKit. - -GPU processes are less sandboxed than content processes, and they are typically shared between multiple origins. -Therefore, they must validate all messages, for example to prevent a compromised content process from being able to look at the GPU memory used by another content process. -Most of WebGPU's validation rules are necessary to ensure it is secure to use, so all the validation needs to happen in the GPU process. - -Likewise, all GPU driver objects only live in the GPU process, including large allocations (like buffers and textures) and complex objects (like pipelines). -In the content process, WebGPU types (`GPUBuffer`, `GPUTexture`, `GPURenderPipeline`, ...) are mostly just "handles" that identify objects that live in the GPU process. -This means that the CPU and GPU memory used by WebGPU object isn't necessarily known in the content process. -A `GPUBuffer` object can use maybe 150 bytes of CPU memory in the content process but hold a 1GB allocation of GPU memory. - -See also the description of [the content and device timelines in the specification](https://gpuweb.github.io/gpuweb/#programming-model-timelines). - - -# JavaScript API # {#api} - - -## Adapter Selection and Device Init ## {#initialization} - -Issue: Some changes are expected here. - -A WebGPU "adapter" (`GPUAdapter`) is an object which provides a connection to a particular WebGPU -implementation on the system (e.g. a hardware accelerated implementation on an integrated or -discrete GPU, or software implementation). -To get a `GPUAdapter` is to select which implementation to use (if multiple are available). - -Each adapter may have different optional capabilities called "features" and "limits". -These are the maximum possible capabilities that can be requested when a device is created. - -To get an adapter, an application calls `navigator.gpu.requestAdapter()`, optionally passing -options which may influence what adapter is chosen. - -`requestAdapter()` always resolves, but may resolve to null if an adapter can't be returned with -the specified options. - -
-const adapter = await navigator.gpu.requestAdapter(options); -if (!adapter) return goToFallback(); -- -A WebGPU "device" (`GPUDevice`) represents a connection to a WebGPU implementation, as well as -an arena for all WebGPU objects created from it (textures, command buffers, etc.) -All WebGPU usage is done through a WebGPU "device" (`GPUDevice`) or objects created from it. -In this sense, it serves a subset of the purpose of `WebGLRenderingContext`; however, unlike -`WebGLRenderingContext`, it is not associated with a canvas object, and most commands are -issued through "child" objects. - -To get a device, an application calls `adapter.requestDevice()`, optionally passing a descriptor -which enables additional optional capabilities (features and limits). -When any work is issued to the device, it is strictly validated against the capabilities passed -to `requestDevice()` - not the capabilities of the adapter. - -`requestDevice()` will reject (only) if the request exceeds the capabilities of the adapter. -It may *not* resolve to `null`; instead, to simplify the number of different cases an app must -handle, it may resolve to a `GPUDevice` which has already been lost - see [[#device-loss]]. - -
-const device = await adapter.requestDevice(descriptor); -device.lost.then(recoverFromDeviceLoss); -- -An adapter may become unavailable, e.g. if it is unplugged from the system, disabled to save -power, or marked "stale" (`[[current]]` becomes false). -Such an adapter can no longer vend valid devices, and always returns already-lost `GPUDevice`s. - - -## Object Validity and Destroyed-ness ## {#invalid-and-destroyed} - -### WebGPU's Error Monad ### {#error-monad} - -A.k.a. Contagious Internal Nullability. -A.k.a. transparent [promise pipelining](http://erights.org/elib/distrib/pipeline.html). - -WebGPU is a very chatty API, with some applications making tens of thousands of calls per frame to render complex scenes. -We have seen that the GPU processes needs to validate the commands to satisfy their security property. -To avoid the overhead of validating commands twice in both the GPU and content process, WebGPU is designed so Javascript calls can be forwarded directly to the GPU process and validated there. -See the error section for more details on what's validated where and how errors are reported. - -At the same time, during a single frame WebGPU objects can be created that depend on one another. -For example a `GPUCommandBuffer` can be recorded with commands that use temporary `GPUBuffer`s created in the same frame. -In this example, because of the performance constraint of WebGPU, it is not possible to send the message to create the `GPUBuffer` to the GPU process and synchronously wait for its processing before continuing Javascript execution. - -Instead, in WebGPU all objects (like `GPUBuffer`) are created immediately on the content timeline and returned to JavaScript. -The validation is almost all done asynchronously on the "device timeline". -In the good case, when no errors occur (validation or out-of-memory), everything looks to JS as if it is synchronous. -However, when an error occurs in a call, it becomes a no-op (aside from error reporting). -If the call returns an object (like `createBuffer`), the object is tagged as "invalid" on the GPU process side. - -All WebGPU calls validate that all their arguments are valid objects. -As a result, if a call takes one WebGPU object and returns a new one, the new object is also invalid (hence the term "contagious"). - - - -
- const srcBuffer = device.createBuffer({ - size: 4, - usage: GPUBufferUsage.COPY_SRC - }); - - const dstBuffer = ...; - - const encoder = device.createCommandEncoder(); - encoder.copyBufferToBuffer(srcBuffer, 0, dstBuffer, 0, 4); - - const commands = encoder.finish(); - device.queue.submit([commands]); --
- // The size of the buffer is too big, this causes an OOM and srcBuffer is invalid. - const srcBuffer = device.createBuffer({ - size: BIG_NUMBER, - usage: GPUBufferUsage.COPY_SRC - }); - - const dstBuffer = ...; - - // The encoder starts as a valid object. - const encoder = device.createCommandEncoder(); - // Special case: an invalid object is used when encoding commands so the encoder - // becomes invalid. - encoder.copyBufferToBuffer(srcBuffer, 0, dstBuffer, 0, 4); - - // commands, the this argument to GPUCommandEncoder.finish is invalid - // so the call returns an invalid object. - const commands = encoder.finish(); - // The command references an invalid object so it becomes a noop. - device.queue.submit([commands]); --
- const dstBuffer = device.createBuffer({ - size: 4 - usage: GPUBufferUsage.COPY_DST - }); - - // The buffer is not destroyed (and valid), success! - device.queue.writeBuffer(dstBuffer, 0, myData); - - buffer.destroy(); - - // The buffer is now destroyed, commands using that would use its - // content produce validation errors. - device.queue.writeBuffer(dstBuffer, 0, myData); --
- | Regular `ArrayBuffer` - | Shared Memory - | Mappable GPU buffer - | Non-mappable GPU buffer (or texture) - |
---|---|---|---|---|
CPU, in the content process - | **Visible** - | **Visible** - | Not visible - | Not visible - |
CPU, in the GPU process - | Not visible - | **Visible** - | **Visible** - | Not visible - |
GPU - | Not visible - | Not visible - | **Visible** - | **Visible** - |
- // Mapping a buffer for writing. Here offset and size are defaulted t - // so the whole buffer is mapped. - const myMapWriteBuffer = ...; - await myMapWriteBuffer.mapAsync(GPUMapMode.WRITE); - - // Mapping a buffer for reading. Only the first four bytes are mapped. - const myMapReadBuffer = ...; - await myMapReadBuffer.mapAsync(GPUMapMode.READ, 0, 4); --
- const myMapReadBuffer = ...; - await myMapReadBuffer.mapAsync(GPUMapMode.READ, 0, 4); - // Do something with the mapped buffer. - buffer.unmap(); --
- const myMapReadBuffer = device.createBuffer({ - usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST, - size: 1000, - }); - const myMapWriteBuffer = device.createBuffer({ - usage: GPUBufferUsage.MAP_WRITE | GPUBufferUsage.COPY_SRC, - size: 1000, - }); --
- const myMapReadBuffer = ...; - await myMapReadBuffer.mapAsync(GPUMapMode.READ); - const data = myMapReadBuffer.getMappedRange(); - // Do something with the data - myMapReadBuffer.unmap(); --
- const buffer = device.createBuffer({ - usage: GPUBufferUsage.UNIFORM, - size: 256, - mappedAtCreation: true, - }); - const data = buffer.getMappedRange(); - // write to data - buffer.unmap(); --
- const dracoDecoder = ...; - - const buffer = device.createBuffer({ - usage: GPUBuffer.VERTEX | GPUBuffer.INDEX, - size: dracoDecoder.decompressedSize, - mappedAtCreation: true, - }); - - dracoDecoder.decodeIn(buffer.getMappedRange()); - buffer.unmap(); --
- const texture = getTheRenderedTexture(); - - const readbackBuffer = device.createBuffer({ - usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ, - size: 4 * textureWidth * textureHeight, - }); - - // Copy data from the texture to the buffer. - const encoder = device.createCommandEncoder(); - encoder.copyTextureToBuffer( - { texture }, - { buffer, rowPitch: textureWidth * 4 }, - [textureWidth, textureHeight], - ); - device.submit([encoder.finish()]); - - // Get the data on the CPU. - await buffer.mapAsync(GPUMapMode.READ); - saveScreenshot(buffer.getMappedRange()); - buffer.unmap(); --
- void frame() { - // Create a new buffer for our updates. In practice we would - // reuse buffers from frame to frame by re-mapping them. - const stagingBuffer = device.createBuffer({ - usage: GPUBufferUsage.MAP_WRITE | GPUBufferUsage.COPY_SRC, - size: 16 * objectCount, - mappedAtCreation: true, - }); - const stagingData = new Float32Array(stagingBuffer.getMappedRange()); - - // For each draw we are going to: - // - Put the data for the draw in stagingData. - // - Record a copy from the stagingData to the uniform buffer for the draw - // - Encoder the draw - const copyEncoder = device.createCommandEncoder(); - const drawEncoder = device.createCommandEncoder(); - const renderPass = myCreateRenderPass(drawEncoder); - for (var i = 0; i < objectCount; i++) { - stagingData[i * 4 + 0] = ...; - stagingData[i * 4 + 1] = ...; - stagingData[i * 4 + 2] = ...; - stagingData[i * 4 + 3] = ...; - - const {uniformBuffer, uniformOffset} = getUniformsForDraw(i); - copyEncoder.copyBufferToBuffer( - stagingData, i * 16, - uniformBuffer, uniformOffset, - 16); - - encodeDraw(renderPass, {uniformBuffer, uniformOffset}); - } - renderPass.endPass(); - - // We are finished filling the staging buffer, unmap() it so - // we can submit commands that use it. - stagingBuffer.unmap(); - - // Submit all the copies and then all the draws. The copies - // will happen before the draw such that each draw will use - // the data that was filled inside the for-loop above. - device.queue.submit([ - copyEncoder.finish(), - drawEncoder.finish() - ]); - } --
Limit name | Type | [=limit/Better=] | [=limit/Default=] - |
---|---|---|---|
maxTextureDimension1D - | {{GPUSize32}} | Higher | 8192 - |
- The maximum allowed value for the {{GPUTextureDescriptor/size}}.[=Extent3D/width=] - of a [=texture=] created with {{GPUTextureDescriptor/dimension}} {{GPUTextureDimension/"1d"}}. - - | |||
maxTextureDimension2D - | {{GPUSize32}} | Higher | 8192 - |
- The maximum allowed value for the {{GPUTextureDescriptor/size}}.[=Extent3D/width=] and {{GPUTextureDescriptor/size}}.[=Extent3D/height=] - of a [=texture=] created with {{GPUTextureDescriptor/dimension}} {{GPUTextureDimension/"2d"}}. - - | |||
maxTextureDimension3D - | {{GPUSize32}} | Higher | 2048 - |
- The maximum allowed value for the {{GPUTextureDescriptor/size}}.[=Extent3D/width=], {{GPUTextureDescriptor/size}}.[=Extent3D/height=] and {{GPUTextureDescriptor/size}}.[=Extent3D/depthOrArrayLayers=] - of a [=texture=] created with {{GPUTextureDescriptor/dimension}} {{GPUTextureDimension/"3d"}}. - - | |||
maxTextureArrayLayers - | {{GPUSize32}} | Higher | 2048 - |
- The maximum allowed value for the {{GPUTextureDescriptor/size}}.[=Extent3D/depthOrArrayLayers=] - of a [=texture=] created with {{GPUTextureDescriptor/dimension}} {{GPUTextureDimension/"1d"}} or {{GPUTextureDimension/"2d"}}. - - | |||
maxBindGroups - | {{GPUSize32}} | Higher | 4 - |
- The maximum number of {{GPUBindGroupLayout|GPUBindGroupLayouts}} - allowed in {{GPUPipelineLayoutDescriptor/bindGroupLayouts}} - when creating a {{GPUPipelineLayout}}. - - | |||
maxDynamicUniformBuffersPerPipelineLayout - | {{GPUSize32}} | Higher | 8 - |
- The maximum number of {{GPUBindGroupLayoutDescriptor/entries}} for which: - - - [$layout entry binding type$] is {{GPUBufferBindingType/"uniform"}}, and - - {{GPUBindGroupLayoutEntry/buffer}}.{{GPUBufferBindingLayout/hasDynamicOffset}} is `true`, - - across all {{GPUPipelineLayoutDescriptor/bindGroupLayouts}} - when creating a {{GPUPipelineLayout}}. - - | |||
maxDynamicStorageBuffersPerPipelineLayout - | {{GPUSize32}} | Higher | 4 - |
- The maximum number of {{GPUBindGroupLayoutDescriptor/entries}} for which: - - - [$layout entry binding type$] is {{GPUBufferBindingType/"storage"}}, and - - {{GPUBindGroupLayoutEntry/buffer}}.{{GPUBufferBindingLayout/hasDynamicOffset}} is `true`, - - across all {{GPUPipelineLayoutDescriptor/bindGroupLayouts}} - when creating a {{GPUPipelineLayout}}. - - | |||
maxSampledTexturesPerShaderStage - | {{GPUSize32}} | Higher | 16 - |
- For each possible {{GPUShaderStage}} `stage`, - the maximum number of {{GPUBindGroupLayoutDescriptor/entries}} for which: - - - {{GPUBindGroupLayoutEntry/texture}} is not `undefined`, and - - {{GPUBindGroupLayoutEntry/visibility}} includes `stage`, - - across all {{GPUPipelineLayoutDescriptor/bindGroupLayouts}} - when creating a {{GPUPipelineLayout}}. - - | |||
maxSamplersPerShaderStage - | {{GPUSize32}} | Higher | 16 - |
- For each possible {{GPUShaderStage}} `stage`, - the maximum number of {{GPUBindGroupLayoutDescriptor/entries}} for which: - - - [=Binding member=] is {{GPUBindGroupLayoutEntry/sampler}}, and - - {{GPUBindGroupLayoutEntry/visibility}} includes `stage`, - - across all {{GPUPipelineLayoutDescriptor/bindGroupLayouts}} - when creating a {{GPUPipelineLayout}}. - - | |||
maxStorageBuffersPerShaderStage - | {{GPUSize32}} | Higher | 4 - |
- For each possible {{GPUShaderStage}} `stage`, - the maximum number of {{GPUBindGroupLayoutDescriptor/entries}} for which: - - - [$layout entry binding type$] is {{GPUBufferBindingType/"storage"}}, and - - {{GPUBindGroupLayoutEntry/visibility}} includes `stage`, - - across all {{GPUPipelineLayoutDescriptor/bindGroupLayouts}} - when creating a {{GPUPipelineLayout}}. - - | |||
maxStorageTexturesPerShaderStage - | {{GPUSize32}} | Higher | 4 - |
- For each possible {{GPUShaderStage}} `stage`, - the maximum number of {{GPUBindGroupLayoutDescriptor/entries}} for which: - - - [=Binding member=] is {{GPUBindGroupLayoutEntry/storageTexture}}, and - - {{GPUBindGroupLayoutEntry/visibility}} includes `stage`, - - across all {{GPUPipelineLayoutDescriptor/bindGroupLayouts}} - when creating a {{GPUPipelineLayout}}. - - | |||
maxUniformBuffersPerShaderStage - | {{GPUSize32}} | Higher | 12 - |
- For each possible {{GPUShaderStage}} `stage`, - the maximum number of {{GPUBindGroupLayoutDescriptor/entries}} for which: - - - [$layout entry binding type$] is {{GPUBufferBindingType/"uniform"}}, and - - {{GPUBindGroupLayoutEntry/visibility}} includes `stage`, - - across all {{GPUPipelineLayoutDescriptor/bindGroupLayouts}} - when creating a {{GPUPipelineLayout}}. - - | |||
maxUniformBufferBindingSize - | {{GPUSize32}} | Higher | 16384 - |
- The maximum {{GPUBufferBinding}}.{{GPUBufferBinding/size}} for bindings for which the - [$layout entry binding type$] is {{GPUBufferBindingType/"uniform"}}. - - | |||
maxStorageBufferBindingSize - | {{GPUSize32}} | Higher | 134217728 (128 MiB) - |
- The maximum {{GPUBufferBinding}}.{{GPUBufferBinding/size}} for bindings for which the - [$layout entry binding type$] is {{GPUBufferBindingType/"storage"}} or {{GPUBufferBindingType/"read-only-storage"}}. - - | |||
maxVertexBuffers - | {{GPUSize32}} | Higher | 8 - |
- The maximum number of {{GPUVertexState/buffers}} - when creating a {{GPURenderPipeline}}. - - | |||
maxVertexAttributes - | {{GPUSize32}} | Higher | 16 - |
- The maximum number of {{GPUVertexBufferLayout/attributes}} - in total across {{GPUVertexState/buffers}} - when creating a {{GPURenderPipeline}}. - - | |||
maxVertexBufferArrayStride - | {{GPUSize32}} | Higher | 2048 - |
- The maximum allowed {{GPUVertexBufferLayout/arrayStride}} - when creating a {{GPURenderPipeline}}. - |
- |options|: Criteria used to select the adapter. -- - **Returns:** {{Promise}}<{{GPUAdapter}}?> - - 1. Let |promise| be [=a new promise=]. - 1. Issue the following steps on the [=Device timeline=] of |this|: -
- const adapter = await navigator.gpu.requestAdapter(/* ... */); - const features = adapter.features; - // ... --
- |descriptor|: Description of the {{GPUDevice}} to request. -- - **Returns:** {{Promise}}<{{GPUDevice}}?> - - 1. Let |promise| be [=a new promise=]. - 1. Let |adapter| be |this|.{{GPUAdapter/[[adapter]]}}. - 1. Issue the following steps to the [=Device timeline=]: -
- |descriptor|: Description of the {{GPUBuffer}} to create. -- - **Returns:** {{GPUBuffer}} - - 1. If any of the following conditions are unsatisfied, return an error buffer and stop. -
- |mode|: Whether the buffer should be mapped for reading or writing. - |offset|: Offset in bytes into the buffer to the start of the range to map. - |size|: Size in bytes of the range to map. -- - **Returns:** {{Promise}}<{{undefined}}> - - Issue(gpuweb/gpuweb#605): Handle error buffers once we have a description of the error monad. - - 1. If |size| is unspecified: - 1. Let |rangeSize| be max(0, |this|.{{GPUBuffer/[[size]]}} - |offset|). - - Otherwise, let |rangeSize| be |size|. - - 1. If any of the following conditions are unsatisfied: -
[|offset|, |offset| + |rangeSize|]
.
- 1. Set |this|.{{GPUBuffer/[[mapped_ranges]]}} to `[]`.
-
- 1. Resolve |p|.
- - |offset|: Offset in bytes into the buffer to return buffer contents from. - |size|: Size in bytes of the {{ArrayBuffer}} to return. -- - **Returns:** {{ArrayBuffer}} - - 1. If |size| is unspecified: - 1. Let |rangeSize| be max(0, |this|.{{GPUBuffer/[[size]]}} - |offset|). - - Otherwise, let |rangeSize| be |size|. - - 1. If any of the following conditions are unsatisfied, throw an {{OperationError}} and stop. -
- descriptor: Description of the {{GPUTexture}} to create. -- - **Returns:** {{GPUTexture}} - - Issue: Describe {{GPUDevice/createTexture()}} algorithm steps. -
- |descriptor|: Description of the {{GPUTextureView}} to create. -- - **Returns:** |view|, of type {{GPUTextureView}}. - - 1. Set |descriptor| to the result of [$resolving GPUTextureViewDescriptor defaults$] with |descriptor|. - 1. Issue the following steps on the [=Device timeline=] of |this|: -
- |descriptor|: Description of the {{GPUSampler}} to create. -- - **Returns:** {{GPUSampler}} - - 1. Let |s| be a new {{GPUSampler}} object. - 1. Set |s|.{{GPUSampler/[[descriptor]]}} to |descriptor|. - 1. Set |s|.{{GPUSampler/[[isComparison]]}} to `false` if the {{GPUSamplerDescriptor/compare}} attribute - of |s|.{{GPUSampler/[[descriptor]]}} is `null` or undefined. Otherwise, set it to `true`. - 1. Set |s|.{{GPUSampler/[[isFiltering]]}} to `false` if none of {{GPUSamplerDescriptor/minFilter}}, - {{GPUSamplerDescriptor/magFilter}}, or {{GPUSamplerDescriptor/mipmapFilter}} has the value of - {{GPUFilterMode/"linear"}}. Otherwise, set it to `true`. - 1. Return |s|. - -
Binding member - | Resource type - | Binding type - | Binding usage - |
---|---|---|---|
{{GPUBindGroupLayoutEntry/buffer}} - | {{GPUBufferBinding}} - | {{GPUBufferBindingType/"uniform"}} - | [=internal usage/constant=] - |
{{GPUBufferBindingType/"storage"}} - | [=internal usage/storage=] - | ||
{{GPUBufferBindingType/"read-only-storage"}} - | [=internal usage/storage-read=] - - | ||
{{GPUBindGroupLayoutEntry/sampler}} - | {{GPUSampler}} - | {{GPUSamplerBindingType/"filtering"}} - | [=internal usage/constant=] - |
{{GPUSamplerBindingType/"non-filtering"}} - | [=internal usage/constant=] - | ||
{{GPUSamplerBindingType/"comparison"}} - | [=internal usage/constant=] - - | ||
{{GPUBindGroupLayoutEntry/texture}} - | {{GPUTextureView}} - | {{GPUTextureSampleType/"float"}} - | [=internal usage/constant=] - |
{{GPUTextureSampleType/"unfilterable-float"}} - | [=internal usage/constant=] - | ||
{{GPUTextureSampleType/"depth"}} - | [=internal usage/constant=] - | ||
{{GPUTextureSampleType/"sint"}} - | [=internal usage/constant=] - | ||
{{GPUTextureSampleType/"uint"}} - | [=internal usage/constant=] - - | ||
{{GPUBindGroupLayoutEntry/storageTexture}} - | {{GPUTextureView}} - | {{GPUStorageTextureAccess/"read-only"}} - | [=internal usage/storage-read=] - |
{{GPUStorageTextureAccess/"write-only"}} - | [=internal usage/storage-write=] - |
- |descriptor|: Description of the {{GPUBindGroupLayout}} to create. -- - **Returns:** {{GPUBindGroupLayout}} - - 1. Let |layout| be a new valid {{GPUBindGroupLayout}} object. - 1. Let |limits| be |this|.{{GPUDevice/[[device]]}}.{{device/[[limits]]}}. - 1. Issue the following steps on the [=Device timeline=] of |this|: -
- |descriptor|: Description of the {{GPUBindGroup}} to create. -- - **Returns:** {{GPUBindGroup}} - - 1. Let |bindGroup| be a new valid {{GPUBindGroup}} object. - 1. Let |limits| be |this|.{{GPUDevice/[[device]]}}.{{device/[[limits]]}}.{{supported limits/maxUniformBufferBindingSize}}. - 1. Issue the following steps on the [=Device timeline=] of |this|: -
- |descriptor|: Description of the {{GPUPipelineLayout}} to create. -- - **Returns:** {{GPUPipelineLayout}} - - 1. If any of the following conditions are unsatisfied: -
- descriptor: Description of the {{GPUShaderModule}} to create. -- - **Returns:** {{GPUShaderModule}} - - Issue: Describe {{GPUDevice/createShaderModule()}} algorithm steps. -
- |index|: Index into the pipeline layout's {{GPUPipelineLayout/[[bindGroupLayouts]]}} - sequence. -- - **Returns:** {{GPUBindGroupLayout}} - - 1. If |index| ≥ - |this|.{{GPUObjectBase/[[device]]}}.{{device/[[limits]]}}.{{supported limits/maxBindGroups}}: - 1. Throw a {{RangeError}}. - - 1. If |this| is not [=valid=]: - 1. Return a new error {{GPUBindGroupLayout}}. - - 1. Return a new {{GPUBindGroupLayout}} object that references the same internal object as - |this|.{{GPUPipelineBase/[[layout]]}}.{{GPUPipelineLayout/[[bindGroupLayouts]]}}[|index|]. - - Issue: Specify this more properly once we have internal objects for {{GPUBindGroupLayout}}. - Alternatively only spec is as a new internal objects that's [=group-equivalent=] - - Note: Only returning new {{GPUBindGroupLayout}} objects ensures no synchronization is necessary - between the [=Content timeline=] and the [=Device timeline=]. -
- |descriptor|: Description of the {{GPUComputePipeline}} to create. -- - **Returns:** {{GPUComputePipeline}} - - If any of the following conditions are unsatisfied: -
- |descriptor|: Description of the {{GPUComputePipeline}} to create. -- - **Returns:** {{Promise}}<{{GPUComputePipeline}}> - - 1. Let |promise| be [=a new promise=]. - 1. Issue the following steps on the [=Device timeline=] of |this|: -
- |descriptor|: Description of the {{GPURenderPipeline}} to create. -- - **Returns:** {{GPURenderPipeline}} - - 1. Let |pipeline| be a new valid {{GPURenderPipeline}} object. - 1. Issue the following steps on the [=Device timeline=] of |this|: -
- |descriptor|: Description of the {{GPURenderPipeline}} to create. -- - **Returns:** {{Promise}}<{{GPURenderPipeline}}> - - 1. Let |promise| be [=a new promise=]. - 1. Issue the following steps on the [=Device timeline=] of |this|: -
Index format | -Primitive restart value | -
---|---|
{{GPUIndexFormat/"uint16"}} | -0xFFFF | -
{{GPUIndexFormat/"uint32"}} | -0xFFFFFFFF | -
- descriptor: Description of the {{GPUCommandEncoder}} to create. -- - **Returns:** {{GPUCommandEncoder}} - - Issue: Describe {{GPUDevice/createCommandEncoder()}} algorithm steps. -
- |descriptor|: Description of the {{GPURenderPassEncoder}} to create. -- - **Returns:** {{GPURenderPassEncoder}} - - Issue the following steps on the [=Device timeline=] of |this|: -
- descriptor: -- - **Returns:** {{GPUComputePassEncoder}} - - Issue the following steps on the [=Device timeline=] of |this|: -
- |source|: The {{GPUBuffer}} to copy from. - |sourceOffset|: Offset in bytes into |source| to begin copying from. - |destination|: The {{GPUBuffer}} to copy to. - |destinationOffset|: Offset in bytes into |destination| to place the copied data. - |size|: Bytes to copy. -- - **Returns:** {{undefined}} - - If any of the following conditions are unsatisfied, generate a validation error and stop. -
- |source|: Combined with |copySize|, defines the region of the source buffer. - |destination|: Combined with |copySize|, defines the region of the destination [=texture subresource=]. - |copySize|: -- - **Returns:** {{undefined}} - - If any of the following conditions are unsatisfied, generate a validation error and stop. -
- |source|: Combined with |copySize|, defines the region of the source [=texture subresources=]. - |destination|: Combined with |copySize|, defines the region of the destination buffer. - |copySize|: -- - **Returns:** {{undefined}} - - If any of the following conditions are unsatisfied, generate a validation error and stop. -
- |source|: Combined with |copySize|, defines the region of the source [=texture subresources=]. - |destination|: Combined with |copySize|, defines the region of the destination [=texture subresources=]. - |copySize|: -- - **Returns:** {{undefined}} - - 1. If any of the following conditions are unsatisfied, generate a validation error and stop. -
- |groupLabel|: The label for the command group. -- - **Returns:** {{undefined}} - - Issue the following steps on the [=Device timeline=] of |this|: -
- markerLabel: The label to insert. -- - **Returns:** {{undefined}} - - Issue the following steps on the [=Device timeline=] of |this|: -
- |querySet|: The query set that will store the timestamp values. - |queryIndex|: The index of the query in the query set. -- - **Returns:** {{undefined}} - - 1. If |this|.{{GPUDevice/[[device]]}}.{{device/[[features]]}} does not [=list/contain=] - {{GPUFeatureName/"timestamp-query"}}, throw a {{TypeError}}. - 1. If any of the following conditions are unsatisfied, generate a validation error and stop. -
- querySet: - firstQuery: - queryCount: - destination: - destinationOffset: -- - **Returns:** {{undefined}} - - If any of the following conditions are unsatisfied, generate a {{GPUValidationError}} and stop. -
- descriptor: -- - **Returns:** {{GPUCommandBuffer}} - - 1. Let |commandBuffer| be a new {{GPUCommandBuffer}}. - 1. Issue the following steps on the [=Device timeline=] of |this|: -
- |index|: The index to set the bind group at. - |bindGroup|: Bind group to use for subsequent render or compute commands. - - - -- - Issue: Resolve bikeshed conflict when using `argumentdef` with overloaded functions that prevents us from - defining |dynamicOffsets|. - - **Returns:** {{undefined}} - - Issue the following steps on the [=Device timeline=] of |this|.{{GPUObjectBase/[[device]]}}: -
- |index|: The index to set the bind group at. - |bindGroup|: Bind group to use for subsequent render or compute commands. - |dynamicOffsetsData|: Array containing buffer offsets in bytes for each entry in - |bindGroup| marked as {{GPUBindGroupLayoutEntry/buffer}}.{{GPUBufferBindingLayout/hasDynamicOffset}}. - |dynamicOffsetsDataStart|: Offset in elements into |dynamicOffsetsData| where the - buffer offset data begins. - |dynamicOffsetsDataLength|: Number of buffer offsets to read from |dynamicOffsetsData|. -- - **Returns:** {{undefined}} - - Issue the following steps on the [=Device timeline=] of |this|.{{GPUObjectBase/[[device]]}}: -
- |groupLabel|: The label for the command group. -- - **Returns:** {{undefined}} - - Issue the following steps on the [=Device timeline=] of |this|: -
- markerLabel: The label to insert. -- - **Returns:** {{undefined}} -
- |pipeline|: The compute pipeline to use for subsequent dispatch commands. -- - **Returns:** {{undefined}} - - Issue the following steps on the [=Device timeline=] of |this|.{{GPUObjectBase/[[device]]}}: -
- |x|: X dimension of the grid of workgroups to dispatch. - |y|: Y dimension of the grid of workgroups to dispatch. - |z|: Z dimension of the grid of workgroups to dispatch. -- - **Returns:** {{undefined}} - - Issue the following steps on the [=Device timeline=] of |this|.{{GPUObjectBase/[[device]]}}: -
- let dispatchIndirectParameters = new Uint32Array(3); - dispatchIndirectParameters[0] = x; - dispatchIndirectParameters[1] = y; - dispatchIndirectParameters[2] = z; -- -
- |indirectBuffer|: Buffer containing the [=indirect dispatch parameters=]. - |indirectOffset|: Offset in bytes into |indirectBuffer| where the dispatch data begins. -- - **Returns:** {{undefined}} - - Issue the following steps on the [=Device timeline=] of |this|.{{GPUObjectBase/[[device]]}}: -
- querySet: - queryIndex: -- - **Returns:** {{undefined}} - - 1. If |this|.{{GPUDevice/[[device]]}}.{{device/[[features]]}} does not [=list/contain=] - {{GPUFeatureName/"pipeline-statistics-query"}}, throw a {{TypeError}}. - - Issue: Describe {{GPUComputePassEncoder/beginPipelineStatisticsQuery()}} algorithm steps. -
- |querySet|: The query set that will store the timestamp values. - |queryIndex|: The index of the query in the query set. -- - **Returns:** {{undefined}} - - 1. If |this|.{{GPUDevice/[[device]]}}.{{device/[[features]]}} does not [=list/contain=] - {{GPUFeatureName/"timestamp-query"}}, throw a {{TypeError}}. - 1. If any of the following conditions are unsatisfied, generate a validation error and stop. -
- |pipeline|: The render pipeline to use for subsequent drawing commands. -- - **Returns:** {{undefined}} - - Issue the following steps on the [=Device timeline=] of |this|.{{GPUObjectBase/[[device]]}}: -
- |buffer|: Buffer containing index data to use for subsequent drawing commands. - |indexFormat|: Format of the index data contained in |buffer|. - |offset|: Offset in bytes into |buffer| where the index data begins. - |size|: Size in bytes of the index data in |buffer|. - If `0`, |buffer|.{{GPUBuffer/[[size]]}} - |offset| is used. -- - **Returns:** {{undefined}} - - Issue the following steps on the [=Device timeline=] of |this|.{{GPUObjectBase/[[device]]}}: -
- |slot|: The vertex buffer slot to set the vertex buffer for. - |buffer|: Buffer containing vertex data to use for subsequent drawing commands. - |offset|: Offset in bytes into |buffer| where the vertex data begins. - |size|: Size in bytes of the vertex data in |buffer|. - If `0`, |buffer|.{{GPUBuffer/[[size]]}} - |offset| is used. -- - **Returns:** {{undefined}} - - Issue the following steps on the [=Device timeline=] of |this|.{{GPUObjectBase/[[device]]}}: -
- vertexCount: The number of vertices to draw. - instanceCount: The number of instances to draw. - firstVertex: Offset into the vertex buffers, in vertices, to begin drawing from. - firstInstance: First instance to draw. -- - **Returns:** {{undefined}} - - Issue the following steps on the [=Device timeline=] of |this|.{{GPUObjectBase/[[device]]}}: -
- indexCount: The number of indices to draw. - instanceCount: The number of instances to draw. - firstIndex: Offset into the index buffer, in indices, begin drawing from. - baseVertex: Added to each index value before indexing into the vertex buffers. - firstInstance: First instance to draw. -- - **Returns:** {{undefined}} - - Issue the following steps on the [=Device timeline=] of |this|.{{GPUObjectBase/[[device]]}}: -
- let drawIndirectParameters = new Uint32Array(4); - drawIndirectParameters[0] = vertexCount; - drawIndirectParameters[1] = instanceCount; - drawIndirectParameters[2] = firstVertex; - drawIndirectParameters[3] = firstInstance; -- -
- |indirectBuffer|: Buffer containing the [=indirect draw parameters=]. - |indirectOffset|: Offset in bytes into |indirectBuffer| where the drawing data begins. -- - **Returns:** {{undefined}} - - Issue the following steps on the [=Device timeline=] of |this|.{{GPUObjectBase/[[device]]}}: -
- let drawIndexedIndirectParameters = new Uint32Array(5); - drawIndexedIndirectParameters[0] = indexCount; - drawIndexedIndirectParameters[1] = instanceCount; - drawIndexedIndirectParameters[2] = firstIndex; - drawIndexedIndirectParameters[3] = baseVertex; - drawIndexedIndirectParameters[4] = firstInstance; -- -
- |indirectBuffer|: Buffer containing the [=indirect drawIndexed parameters=]. - |indirectOffset|: Offset in bytes into |indirectBuffer| where the drawing data begins. -- - **Returns:** {{undefined}} - - Issue the following steps on the [=Device timeline=] of |this|.{{GPUObjectBase/[[device]]}}: -
- |x|: Minimum X value of the viewport in pixels. - |y|: Minimum Y value of the viewport in pixels. - |width|: Width of the viewport in pixels. - |height|: Height of the viewport in pixels. - |minDepth|: Minimum depth value of the viewport. - |maxDepth|: Maximum depth value of the viewport. -- - **Returns:** {{undefined}} - - Issue the following steps on the [=Device timeline=] of |this|: -
- |x|: Minimum X value of the scissor rectangle in pixels. - |y|: Minimum Y value of the scissor rectangle in pixels. - |width|: Width of the scissor rectangle in pixels. - |height|: Height of the scissor rectangle in pixels. -- - **Returns:** {{undefined}} - - Issue the following steps on the [=Device timeline=] of |this|: -
- color: The color to use when blending. --
- reference: The stencil reference value. --
- |queryIndex|: The index of the query in the query set. -- - **Returns:** {{undefined}} - - Issue the following steps on the [=Device timeline=] of |this|.{{GPUObjectBase/[[device]]}}: -
- querySet: - queryIndex: -- - **Returns:** {{undefined}} - - 1. If |this|.{{GPUDevice/[[device]]}}.{{device/[[features]]}} does not [=list/contain=] - {{GPUFeatureName/"pipeline-statistics-query"}}, throw a {{TypeError}}. - - Issue: Describe {{GPURenderPassEncoder/beginPipelineStatisticsQuery()}} algorithm steps. -
- querySet: The query set that will store the timestamp values. - queryIndex: The index of the query in the query set. -- - **Returns:** {{undefined}} - - 1. If |this|.{{GPUDevice/[[device]]}}.{{device/[[features]]}} does not [=list/contain=] - {{GPUFeatureName/"timestamp-query"}}, throw a {{TypeError}}. - 1. If any of the following conditions are unsatisfied, generate a validation error and stop. -
- bundles: List of render bundles to execute. -- - **Returns:** {{undefined}} - - Issue: Describe {{GPURenderPassEncoder/executeBundles()}} algorithm steps. -
- descriptor: Description of the {{GPURenderBundleEncoder}} to create. -- - **Returns:** {{GPURenderBundleEncoder}} - - Issue: Describe {{GPUDevice/createRenderBundleEncoder()}} algorithm steps. -
- descriptor: -- - **Returns:** {{GPURenderBundle}} - - Issue: Describe {{GPURenderBundleEncoder/finish()}} algorithm steps. -
- |buffer|: The buffer to write to. - |bufferOffset|: Offset in bytes into |buffer| to begin writing at. - |data|: Data to write into |buffer|. - |dataOffset|: Offset in into |data| to begin writing from. Given in elements if - |data| is a `TypedArray` and bytes otherwise. - |size|: Size of content to write from |data| to |buffer|. Given in elements if - |data| is a `TypedArray` and bytes otherwise. -- - **Returns:** {{undefined}} - - 1. If |data| is an {{ArrayBuffer}} or {{DataView}}, let the element type be "byte". - Otherwise, |data| is a TypedArray; let the element type be the type of the TypedArray. - 1. Let |dataSize| be the size of |data|, in elements. - 1. If |size| is unspecified, - let |contentsSize| be |dataSize| − |dataOffset|. - Otherwise, let |contentsSize| be |size|. - 1. If any of the following conditions are unsatisfied, - throw {{OperationError}} and stop. - -
- |destination|: The [=texture subresource=] and origin to write to. - |data|: Data to write into |destination|. - |dataLayout|: Layout of the content in |data|. - |size|: Extents of the content to write from |data| to |destination|. -- - **Returns:** {{undefined}} - - 1. Let |dataBytes| be [=get a copy of the buffer source|a copy of the bytes held by the buffer source=] |data|. - 1. Let |dataByteSize| be the number of bytes in |dataBytes|. - 1. If any of the following conditions are unsatisfied, - throw {{OperationError}} and stop. -
- |source|: {{ImageBitmap}} and origin to copy to |destination|. - |destination|: The [=texture subresource=] and origin to write to. - |copySize|: Extents of the content to write from |source| to |destination|. -- - **Returns:** {{undefined}} - - If any of the following conditions are unsatisfied, throw an {{OperationError}} and stop. -
- |commandBuffers|: -- - **Returns:** {{undefined}} - - Issue the following steps on the [=Device timeline=] of |this|: -
-- - **Returns:** {{Promise}}<{{undefined}}> - - Issue: Describe {{GPUQueue/onSubmittedWorkDone()}} algorithm steps. -
- descriptor: Description of the {{GPUQuerySet}} to create. -- - **Returns:** {{GPUQuerySet}} - - 1. If |descriptor|.{{GPUQuerySetDescriptor/type}} is {{GPUQueryType/"pipeline-statistics"}}, - but |this|.{{GPUDevice/[[device]]}}.{{device/[[features]]}} does not [=list/contain=] - {{GPUFeatureName/"pipeline-statistics-query"}}, throw a {{TypeError}}. - 1. If |descriptor|.{{GPUQuerySetDescriptor/type}} is {{GPUQueryType/"timestamp"}}, - but |this|.{{GPUDevice/[[device]]}}.{{device/[[features]]}} does not [=list/contain=] - {{GPUFeatureName/"timestamp-query"}}, throw a {{TypeError}}. - 1. If any of the following requirements are unmet, return an error query set and stop. -
- const canvas = document.createElement('canvas'); - const context = canvas.getContext('gpupresent'); - const swapChain = context.configureSwapChain(/* ... */); - // ... --
- |descriptor|: Description of the {{GPUSwapChain}} to configure. -- - **Returns:** {{GPUSwapChain}} - - 1. Issue the following steps on the [=Device timeline=] of |this|: -
- |adapter|: Adapter the swap chain format should be queried for. -- - **Returns:** {{GPUTextureFormat}} - -
GPUCanvasCompositingAlphaMode - | Description - | dst.rgb - | dst.a - |
---|---|---|---|
{{GPUCanvasCompositingAlphaMode/opaque}} - | Paint RGB as opaque and ignore alpha values. - If the content is not already opaque, implementations may need to clear alpha to opaque during presentation. - | |dst.rgb = src.rgb| - | |dst.a = 1| - |
{{GPUCanvasCompositingAlphaMode/premultiplied}} - | Composite assuming color values are premultiplied by their alpha value. - 100% red 50% opaque is [0.5, 0, 0, 0.5]. - Color values must be less than or equal to their alpha value. - [1.0, 0, 0, 0.5] is "super-luminant" and cannot reliably be displayed. - | |dst.rgb = src.rgb + dst.rgb*(1-src.a)| - | |dst.a = src.a + dst.a*(1-src.a)| - |
Format - | {{GPUTextureSampleType}} - | {{GPUTextureUsage/RENDER_ATTACHMENT}} - | {{GPUTextureUsage/STORAGE}} - | |
---|---|---|---|---|
8-bit per component | - | |||
{{GPUTextureFormat/r8unorm}} - | {{GPUTextureSampleType/"float"}}, {{GPUTextureSampleType/"unfilterable-float"}} - | ✓ - | - | |
{{GPUTextureFormat/r8snorm}} - | {{GPUTextureSampleType/"float"}}, {{GPUTextureSampleType/"unfilterable-float"}} - | - | - | |
{{GPUTextureFormat/r8uint}} - | {{GPUTextureSampleType/"uint"}} - | ✓ - | - | |
{{GPUTextureFormat/r8sint}} - | {{GPUTextureSampleType/"sint"}} - | ✓ - | - | |
{{GPUTextureFormat/rg8unorm}} - | {{GPUTextureSampleType/"float"}}, {{GPUTextureSampleType/"unfilterable-float"}} - | ✓ - | - | |
{{GPUTextureFormat/rg8snorm}} - | {{GPUTextureSampleType/"float"}}, {{GPUTextureSampleType/"unfilterable-float"}} - | - | - | |
{{GPUTextureFormat/rg8uint}} - | {{GPUTextureSampleType/"uint"}} - | ✓ - | - | |
{{GPUTextureFormat/rg8sint}} - | {{GPUTextureSampleType/"sint"}} - | ✓ - | - | |
{{GPUTextureFormat/rgba8unorm}} - | {{GPUTextureSampleType/"float"}}, {{GPUTextureSampleType/"unfilterable-float"}} - | ✓ - | ✓ - | |
{{GPUTextureFormat/rgba8unorm-srgb}} - | {{GPUTextureSampleType/"float"}}, {{GPUTextureSampleType/"unfilterable-float"}} - | ✓ - | - | |
{{GPUTextureFormat/rgba8snorm}} - | {{GPUTextureSampleType/"float"}}, {{GPUTextureSampleType/"unfilterable-float"}} - | - | ✓ - | |
{{GPUTextureFormat/rgba8uint}} - | {{GPUTextureSampleType/"uint"}} - | ✓ - | ✓ - | |
{{GPUTextureFormat/rgba8sint}} - | {{GPUTextureSampleType/"sint"}} - | ✓ - | ✓ - | |
{{GPUTextureFormat/bgra8unorm}} - | {{GPUTextureSampleType/"float"}}, {{GPUTextureSampleType/"unfilterable-float"}} - | ✓ - | - | |
{{GPUTextureFormat/bgra8unorm-srgb}} - | {{GPUTextureSampleType/"float"}}, {{GPUTextureSampleType/"unfilterable-float"}} - | ✓ - | - | |
16-bit per component | - | |||
{{GPUTextureFormat/r16uint}} - | {{GPUTextureSampleType/"uint"}} - | ✓ - | - | |
{{GPUTextureFormat/r16sint}} - | {{GPUTextureSampleType/"sint"}} - | ✓ - | - | |
{{GPUTextureFormat/r16float}} - | {{GPUTextureSampleType/"float"}}, {{GPUTextureSampleType/"unfilterable-float"}} - | ✓ - | - | |
{{GPUTextureFormat/rg16uint}} - | {{GPUTextureSampleType/"uint"}} - | ✓ - | - | |
{{GPUTextureFormat/rg16sint}} - | {{GPUTextureSampleType/"sint"}} - | ✓ - | - | |
{{GPUTextureFormat/rg16float}} - | {{GPUTextureSampleType/"float"}}, {{GPUTextureSampleType/"unfilterable-float"}} - | ✓ - | - | |
{{GPUTextureFormat/rgba16uint}} - | {{GPUTextureSampleType/"uint"}} - | ✓ - | ✓ - | |
{{GPUTextureFormat/rgba16sint}} - | {{GPUTextureSampleType/"sint"}} - | ✓ - | ✓ - | |
{{GPUTextureFormat/rgba16float}} - | {{GPUTextureSampleType/"float"}}, {{GPUTextureSampleType/"unfilterable-float"}} - | ✓ - | ✓ - | |
32-bit per component | - | |||
{{GPUTextureFormat/r32uint}} - | {{GPUTextureSampleType/"uint"}} - | ✓ - | ✓ - | |
{{GPUTextureFormat/r32sint}} - | {{GPUTextureSampleType/"sint"}} - | ✓ - | ✓ - | |
{{GPUTextureFormat/r32float}} - | {{GPUTextureSampleType/"unfilterable-float"}} - | ✓ - | ✓ - | |
{{GPUTextureFormat/rg32uint}} - | {{GPUTextureSampleType/"uint"}} - | ✓ - | ✓ - | |
{{GPUTextureFormat/rg32sint}} - | {{GPUTextureSampleType/"sint"}} - | ✓ - | ✓ - | |
{{GPUTextureFormat/rg32float}} - | {{GPUTextureSampleType/"unfilterable-float"}} - | ✓ - | ✓ - | |
{{GPUTextureFormat/rgba32uint}} - | {{GPUTextureSampleType/"uint"}} - | ✓ - | ✓ - | |
{{GPUTextureFormat/rgba32sint}} - | {{GPUTextureSampleType/"sint"}} - | ✓ - | ✓ - | |
{{GPUTextureFormat/rgba32float}} - | {{GPUTextureSampleType/"unfilterable-float"}} - | ✓ - | ✓ - | |
mixed component width | - | |||
{{GPUTextureFormat/rgb10a2unorm}} - | {{GPUTextureSampleType/"float"}}, {{GPUTextureSampleType/"unfilterable-float"}} - | ✓ - | - | |
{{GPUTextureFormat/rg11b10ufloat}} - | {{GPUTextureSampleType/"float"}}, {{GPUTextureSampleType/"unfilterable-float"}} - | - | - - |
Format - | Bytes per texel - | Aspect - | {{GPUTextureSampleType}} - | Copy aspect from Buffer - | Copy aspect into Buffer - |
---|---|---|---|---|---|
{{GPUTextureFormat/stencil8}} - | 1 − 5 - | stencil - | {{GPUTextureSampleType/"uint"}} - | ✓ - | |
{{GPUTextureFormat/depth16unorm}} - | 2 - | depth - | {{GPUTextureSampleType/"depth"}} - | ✓ - | |
{{GPUTextureFormat/depth24plus}} - | 4 - | depth - | {{GPUTextureSampleType/"depth"}} - | ✗ - | |
{{GPUTextureFormat/depth24plus-stencil8}} - | 4 − 8 - | depth - | {{GPUTextureSampleType/"depth"}} - | ✗ - | |
stencil - | {{GPUTextureSampleType/"uint"}} - | ✓ - | |||
{{GPUTextureFormat/depth32float}} - | 4 - | depth - | {{GPUTextureSampleType/"depth"}} - | ✗ - | ✓ - |
Format - | Bytes per block - | {{GPUTextureSampleType}} - | Block Size - | [=Feature=] - |
---|---|---|---|---|
{{GPUTextureFormat/rgb9e5ufloat}} - | 4 - | {{GPUTextureSampleType/"float"}}, {{GPUTextureSampleType/"unfilterable-float"}} - | 1 × 1 - | - |
{{GPUTextureFormat/bc1-rgba-unorm}} - | 8 - | {{GPUTextureSampleType/"float"}}, {{GPUTextureSampleType/"unfilterable-float"}} - | 4 × 4 - | {{GPUFeatureName/texture-compression-bc}} - |
{{GPUTextureFormat/bc1-rgba-unorm-srgb}} - | ||||
{{GPUTextureFormat/bc2-rgba-unorm}} - | 16 - | |||
{{GPUTextureFormat/bc2-rgba-unorm-srgb}} - | ||||
{{GPUTextureFormat/bc3-rgba-unorm}} - | 16 - | |||
{{GPUTextureFormat/bc3-rgba-unorm-srgb}} - | ||||
{{GPUTextureFormat/bc4-r-unorm}} - | 8 - | |||
{{GPUTextureFormat/bc4-r-snorm}} - | ||||
{{GPUTextureFormat/bc5-rg-unorm}} - | 16 - | |||
{{GPUTextureFormat/bc5-rg-snorm}} - | ||||
{{GPUTextureFormat/bc6h-rgb-ufloat}} - | 16 - | |||
{{GPUTextureFormat/bc6h-rgb-float}} - | ||||
{{GPUTextureFormat/bc7-rgba-unorm}} - | 16 - | |||
{{GPUTextureFormat/bc7-rgba-unorm-srgb}} - |
-Title: WebGPU Shading Language -Shortname: WGSL -Level: 1 -Status: w3c/ED -Group: webgpu -URL: https://gpuweb.github.io/gpuweb/wgsl/ -Ignored Vars: i, e, e1, e2, e3, N, M, v, Stride, Offset, Align, Extent, S, T, T1 - -!Participate: File an issue (open issues) - -Editor: David Neto, Google https://www.google.com, dneto@google.com -Editor: Myles C. Maxfield, Apple Inc., mmaxfield@apple.com, w3cid 77180 -Former Editor: dan sinclair, Google https://www.google.com, dsinclair@google.com -Abstract: Shading language for WebGPU. -Markup Shorthands: markdown yes -Markup Shorthands: biblio yes -Markup Shorthands: idl no -- - - -
-{ - "WebGPU": { - "authors": [ - "Dzmitry Malyshau", - "Justin Fan", - "Kai Ninomiya" - ], - "href": "https://gpuweb.github.io/gpuweb/", - "title": "WebGPU", - "status": "Editor's Draft", - "publisher": "W3C", - "deliveredBy": [ - "https://github.com/gpuweb/gpuweb" - ] - }, - "VulkanMemoryModel": { - "authors": [ - "Jeff Bolz", - "Alan Baker", - "Tobias Hector", - "David Neto", - "Robert Simpson", - "Brian Sumner" - ], - "href": "https://www.khronos.org/registry/vulkan/specs/1.2-extensions/html/vkspec.html#memory-model", - "title": "Vulkan Memory Model", - "publisher": "Khronos Group" - } -} -- -# Introduction # {#intro} - -WebGPU Shader Language ([SHORTNAME]) is the shader language for [[!WebGPU]]. -That is, an application using the WebGPU API uses [SHORTNAME] to express the programs, known as shaders, -that run on the GPU. - -
Token | Definition - |
---|---|
`DECIMAL_FLOAT_LITERAL` | `(-?[0-9]*.[0-9]+ | -?[0-9]+.[0-9]*)((e|E)(+|-)?[0-9]+)?` - |
`HEX_FLOAT_LITERAL` | `-?0x([0-9a-fA-F]*.?[0-9a-fA-F]+ | [0-9a-fA-F]+.[0-9a-fA-F]*)(p|P)(+|-)?[0-9]+` - |
`INT_LITERAL` | `-?0x[0-9a-fA-F]+ | 0 | -?[1-9][0-9]*` - |
`UINT_LITERAL` | `0x[0-9a-fA-F]+u | 0u | [1-9][0-9]*u` - |
-const_literal - : INT_LITERAL - | UINT_LITERAL - | FLOAT_LITERAL - | TRUE - | FALSE -- -
-FLOAT_LITERAL - : DECIMAL_FLOAT_LITERAL - | HEX_FLOAT_LITERAL -- - -## Keywords TODO ## {#keywords} - -TODO: *Stub* - -See [[#keyword-summary]] for a list of keywords. - -## Identifiers TODO ## {#identifiers} - -
Token | Definition - |
---|---|
`IDENT` | `[a-zA-Z][0-9a-zA-Z_]*` - |
-attribute_list - : ATTR_LEFT (attribute COMMA)* attribute ATTR_RIGHT - -attribute - : IDENT PAREN_LEFT literal_or_ident PAREN_RIGHT - | IDENT - -literal_or_ident - : FLOAT_LITERAL - | INT_LITERAL - | UINT_LITERAL - | IDENT -- -
Attribute | Valid Values | Description - | ||||||||||||||||||||||||||||||
---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|
`access` - | `read`, `write`, or `read_write` - | Must only be applied to a type used as a store type for a variable in - the [=storage classes/storage=] storage class or a variable of [storage - texture](#texture-storage) type. - - Specifies the access qualification of a storage [=resource=] variable. - - | ||||||||||||||||||||||||||||||
`align` - | positive i32 literal - | Must only be applied to a member of a [=structure=] type. - - Must be a power of 2. - - See memory layout [alignment and size](#alignment-and-size). - - | ||||||||||||||||||||||||||||||
`binding`
- non-negative i32 literal
- | Must only be applied to a [=resource=] variable.
-
- Specifies the binding number of the resource in a bind [=attribute/group=].
- See [[#resource-interface]].
-
- | `block`
- | *None*
- | Must only be applied to a [=structure=] type.
-
- Indicates this structure type represents the contents of a buffer
- resource occupying a single binding slot in the [=resource interface of a
- shader|shader's resource interface=].
-
- The `block` attribute must be applied to a structure type used as the
- [=store type=] of a [=uniform buffer=] or [=storage buffer=] variable.
-
- A structure type with the block attribute must not be:
- * the element type of an [=array=] type
- * the member type in another structure
-
- | `builtin`
- | a builtin variable identifier
- | Must only be applied to an entry point function parameter, entry point
- return type, or member of a [=structure=].
-
- Declares a builtin variable.
- See [[#builtin-variables]].
-
- | `constant_id`
- | non-negative i32 literal
- | Must only be applied to module scope constant declaration of [=scalar=] type.
-
- Specifies a [=pipeline-overridable=] constant.
-
- | `group`
- | non-negative i32 literal
- | Must only be applied to a [=resource=] variable.
-
- Specifies the binding group of the resource.
- See [[#resource-interface]].
-
- | `interpolate`
- | One or two parameters.
-
- The first parameter must be an [=interpolation type=].
- The second parameter, if present, must specify the [=interpolation sampling=].
- | Must only be applied to an entry point function parameter, entry point
- return type, or member of a [=structure=] type.
- Must only be applied to declarations of scalars or vectors of floating-point type.
- Must not be used with the [=compute=] shader stage.
- If the first parameter is `flat`, the second parameter must not be specified.
-
- Specifies how the user-defined IO must be interpolated.
- The attribute is only significant on user-defined [=vertex=] outputs
- and [=fragment=] inputs.
- See [[#interpolation]].
-
- | `location`
- | non-negative i32 literal
- | Must only be applied to an entry point function parameter, entry point
- return type, or member of a [=structure=] type.
- Must only be applied to declarations of [=numeric scalar=] or [=numeric
- vector=] type.
- Must not be used with the [=compute=] shader stage.
-
- Specifies a part of the user-defined IO of an entry point.
- See [[#input-output-locations]].
-
- | `size`
- | positive i32 literal
- | Must only be applied to a member of a [=structure=] type.
-
- The number of bytes reserved in the struct for this member.
-
- | `stage`
- | `compute`, `vertex`, or `fragment`
- | Must only be applied to a function declaration.
-
- Declares an entry point by specifying its pipeline stage.
-
- | `stride`
- | positive i32 literal
- | Must only be applied to an [=array=] type.
-
- The number of bytes from the start of one element of the array to the
- start of the next element.
-
- | `workgroup_size`
- | One, two or three parameters.
-
- Each parameter is either a positive i32 literal or the name of a
- [=pipeline-overridable=] constant of i32 type.
- | Must only be applied to a [=compute shader stage=] function declaration.
-
- Specifies the x, y, and z dimensions of the [=workgroup grid=] for the compute shader.
-
- The first parameter specifies the x dimension.
- The second parameter, if provided, specifies the y dimension, otherwise is assumed to be 1.
- The third parameter, if provided, specifies the z dimension, otherwise is assumed to be 1.
- Each dimension must be at least 1 and at most an upper bound specified by the WebGPU API.
-
- | |
*Scalar* | [=scalar=] types: one of bool, i32, u32, f32 - |
*BoolVec* | [[#vector-types]] with bool component - |
*Int* | i32 or u32 - |
*IntVec* | [[#vector-types]] with an *Int* component - |
*Integral* | *Int* or [[#vector-types]] with an *Int* component - |
*SignedIntegral* | i32 or [[#vector-types]] with an i32 component - |
*FloatVec* | [[#vector-types]] with f32 component - |
*Floating* | f32 or *FloatVec* - |
*Arity(T)* | number of components in [[#vector-types]] *T* - |
Type | Description - |
---|---|
vec*N*<*T*> | Vector of *N* elements of type *T*. - *N* must be in {2, 3, 4} and *T* - must be one of the [=scalar=] types. - We say *T* is the component type of the vector. - |
Type | Description - |
---|---|
mat|N|x|M|<f32> - | Matrix of |N| columns and |M| rows, where |N| and |M| are both in {2, 3, 4}. - Equivalently, it can be viewed as |N| column vectors of type vec|M|<f32>. - |
Type | Description - |
---|---|
array<|E|,|N|> | An |N|-element array of elements of type |E|. - |N| must be 1 or larger. - |
array<|E|> | A runtime-sized array of elements of type |E|,
- also known as a runtime array.
- These may only appear in specific contexts. - |
Type | Description - |
---|---|
struct<|T|1,...,|T|N> - | An ordered tuple of *N* members of types - |T|n through |T|N, with |N| being an integer greater than 0. - A structure type declaration specifies an identifier name for each member. - Two members of the same structure type must not have the same name. - |
-struct_decl - : attribute_list* STRUCT IDENT struct_body_decl -- -
-struct_body_decl - : BRACE_LEFT struct_member* BRACE_RIGHT - -struct_member - : attribute_list* variable_ident_decl SEMICOLON -- -[SHORTNAME] defines the following attributes that can be applied to structure types: - * [=attribute/block=] - -[SHORTNAME] defines the following attributes that can be applied to structure members: - * [=attribute/builtin=] - * [=attribute/location=] - * [=attribute/stride=] - * [=attribute/align=] - * [=attribute/size=] - -Note: Layout attributes may be required if the structure type is used -to define a [=uniform buffer=] or a [=storage buffer=]. See [[#memory-layouts]]. - -
Storage class - | Readable by shader? Writable by shader? - | Sharing among invocations - | Variable scope - | Restrictions on stored values - | Notes - |
---|---|---|---|---|---|
function - | Read-write - | Same invocation only - | [=Function scope=] - | [=Storable=] - | - |
private - | Read-write - | Same invocation only - | [=Module scope=] - | [=Storable=] - | - |
workgroup - | Read-write - | Invocations in the same [=compute shader stage|compute shader=] [=compute shader stage/workgroup=] - | [=Module scope=] - | [=Storable=] - | - |
uniform - | Read-only - | Invocations in the same [=shader stage=] - | [=Module scope=] - | [=Host-shareable=] - | For [=uniform buffer=] variables - |
storage - | Readable. - Also writable if the variable is not read-only. - | Invocations in the same [=shader stage=] - | [=Module scope=] - | [=Host-shareable=] - | For [=storage buffer=] variables - |
handle - | Read-only - | Invocations in the same shader stage - | [=Module scope=] - | Opaque representation of handle to a sampler or texture - | Used for sampler and texture variables - The token `handle` is reserved: it is never used in a [SHORTNAME] program. - |
-storage_class - : IN - | OUT - | FUNCTION - | PRIVATE - | WORKGROUP - | UNIFORM - | STORAGE -- -
WGSL storage class | SPIR-V storage class - |
---|---|
uniform | Uniform - |
workgroup | Workgroup - |
handle | UniformConstant - |
storage | StorageBuffer - |
private | Private - |
function | Function - |
Host-shareable type |T| - | [=AlignOf=](|T|) - | [=SizeOf=](|T|) - |
---|---|---|
[=i32=], [=u32=], or [=f32=] - | 4 - | 4 - |
vec2<|T|> - | 8 - | 8 - |
vec3<|T|> - | 16 - | 12 - |
vec4<|T|> - | 16 - | 16 - |
mat|N|x|M| (col-major) - (General form) - | [=AlignOf=](vec|M|) - | [=SizeOf=](array<vec|M|, |N|>) - |
mat2x2<f32> - | 8 - | 16 - |
mat3x2<f32> - | 8 - | 24 - |
mat4x2<f32> - | 8 - | 32 - |
mat2x3<f32> - | 16 - | 32 - |
mat3x3<f32> - | 16 - | 48 - |
mat4x3<f32> - | 16 - | 64 - |
mat2x4<f32> - | 16 - | 32 - |
mat3x4<f32> - | 16 - | 48 - |
mat4x4<f32> - | 16 - | 64 - |
struct |S| - | max([=AlignOf=](S, M1), ... , [=AlignOf=](S, Mn)) - | [=roundUp=]([=AlignOf=](|S|), [=OffsetOf=](|S|, |L|) + [=SizeOf=](|S|, |L|)) - Where |L| is the last member of the structure - |
array<|E|, |N|> - (Implicit stride) - | [=AlignOf=](|E|) - | |N| * [=roundUp=]([=AlignOf=](|E|), [=SizeOf=](|E|)) - |
array<|E|> - (Implicit stride) - | [=AlignOf=](|E|) - | Nruntime * [=roundUp=]([=AlignOf=](|E|), [=SizeOf=](|E|)) - Where Nruntime is the runtime-determined number of elements of |T| - |
[[[=stride=](|Q|)]] array<|E|, |N|> - | [=AlignOf=](|E|) - | |N| * |Q| - |
[[[=stride=](|Q|)]] array<|E|> - | [=AlignOf=](|E|) - | Nruntime * |Q| - |
atomic<|T|> - | [=AlignOf=](|T|) - | [=SizeOf=](|T|) - |
- [=OffsetOf=](|S|, MN) = [=roundUp=]([=AlignOf=](|S|, MN), [=OffsetOf=](|S|, MN-1) + [=SizeOf=](|S|, MN-1)
- Where MN is the current member and MN-1 is the previous member
-
- [=AlignOf=](|S|) = max([=AlignOf=](|S|, M1), ... , [=AlignOf=](|S|, MN)) -
- -The size of a structure is equal to the offset plus the size of its last member, -rounded to the next multiple of the structure's alignment: -
- [=SizeOf=](|S|) = [=roundUp=]([=AlignOf=](|S|), [=OffsetOf=](|S|, |L|) + [=SizeOf=](|S|, |L|))
- Where |L| is the last member of the structure
-
- [=StrideOf=](array<|T|[, |N|]>) = [=roundUp=]([=AlignOf=](T), [=SizeOf=](T)) -
- -In all cases, the array stride must be a multiple of the element alignment. - -
- [=SizeOf=](array<|T|, |N|>) = [=StrideOf=](array<|T|, |N|>) × |N|
- [=SizeOf=](array<|T|>) = [=StrideOf=](array<|T|>) × Nruntime
-
- [=AlignOf=](array<|T|[, N]>) = [=AlignOf=](|T|) -
- -For example, the layout for a `[[stride(S)]] arrayHost-shareable type |S| - | [=RequiredAlignOf=](|S|, [=storage classes/storage=]) - | [=RequiredAlignOf=](|S|, [=storage classes/uniform=]) - |
---|---|---|
[=i32=], [=u32=], or [=f32=] - | [=AlignOf=](|S|) - | [=AlignOf=](|S|) - |
vec|N|<`T`> - | [=AlignOf=](|S|) - | [=AlignOf=](|S|) - |
mat|N|x|M|<f32> - | [=AlignOf=](|S|) - | [=AlignOf=](|S|) - |
array<|T|,|N|> - | [=AlignOf=](|T|) - | [=roundUp=](16, [=AlignOf=](|T|)) - |
array<|T|> - | [=AlignOf=](|T|) - | [=roundUp=](16, [=AlignOf=](|T|)) - |
struct<T0, ..., TN> - | max([=AlignOf=](T0), ..., [=AlignOf=](TN)) - | [=roundUp=](16, max([=AlignOf=](T0), ..., [=AlignOf=](TN))) - |
atomic<|T|> - | [=AlignOf=](|T|) - | [=AlignOf=](|T|) - |
- [=OffsetOf=](|S|, |M|) = |k| × [=RequiredAlignOf=](|T|, C)
- Where |k| is a non-negative integer and |M| is a member of structure |S| with type |T|
-
- [=StrideOf=](array<|T|[, |N|]>) = |k| × [=RequiredAlignOf=](|T|, C)
- Where |k| is a non-negative integer
-
Constraint | Type | Description - |
---|---|---|
|SC| is a [=storage class=], |T| is a [=storable=] type - | ref<|SC|,|T|> - | The reference type
- identified with the set of [=memory views=] for memory locations in |SC| holding values of type |T|. - In this context |T| is known as the store type. - Reference types are not written [SHORTNAME] progam source; instead they are used to analyze a [SHORTNAME] program. - |
|SC| is a [=storage class=], |T| is a [=storable=] type - | ptr<|SC|,|T|> - | The pointer type
- identified with the set of [=memory views=] for memory locations in |SC| holding values of type |T|. - In this context |T| is known as the pointee type. - Pointer types appear in [SHORTNAME] progam source. - |
-Each pointer value |p| of type ptr<|SC|,|T|> corresponds to a unique reference value |r| of type ref<|SC|,|T|>, -and vice versa, -where |p| and |r| describe the same memory view. -- -In [SHORTNAME] a reference value always corresponds to the memory view -for some or all of the memory locations for some variable. -This defines the originating variable for the reference value. -A pointer value always corresponds to a reference value, and so the originating variable -of a pointer is the same as the originating variable of the corresponding reference. - -Note: The originating variable is a dynamic concept. -The originating variable for a formal parameter of a function depends on the call sites for the function. -Different call sites may supply pointers into different originating variables. - -References and pointers are distinguished by how they are used: - -* The type of a [=variable=] is a reference type. -* The [=address-of=] operation (unary `&`) converts a reference value to its corresponding pointer value. -* The [=indirection=] operation (unary `*`) converts a pointer value to its corresponding reference value. -* A const declaration can be of pointer type, but not of reference type. -* A [=formal parameter=] can be of pointer type, but not of reference type. -* An [=assignment statement=] updates the contents of memory via a reference: - * The left-hand side of the assignment statement must be of reference type. - * The right-hand side of the assignment statement must evaluate to the store type of the left-hand side. -* The Load Rule: Inside a function, a reference is automatically dereferenced (read from) to satisfy type rules: - * In a function, when a reference expression |r| with store type |T| is used in a statement or an expression, where - * The only potentially matching type rules require |r| to have a value of type |T|, then - * That type rule requirement is considered to have been met, and - * The result of evaluating |r| in that context is the value (of type |T|) stored in the memory locations - referenced by |r| at the time of evaluation. - -Defining references in this way enables simple idiomatic use of variables: - -
Channel format - | Number of stored bits - | Interpetation of stored bits - | Shader type | Shader value -(Channel Transfer Function) - |
---|---|---|---|---|
8unorm | 8 | unsigned integer |v| ∈ {0,...,255} | f32 | |v| ÷ 255 - |
8snorm | 8 | signed integer |v| ∈ {-128,...,127} | f32 | max(-1, |v| ÷ 127) - |
8uint | 8 | unsigned integer |v| ∈ {0,...,255} | u32 | |v| ÷ 255 - |
8sint | 8 | signed integer |v| ∈ {-128,...,127} | i32 | max(-1, |v| ÷ 127) - |
16uint | 16 | unsigned integer |v| ∈ {0,...,65535} | u32 | |v| - |
16sint | 16 | signed integer |v| ∈ {-32768,...,32767} | i32 | |v| - |
16float | 16 | IEEE 754 16-bit floating point value |v|, with 1 sign bit, 5 exponent bits, 10 mantissa bits | f32 | |v| - |
32uint | 32 | 32-bit unsigned integer value |v| | u32 | |v| - |
32sint | 32 | 32-bit signed integer value |v| | i32 | |v| - |
32float | 32 | IEEE 754 32-bit floating point value |v| | f32 | |v| - |
Texel format - | Channel format - | Channels in memory order - | Corresponding shader value - |
---|---|---|---|
rgba8unorm | 8unorm | r, g, b, a | vec4<f32>(CTF(r), CTF(g), CTF(b), CTF(a)) - |
rgba8snorm | 8snorm | r, g, b, a | vec4<f32>(CTF(r), CTF(g), CTF(b), CTF(a)) - |
rgba8uint | 8uint | r, g, b, a | vec4<u32>(CTF(r), CTF(g), CTF(b), CTF(a)) - |
rgba8sint | 8sint | r, g, b, a | vec4<i32>(CTF(r), CTF(g), CTF(b), CTF(a)) - |
rgba16uint | 16uint | r, g, b, a | vec4<u32>(CTF(r), CTF(g), CTF(b), CTF(a)) - |
rgba16sint | 16sint | r, g, b, a | vec4<i32>(CTF(r), CTF(g), CTF(b), CTF(a)) - |
rgba16float | 16float | r, g, b, a | vec4<f32>(CTF(r), CTF(g), CTF(b), CTF(a)) - |
r32uint | 32uint | r | vec4<u32>(CTF(r), 0u, 0u, 1u) - |
r32sint | 32sint | r | vec4<i32>(CTF(r), 0, 0, 1) - |
r32float | 32float | r | vec4<f32>(CTF(r), 0.0, 0.0, 1.0) - |
rg32uint | 32uint | r, g | vec4<u32>(CTF(r), CTF(g), 0.0, 1.0) - |
rg32sint | 32sint | r, g | vec4<i32>(CTF(r), CTF(g), 0.0, 1.0) - |
rg32float | 32float | r, g | vec4<f32>(CTF(r), CTF(g), 0.0, 1.0) - |
rgba32uint | 32uint | r, g, b, a | vec4<u32>(CTF(r), CTF(g), CTF(b), CTF(a)) - |
rgba32sint | 32sint | r, g, b, a | vec4<i32>(CTF(r), CTF(g), CTF(b), CTF(a)) - |
rgba32float | 32float | r, g, b, a | vec4<f32>(CTF(r), CTF(g), CTF(b), CTF(a)) - |
Texel format - | SPIR-V Image Format - | SPIR-V Enabling Capability - |
---|---|---|
rgba8unorm | Rgba8 | Shader - |
rgba8snorm | Rgba8Snorm | Shader - |
rgba8uint | Rgba8ui | Shader - |
rgba8sint | Rgba8i | Shader - |
rgba16uint | Rgba16ui | Shader - |
rgba16sint | Rgba16i | Shader - |
rgba16float | Rgba16f | Shader - |
r32uint | R32ui | Shader - |
r32sint | R32i | Shader - |
r32float | R32f | Shader - |
rg32uint | Rg32ui | StorageImageExtendedFormats - |
rg32sint | Rg32i | StorageImageExtendedFormats - |
rg32float | Rg32f | StorageImageExtendedFormats - |
rgba32uint | Rgba32ui | Shader - |
rgba32sint | Rgba32i | Shader - |
rgba32float | Rgba32f | Shader - |
-`texture_1d-* type must be `f32`, `i32` or `u32` -* The parameterized type for the images is the type after conversion from sampling. - E.g. you can have an image with texels with 8bit unorm components, but when you sample - them you get a 32-bit float result (or vec-of-f32). - -### Multisampled Texture Types ### {#multisampled-texture-type} - -` - %1 = OpTypeImage %type 1D 0 0 0 1 Unknown - -`texture_2d ` - %1 = OpTypeImage %type 2D 0 0 0 1 Unknown - -`texture_2d_array ` - %1 = OpTypeImage %type 2D 0 1 0 1 Unknown - -`texture_3d ` - %1 = OpTypeImage %type 3D 0 0 0 1 Unknown - -`texture_cube ` - %1 = OpTypeImage %type Cube 0 0 0 1 Unknown - -`texture_cube_array ` - %1 = OpTypeImage %type Cube 0 1 0 1 Unknown -
-`texture_multisampled_2d-* type must be `f32`, `i32` or `u32` - -### Storage Texture Types ### {#texture-storage} - -A read-only storage texture supports reading a single texel without the use of a sampler, -with automatic conversion of the stored texel value to a usable shader value. A write-only storage -texture supports writing a single texel, with automatic conversion -of the shader value to a stored texel value. -See [[#texture-builtin-functions]]. - -A storage texture type must be parameterized by one of the -[=storage-texel-format|texel formats for storage textures=]. -The texel format determines the conversion function as specified in [[#texel-formats]]. - -For a write-only storage texture the *inverse* of the conversion function is used to convert the shader value to -the stored texel. - -TODO(dneto): Move description of the conversion to the builtin function that actually does the reading. - -` - %1 = OpTypeImage %type 2D 0 0 1 1 Unknown -
-`texture_storage_1d- -In the SPIR-V mapping: -* The *Image Format* parameter of the image type declaration is - as specified by the SPIR-V texel format correspondence table in [[#texel-formats]]. -* The *Sampled Type* parameter of the image type declaration is - the SPIR-V scalar type corresponding to the channel format for the texel format. - -When mapping to SPIR-V, a read-only storage texture variable must have a `NonWritable` decoration and -a write-only storage texture variable must have a `NonReadable` decoration. - -For example: - -` - // %1 = OpTypeImage sampled_type 1D 0 0 0 2 image_format - -`texture_storage_2d ` - // %1 = OpTypeImage sampled_type 2D 0 0 0 2 image_format - -`texture_storage_2d_array ` - // %1 = OpTypeImage sampled_type 2D 0 1 0 2 image_format - -`texture_storage_3d ` - // %1 = OpTypeImage sampled_type 3D 0 0 0 2 texel_format -
-`texture_depth_2d` - %1 = OpTypeImage %f32 2D 1 0 0 1 Unknown - -`texture_depth_2d_array` - %1 = OpTypeImage %f32 2D 1 1 0 1 Unknown - -`texture_depth_cube` - %1 = OpTypeImage %f32 Cube 1 0 0 1 Unknown - -`texture_depth_cube_array` - %1 = OpTypeImage %f32 Cube 1 1 0 1 Unknown -- -### Sampler Type ### {#sampler-type} -
-sampler - OpTypeSampler - -sampler_comparison - OpTypeSampler -- -### Texture Types Grammar ### {#texture-types-grammar} -TODO: Add texture usage validation rules. - -
-texture_sampler_types - : sampler_type - | depth_texture_type - | sampled_texture_type LESS_THAN type_decl GREATER_THAN - | multisampled_texture_type LESS_THAN type_decl GREATER_THAN - | storage_texture_type LESS_THAN texel_format GREATER_THAN - -sampler_type - : SAMPLER - | SAMPLER_COMPARISON - -sampled_texture_type - : TEXTURE_1D - | TEXTURE_2D - | TEXTURE_2D_ARRAY - | TEXTURE_3D - | TEXTURE_CUBE - | TEXTURE_CUBE_ARRAY - -multisampled_texture_type - : TEXTURE_MULTISAMPLED_2D - -storage_texture_type - : TEXTURE_STORAGE_1D - | TEXTURE_STORAGE_2D - | TEXTURE_STORAGE_2D_ARRAY - | TEXTURE_STORAGE_3D - -depth_texture_type - : TEXTURE_DEPTH_2D - | TEXTURE_DEPTH_2D_ARRAY - | TEXTURE_DEPTH_CUBE - | TEXTURE_DEPTH_CUBE_ARRAY - -texel_format - : R8UNORM - R8 -- Capability: StorageImageExtendedFormats - | R8SNORM - R8Snorm -- Capability: StorageImageExtendedFormats - | R8UINT - R8ui -- Capability: StorageImageExtendedFormats - | R8SINT - R8i -- Capability: StorageImageExtendedFormats - | R16UINT - R16ui -- Capability: StorageImageExtendedFormats - | R16SINT - R16i -- Capability: StorageImageExtendedFormats - | R16FLOAT - R16f -- Capability: StorageImageExtendedFormats - | RG8UNORM - Rg8 -- Capability: StorageImageExtendedFormats - | RG8SNORM - Rg8Snorm -- Capability: StorageImageExtendedFormats - | RG8UINT - Rg8ui -- Capability: StorageImageExtendedFormats - | RG8SINT - Rg8i -- Capability: StorageImageExtendedFormats - | R32UINT - R32ui - | R32SINT - R32i - | R32FLOAT - R32f - | RG16UINT - Rg16ui -- Capability: StorageImageExtendedFormats - | RG16SINT - Rg16i -- Capability: StorageImageExtendedFormats - | RG16FLOAT - Rg16f -- Capability: StorageImageExtendedFormats - | RGBA8UNORM - Rgba8 - | RGBA8UNORM-SRGB - ??? - | RGBA8SNORM - Rgba8Snorm - | RGBA8UINT - Rgba8ui - | RGBA8SINT - Rgba8i - | BGRA8UNORM - Rgba8 ??? - | BGRA8UNORM-SRGB - ??? - | RGB10A2UNORM - Rgb10A2 -- Capability: StorageImageExtendedFormats - | RG11B10FLOAT - R11fG11fB10f -- Capability: StorageImageExtendedFormats - | RG32UINT - Rg32ui -- Capability: StorageImageExtendedFormats - | RG32SINT - Rg32i -- Capability: StorageImageExtendedFormats - | RG32FLOAT - Rg32f -- Capability: StorageImageExtendedFormats - | RGBA16UINT - Rgba16ui - | RGBA16SINT - Rgba16i - | RGBA16FLOAT - Rgba16f - | RGBA32UINT - Rgba32ui - | RGBA32SINT - Rgba32i - | RGBA32FLOAT - Rgba32f - -- -## Atomic Types ## {#atomic-types} - -Operations on atomic objects in [SHORTNAME] are mutually ordered for each object. -That is, during execution of a shader stage, for each atomic object A, all -agents observe the same order of operations applied to A. -The ordering for distinct atomic objects may not be related in any way; no -causality is implied. -Note that variables in [=storage classes/workgroup=] storage are shared within a -[=compute shader stage/workgroup=], but are not shared between different -workgroups. - -Atomic objects may only be operated on by the -[[#atomic-builtin-functions|atomic builtin functions]]. - -Atomic types may only be instantiated by variables in the [=storage -classes/workgroup=] storage class or `read_write` [=attribute/access=] variables in the -[=storage classes/storage=] storage class. - -
Type | Description - |
---|---|
atomic<|T|> - | Atomic of type |T|. |T| must be either [=u32=] or [=i32=]. - |
-- -- [[block]] struct S { - a : atomic -; - b : atomic ; - }; - - [[group(0), binding(0)]] - var x : [[access(read_write)]] S; - - // Maps to the following SPIR-V: - // - When atomic types are members of a struct, the Volatile decoration - // is annotated on the member. - // OpDecorate %S Block - // OpMemberDecorate %S 0 Volatile - // OpMemberDecorate %S 1 Volatile - // ... - // %i32 = OpTypeInt 32 1 - // %u32 = OpTypeInt 32 0 - // %S = OpTypeStruct %i32 %u32 - // %ptr_storage_S = OpTypePointer StorageBuffer %S - // %x = OpVariable %ptr_storage_S StorageBuffer -
-- -## Type Aliases TODO ## {#type-aliases} - -- var -x : atomic ; - - // Maps to the following SPIR-V: - // - When atomic types are directly instantiated by a variable, the Volatile - // decoration is annotated on the OpVariable. - // OpDecorate %x Volatile - // ... - // %u32 = OpTypeInt 32 0 - // %ptr_workgroup_u32 = OpTypePointer Workgroup %S - // %x = OpVariable %ptr_workgroup_u32 Workgroup -
-type_alias - : TYPE IDENT EQUAL type_decl -- -
-type_decl - : IDENT - | BOOL - | FLOAT32 - | INT32 - | UINT32 - | VEC2 LESS_THAN type_decl GREATER_THAN - | VEC3 LESS_THAN type_decl GREATER_THAN - | VEC4 LESS_THAN type_decl GREATER_THAN - | POINTER LESS_THAN storage_class COMMA type_decl GREATER_THAN - | attribute_list* ARRAY LESS_THAN type_decl (COMMA INT_LITERAL)? GREATER_THAN - | MAT2x2 LESS_THAN type_decl GREATER_THAN - | MAT2x3 LESS_THAN type_decl GREATER_THAN - | MAT2x4 LESS_THAN type_decl GREATER_THAN - | MAT3x2 LESS_THAN type_decl GREATER_THAN - | MAT3x3 LESS_THAN type_decl GREATER_THAN - | MAT3x4 LESS_THAN type_decl GREATER_THAN - | MAT4x2 LESS_THAN type_decl GREATER_THAN - | MAT4x3 LESS_THAN type_decl GREATER_THAN - | MAT4x4 LESS_THAN type_decl GREATER_THAN - | texture_sampler_types -- -When the type declaration is an identifer, then the expression must be in scope of a -declaration of the identifier as a type alias or structure type. - -
-variable_statement - : variable_decl - | variable_decl EQUAL short_circuit_or_expression - | LET (IDENT | variable_ident_decl) EQUAL short_circuit_or_expression - -variable_decl - : VAR variable_storage_decoration? variable_ident_decl - -variable_ident_decl - : IDENT COLON attribute_list* type_decl - -variable_storage_decoration - : LESS_THAN storage_class GREATER_THAN - -- -The `let` identifiers denote values that are immutable. -When a `let` identifier is declared without the corresponding type, -e.g. `let foo = 4`, the type is automatically inferred from the expression to the right of `=`. -If the type is provided, e.g `let foo: i32 = 4`, it has to match exactly to the type of the initializer expression. - -Variables in the [=storage classes/storage=] storage class and variables with a -[storage texture](#texture-storage) type must have an [=access=] attribute -applied to the store type. - -Two variables with overlapping lifetimes will not have overlapping storage. - -When a variable is created, its storage contains an initial value as follows: - -* For variables in the [=storage classes/private=] or [=storage classes/function=] storage classes: - * The zero value for the store type, if the variable declaration has no initializer. - * Otherwise, it is the result of evaluating the initializer expression at that point in the program execution. -* For variables in other storage classes, the execution environment provides the initial value. - -Consider the following snippet of WGSL: -
-global_variable_decl - : attribute_list* variable_decl (EQUAL const_expr)? -- -
-global_constant_decl - : attribute_list* LET variable_ident_decl global_const_initializer? - -global_const_initializer - : EQUAL const_expr - -const_expr - : type_decl PAREN_LEFT (const_expr COMMA)* const_expr PAREN_RIGHT - | const_literal -- -
Precondition | Conclusion | Notes - |
---|---|---|
`true` : bool | OpConstantTrue %bool - | |
`false` : bool | OpConstantFalse %bool - | |
*INT_LITERAL* : i32 | OpConstant %int *literal* - | |
*UINT_LITERAL* : u32 | OpConstant %uint *literal* - | |
*FLOAT_LITERAL* : f32 | OpConstant %float *literal* - |
Precondition | Conclusion | Notes - |
---|---|---|
*e* : bool | `bool(e)` : bool | Identity. In the SPIR-V translation, the ID of this expression reuses the ID of the operand. - |
*e* : i32 | `i32(e)` : i32 | Identity. In the SPIR-V translation, the ID of this expression reuses the ID of the operand. - |
*e* : u32 | `u32(e)` : u32 | Identity. In the SPIR-V translation, the ID of this expression reuses the ID of the operand. - |
*e* : f32 | `f32(e)` : f32 | Identity. In the SPIR-V translation, the ID of this expression reuses the ID of the operand. - |
Precondition | Conclusion | Notes - | |||||||||||||||||||||||||||
---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|
|e| : |T| - | `vec`|N|`<`|T|`>(`|e|`)` : vec|N|<|T|> - | Evaluates |e| once. Results in the |N|-element vector where each component has the value of |e|. - | |||||||||||||||||||||||||||
*e1* : *T* - *e2* : *T* - | `vec2OpCompositeConstruct
- | *e* : vec2<T>
- | `vec2 | Identity. The result is |e|.
- | *e1* : *T* | - *e2* : *T* - *e3* : *T* - `vec3 | OpCompositeConstruct
- | *e1* : *T* | - *e2* : vec2<*T*> - `vec3 | - `vec3 OpCompositeConstruct
- | *e* : vec3<T>
- | `vec3 | Identity. The result is |e|.
- | *e1* : *T* | - *e2* : *T* - *e3* : *T* - *e4* : *T* - `vec4 | OpCompositeConstruct
- | *e1* : *T* | - *e2* : *T* - *e3* : vec2<*T*> - `vec4 | - `vec4 - `vec4 OpCompositeConstruct
- | *e1* : vec2<*T*> | - *e2* : vec2<*T*> - `vec4 | OpCompositeConstruct
- | *e1* : *T* | - *e2* : vec3<*T*> - `vec4 | - `vec4 OpCompositeConstruct
- | *e* : vec4<T>
- | `vec4 | Identity. The result is |e|.
- | |
Precondition | Conclusion | Notes - | ||||||
---|---|---|---|---|---|---|---|---|
*e1* : vec2 - *e2* : vec2 - *e3* : vec2 - *e4* : vec2 `mat2x2 | - `mat3x2 - `mat4x2 Column by column construction. | - OpCompositeConstruct - *e1* : vec3 | - *e2* : vec3 - *e3* : vec3 - *e4* : vec3 `mat2x3 | - `mat3x3 - `mat4x3 Column by column construction. | - OpCompositeConstruct - *e1* : vec4 | - *e2* : vec4 - *e3* : vec4 - *e4* : vec4 `mat2x4 | - `mat3x4 - `mat4x4 Column by column construction. | - OpCompositeConstruct - |
Precondition | Conclusion | Notes - |
---|---|---|
*e1* : *T* - ... - *eN* : *T* - | `array<`*T*,*N*`>(e1,...,eN)` : array<*T*, *N*> - | Construction of an array from elements - |
Precondition | Conclusion | Notes - |
---|---|---|
*e1* : *T1* - ... - *eN* : *TN* - *T1* is storable - ... - *TN* is storable - S is a structure type with members having types *T1* ... *TN*. - The expression is in the scope of declaration of S. - | `S(e1,...,eN)` : S - | Construction of a structure from members - |
Precondition | Conclusion | Notes - |
---|---|---|
`bool()` : bool | false Zero value (OpConstantNull for bool) - | |
`i32()` : i32 | 0 Zero value (OpConstantNull for i32) - | |
`u32()` : u32 | 0u Zero value (OpConstantNull for u32) - | |
`f32()` : f32 | 0.0 Zero value (OpConstantNull for f32) - |
Precondition | Conclusion | Notes - | ||||||
---|---|---|---|---|---|---|---|---|
- | `vec2Zero value (OpConstantNull)
- |
- | `vec3 | Zero value (OpConstantNull)
- |
- | `vec4 | Zero value (OpConstantNull)
- | |
Precondition | Conclusion | Notes - | ||||||
---|---|---|---|---|---|---|---|---|
- | `mat2x2 - `mat3x2 - `mat4x2 Zero value (OpConstantNull)
- |
- | `mat2x3 | - `mat3x3 - `mat4x3 Zero value (OpConstantNull)
- |
- | `mat2x4 | - `mat3x4 - `mat4x4 Zero value (OpConstantNull)
- | |
Precondition | Conclusion | Notes - |
---|---|---|
*T* is storable - | `array<`*T*,*N*`>()` : array<*T*, *N*> - | Zero-valued array (OpConstantNull) - |
Precondition | Conclusion | Notes - |
---|---|---|
`S` is a storable structure type. - The expression is in the scope of declaration of S. - | `S()` : S - | Zero-valued structure: a structure of type S where each member is the zero value for its member type.
- - (OpConstantNull) - |
Precondition | Conclusion | Notes - |
---|---|---|
|e| : u32 | `bool(`|e|`)` : bool - | Coercion to boolean. - The result is false if |e| is 0, and true otherwise. - (Use OpINotEqual to compare |e| against 0.) - |
|e| : i32 | `bool(`|e|`)` : bool - | Coercion to boolean. - The result is false if |e| is 0, and true otherwise. - (Use OpINotEqual to compare |e| against 0.) - |
|e| : f32 | `bool(`|e|`)` : bool - | Coercion to boolean. - The result is false if |e| is 0.0 or -0.0, and true otherwise. - In particular NaN and infinity values map to true. - (Use OpFUnordNotEqual to compare |e| against `0.0`.) - |
|e| : u32 | `i32(`|e|`)` : i32 - | Reinterpretation of bits. - The result is the unique value in [=i32=] that is equal to (|e| mod 232). - (OpBitcast) - |
|e| : f32 | `i32(`|e|`)` : i32 | Value conversion, including invalid cases. (OpConvertFToS) - |
|e| : i32 | `u32(`|e|`)` : u32 - | Reinterpretation of bits. - The result is the unique value in [=u32=] that is equal to (|e| mod 232). - (OpBitcast) - |
|e| : f32 | `u32(`|e|`)` : u32 - | Value conversion, including invalid cases. (OpConvertFToU) - |
|e| : i32 | `f32(`|e|`)` : f32 | Value conversion, including invalid cases. (OpConvertSToF) - |
|e| : u32 | `f32(`|e|`)` : f32 | Value conversion, including invalid cases. (OpConvertUToF) - |
Precondition | Conclusion | Notes - |
---|---|---|
|e| : vec|N|<u32> - | `vec`|N|<`bool`>`(`|e|`)` : vec|N|<bool> - | Component-wise coercion of a unsigned integer vector to a boolean vector. - Component |i| of the result is `bool(`|e|`[`|i|`])` - (OpINotEqual to compare |e| against a zero vector.) - - |
|e| : vec|N|<i32> - | `vec`|N|<`bool`>`(`|e|`)` : vec|N|<bool> - | Component-wise coercion of a signed integer vector to a boolean vector. - Component |i| of the result is `bool(`|e|`[`|i|`])` - (OpINotEqual to compare |e| against a zero vector.) - - |
|e| : vec|N|<f32> - | `vec`|N|<`bool`>`(`|e|`)` : vec|N|<bool> - | Component-wise coercion of a floating point vector to a boolean vector. - Component |i| of the result is `bool(`|e|`[`|i|`])` - (OpFUnordNotEqual to compare |e| against a zero vector.) - - |
|e| : vec|N|<u32> - | `vec`|N|<`i32`>`(`|e|`)` : vec|N|<i32> - | Component-wise reinterpretation of bits. - Component |i| of the result is `i32(`|e|`[`|i|`])` - (OpBitcast) - - |
|e| : vec|N|<f32> - | `vec`|N|<`i32`>`(`|e|`)` : vec|N|<i32> - | Component-wise value conversion to signed integer, including invalid cases. - Component |i| of the result is `i32(`|e|`[`|i|`])` - (OpConvertFToS) - - |
|e| : vec|N|<i32> - | `vec`|N|<`u32`>`(`|e|`)` : vec|N|<u32> - | Component-wise reinterpretation of bits. - Component |i| of the result is `u32(`|e|`[`|i|`])` - (OpBitcast) - - |
|e| : vec|N|<f32> - | `vec`|N|<`u32`>`(`|e|`)` : vec|N|<u32> - | Component-wise value conversion to unsigned integer, including invalid cases. - Component |i| of the result is `u32(`|e|`[`|i|`])` - (OpConvertFToU) - - |
|e| : vec|N|<i32> - | `vec`|N|<`f32`>`(`|e|`)` : vec|N|<f32> - | Component-wise value conversion to floating point, including invalid cases. - Component |i| of the result is `f32(`|e|`[`|i|`])` - (OpConvertSToF) - - |
|e| : vec|N|<u32> - | `vec`|N|<`f32`>`(`|e|`)` : vec|N|<f32> - | Component-wise value conversion to floating point, including invalid cases. - Component |i| of the result is `f32(`|e|`[`|i|`])` - (ConvertUToF) - - |
Precondition | Conclusion | Notes - |
---|---|---|
|e| : |T|, - |T| is one of i32, u32, f32 - | bitcast<|T|>(|e|) : |T| - | Identity transform. - The result is |e|. - In the SPIR-V translation, the ID of this expression reuses the ID of the operand. - |
|e| : |T|, - |T| is one of u32, f32 - | bitcast<i32>(|e|) : i32 - | Reinterpretation of bits as a signed integer. - The result is the reinterpretation of the 32 bits in the representation of |e| as a [=i32=] value. - (OpBitcast) - |
|e| : |T|, - |T| is one of i32, f32 - | bitcast<u32>(|e|) : u32 - | Reinterpretation of bits as an unsigned integer. - The result is the reinterpretation of the 32 bits in the representation of |e| as a [=u32=] value. - (OpBitcast) - |
|e| : |T|, - |T| is one of i32, u32 - | bitcast<f32>(|e|) : f32 - | Reinterpretation of bits as a floating point value. - The result is the reinterpretation of the 32 bits in the representation of |e| as a [=f32=] value. - (OpBitcast) - |
Precondition | Conclusion | Notes - | ||||||
---|---|---|---|---|---|---|---|---|
|e| : vec<|N|>|T|>, - |T| is one of i32, u32, f32 - | bitcast<vec|N|<|T|>>(|e|) : |T| - | Identity transform. - The result is |e|. - In the SPIR-V translation, the ID of this expression reuses the ID of the operand. - | ||||||
|e| : vec<|N|>|T|>, - |T| is one of u32, f32 - | bitcast<vec|N|<i32>>(|e|) : vec|N|<i32> - | Component-wise reinterpretation of bits. - Component |i| of the result is `bitcast - (OpBitcast) - |e| : vec<|N|>|T|>, | - |T| is one of i32, f32 - bitcast<vec|N|<u32>>(|e|) : vec|N|<u32>
- | Component-wise reinterpretation of bits. | - Component |i| of the result is `bitcast - (OpBitcast) - |e| : vec<|N|>|T|>, | - |T| is one of i32, u32 - bitcast<vec|N|<f32>>(|e|) : vec|N|<f32>
- | Component-wise Reinterpretation of bits. | - Component |i| of the result is `bitcast - (OpBitcast) - - |
Accessor | Result type - | ||||
---|---|---|---|---|---|
r | `f32` - | ||||
rg | `vec2rgb | `vec3 | rgba | `vec4 | |
Precondition | Conclusion | Description - |
---|---|---|
|e| : vec|N|<|T|> - |
- |e|`.x` : |T| - |e|`.r` : |T| - | Select the first component of |e| - (OpCompositeExtract with selection index 0) - |
|e| : vec|N|<|T|> - |
- |e|`.y` : |T| - |e|`.g` : |T| - | Select the second component of |e| - (OpCompositeExtract with selection index 1) - |
|e| : vec|N|<|T|> - |N| is 3 or 4 - |
- |e|`.z` : |T| - |e|`.b` : |T| - | Select the third component of |e| - (OpCompositeExtract with selection index 2) - |
|e| : vec4<|T|> - |
- |e|`.w` : |T| - |e|`.a` : |T| - | Select the fourth component of |e| - (OpCompositeExtract with selection index 3) - |
|e| : vec|N|<|T|> - |i| : *Int* - | - |e|[|i|] : |T| - | Select the |i|'th component of vector - The first component is at index |i|=0. - If |i| is outside the range [0,|N|-1], then an index in the range [0, |N|-1] is used instead. - (OpVectorExtractDynamic) - |
Precondition | Conclusion | Description - |
---|---|---|
- |e| : vec|N|<|T|> - |I| is the letter `x`, `y`, `z`, or `w` - |J| is the letter `x`, `y`, `z`, or `w` - |
- |e|`.`|I||J| : vec2<|T|> - | Computes the two-element vector with first component |e|.|I|, and second component |e|.|J|. - Letter `z` is valid only when |N| is 3 or 4. - Letter `w` is valid only when |N| is 4. - (OpVectorShuffle) - |
- |e| : vec|N|<|T|> - |I| is the letter `r`, `g`, `b`, or `a` - |J| is the letter `r`, `g`, `b`, or `a` - |
- |e|`.`|I||J| : vec2<|T|> - | Computes the two-element vector with first component |e|.|I|, and second component |e|.|J|. - Letter `b` is valid only when |N| is 3 or 4. - Letter `a` is valid only when |N| is 4. - (OpVectorShuffle) - |
- |e| : vec|N|<|T|> - |I| is the letter `x`, `y`, `z`, or `w` - |J| is the letter `x`, `y`, `z`, or `w` - |K| is the letter `x`, `y`, `z`, or `w` - |
- |e|`.`|I||J||K| : vec3<|T|> - | Computes the three-element vector with first component |e|.|I|, second component |e|.|J|, and third component |e|.|K|. - Letter `z` is valid only when |N| is 3 or 4. - Letter `w` is valid only when |N| is 4. - (OpVectorShuffle) - |
- |e| : vec|N|<|T|> - |I| is the letter `r`, `g`, `b`, or `a` - |J| is the letter `r`, `g`, `b`, or `a` - |K| is the letter `r`, `g`, `b`, or `a` - |
- |e|`.`|I||J||K| : vec3<|T|> - | Computes the three-element vector with first component |e|.|I|, second component |e|.|J|, and third component |e|.|K|. - Letter `b` is only valid when |N| is 3 or 4. - Letter `a` is only valid when |N| is 4. - (OpVectorShuffle) - |
- |e| : vec|N|<|T|> - |I| is the letter `x`, `y`, `z`, or `w` - |J| is the letter `x`, `y`, `z`, or `w` - |K| is the letter `x`, `y`, `z`, or `w` - |L| is the letter `x`, `y`, `z`, or `w` - |
- |e|`.`|I||J||K||L| : vec4<|T|> - | Computes the four-element vector with first component |e|.|I|, second component |e|.|J|, third component |e|.|K|, and fourth component |e|.|L|. - Letter `z` is valid only when |N| is 3 or 4. - Letter `w` is valid only when |N| is 4. - (OpVectorShuffle) - |
- |e| : vec|N|<|T|> - |I| is the letter `r`, `g`, `b`, or `a` - |J| is the letter `r`, `g`, `b`, or `a` - |K| is the letter `r`, `g`, `b`, or `a` - |L| is the letter `r`, `g`, `b`, or `a` - |
- |e|`.`|I||J||K||L| : vec4<|T|> - | Computes the four-element vector with first component |e|.|I|, second component |e|.|J|, third component |e|.|K|, and fourth component |e|.|L|. - Letter `b` is only valid when |N| is 3 or 4. - Letter `a` is only valid when |N| is 4. - (OpVectorShuffle) - |
Precondition | Conclusion | Description - |
---|---|---|
|r| : ref<|SC|,vec|N|<|T|>> - |
- |r|`.x` : ref<|SC|,|T|> - |r|`.r` : ref<|SC|,|T|> - | Compute a reference to the first component of the vector referenced by the reference |r|. - The [=originating variable=] of the resulting reference is - the same as the originating variable of |r|. - (OpAccessChain with index value 0) - |
|r| : ref<|SC|,vec|N|<|T|>> - |
- |r|`.y` : ref<|SC|,|T|> - |r|`.g` : ref<|SC|,|T|> - | Compute a reference to the second component of the vector referenced by the reference |r|. - The [=originating variable=] of the resulting reference is - the same as the originating variable of |r|. - (OpAccessChain with index value 1) - |
|r| : ref<|SC|,vec|N|<|T|>> - |N| is 3 or 4 - |
- |r|`.z` : ref<|SC|,|T|> - |r|`.b` : ref<|SC|,|T|> - | Compute a reference to the third component of the vector referenced by the reference |r|. - The [=originating variable=] of the resulting reference is - the same as the originating variable of |r|. - (OpAccessChain with index value 2) - |
|r| : ref<|SC|,vec4<|T|>> - |
- |r|`.w` : ref<|SC|,|T|> - |r|`.a` : ref<|SC|,|T|> - | Compute a reference to the fourth component of the vector referenced by the reference |r|. - The [=originating variable=] of the resulting reference is - the same as the originating variable of |r|. - (OpAccessChain with index value 3) - |
|r| : ref<|SC|,vec|N|<|T|>> - |i| : *Int* - | - |r|[|i|] : ref<|SC|,|T|> - | Compute a reference to the |i|'th component of the vector referenced by the reference |r|. - If |i| is outside the range [0,|N|-1], then an index in the range [0, |N|-1] is used instead. - The [=originating variable=] of the resulting reference is - the same as the originating variable of |r|. - (OpAccessChain) - |
Precondition | Conclusion | Description - |
---|---|---|
- |e| : mat|N|x|M|<|T|> - |i| : *Int* - | - |e|[|i|] : vec|M|<|T|> - | The result is the |i|'th column vector of |e|. - If |i| is outside the range [0,|N|-1], then an index in the range [0, |N|-1] is used instead. - (OpCompositeExtract) - |
Precondition | Conclusion | Description - |
---|---|---|
- |r| : ref<|SC|,mat|N|x|M|<|T|>> - |i| : *Int* - | - |r|[|i|] : ref<vec|M|<|SC|,|T|>> - | Compute a reference to the |i|'th column vector of the matrix referenced by the reference |r|. - If |i| is outside the range [0,|N|-1], then an index in the range [0, |N|-1] is used instead. - The [=originating variable=] of the resulting reference is - the same as the originating variable of |r|. - (OpAccessChain) - |
Precondition | Conclusion | Description - |
---|---|---|
- |e| : array<|T|,|N|> - |i| : *Int* - | - |e|[|i|] : |T| - | The result is the value of the |i|'th element of the array value |e|. - If |i| is outside the range [0,|N|-1], then an index in the range [0, |N|-1] is used instead. - (OpCompositeExtract) - |
Precondition | Conclusion | Description - |
---|---|---|
- |r| : ref<|SC|,array<|T|,|N|>> - |i| : *Int* - | - |r|[|i|] : ref<|SC|,|T|> - | Compute a reference to the |i|'th element of the array referenced by the reference |r|. - If |i| is outside the range [0,|N|-1], then an index in the range [0, |N|-1] is used instead. - The [=originating variable=] of the resulting reference is - the same as the originating variable of |r|. - (OpAccessChain) - |
|r| : ref<|SC|,array<|T|>> - |i| : *Int* - | - |r|[|i|] : ref<|SC|,|T|> - | Compute a reference to the |i|'th element of the runtime-sized array referenced by the reference |r|. - If at runtime the array has |N| elements, and |i| is outside the range [0,|N|-1], then an index in the - range [0, |N|-1] is used instead. - The [=originating variable=] of the resulting reference is - the same as the originating variable of |r|. - (OpAccessChain) - |
Precondition | Conclusion | Description - |
---|---|---|
- |S| is a structure type - |M| is the identifier name of a member of |S|, having type |T| - |e| : |S| - | - |e|.|M| : |T| - | The result is the value of the member with name |M| from the structure value |e|. - (OpCompositeExtract, using the member index) - |
Precondition | Conclusion | Description - |
---|---|---|
- |S| is a structure type - |M| is the name of a member of |S|, having type |T| - |r| : ref<|SC|,|S|> - | - |r|.|M| : ref<|SC|,|T|> - | Given a reference to a structure, the result is a reference to the structure member with identifier name |M|. - The [=originating variable=] of the resulting reference is - the same as the originating variable of |r|. - (OpAccessChain, using the index of the structure member) - |
Precondition | Conclusion | Notes - |
---|---|---|
|e| : bool | `!`|e| : *bool* - | Logical negation. Yields true when |e| is false, and false when |e| is true. (OpLogicalNot) - |
|e| : vec|N|<bool> | `!`|e| : vec|N|<bool> - | Component-wise logical negation. Component |i| of the result is `!(`|e|`[`|i|`])`. (OpLogicalNot) - |
Precondition | Conclusion | Notes - |
---|---|---|
*e1* : bool *e2* : bool | `e1 || e2` : bool | - Short-circuiting "or". Yields `true` if either `e1` or `e2` are true; evaluates `e2` only if `e1` is false. - |
*e1* : bool *e2* : bool | `e1 && e2` : bool | - Short-circuiting "and". Yields `true` if both `e1` and `e2` are true; evaluates `e2` only if `e1` is true. - |
*e1* : bool *e2* : bool | `e1 | e2` : bool | - Logical "or". Evaluates both `e1` and `e2`; yields `true` if either are `true`. - |
*e1* : bool *e2* : bool | `e1 & e2` : bool | - Logical "and". Evaluates both `e1` and `e2`; yields `true` if both are `true`. - |
*e1* : *T* *e2* : *T* *T* is *BoolVec* | `e1 | e2` : *T* | Component-wise logical "or" - |
*e1* : *T* *e2* : *T* *T* is *BoolVec* | `e1 & e2` : *T* | Component-wise logical "and" - |
Precondition | Conclusion | Notes - |
---|---|---|
*e* : *T*, *T* is *SignedIntegral* | `-e` : *T* | Signed integer negation. OpSNegate - |
*e* : *T*, *T* is *Floating* | `-e` : *T* | Floating point negation. OpFNegate - |
Precondition | Conclusion | Notes - |
---|---|---|
*e1* : u32 *e2* : u32 | `e1 + e2` : u32 | Integer addition, modulo 232 (OpIAdd) - |
*e1* : i32 *e2* : i32 | `e1 + e2` : i32 | Integer addition, modulo 232 (OpIAdd) - |
*e1* : f32 *e2* : f32 | `e1 + e2` : f32 | Floating point addition (OpFAdd) - |
*e1* : u32 *e2* : u32 | `e1 - e2` : u32 | Integer subtraction, modulo 232 (OpISub) - |
*e1* : i32 *e2* : i32 | `e1 - e2` : i32 | Integer subtraction, modulo 232 (OpISub) - |
*e1* : f32 *e2* : f32 | `e1 - e2` : f32 | Floating point subtraction (OpFSub) - |
*e1* : u32 *e2* : u32 | `e1 * e2` : u32 | Integer multiplication, modulo 232 (OpIMul) - |
*e1* : i32 *e2* : i32 | `e1 * e2` : i32 | Integer multiplication, modulo 232 (OpIMul) - |
*e1* : f32 *e2* : f32 | `e1 * e2` : f32 | Floating point multiplication (OpFMul) - |
*e1* : u32 *e2* : u32 | `e1 / e2` : u32 | Unsigned integer division (OpUDiv) - |
*e1* : i32 *e2* : i32 | `e1 / e2` : i32 | Signed integer division (OpSDiv) - |
*e1* : f32 *e2* : f32 | `e1 / e2` : f32 | Floating point division (OpFDiv) - |
*e1* : u32 *e2* : u32 | `e1 % e2` : u32 | Unsigned integer modulus (OpUMod) - |
*e1* : i32 *e2* : i32 | `e1 % e2` : i32 | Signed integer remainder, where sign of non-zero result matches sign of *e2* (OpSMod) - |
*e1* : f32 *e2* : f32 | `e1 % e2` : f32 | Floating point modulus, where sign of non-zero result matches sign of *e2* (OpFMod) - |
Precondition | Conclusion | Notes - |
---|---|---|
*e1* : *T* *e2* : *T* *T* is *IntVec* | `e1 + e2` : *T* | Component-wise integer addition (OpIAdd) - |
*e1* : *T* *e2* : *T* *T* is *FloatVec* | `e1 + e2` : *T* | Component-wise floating point addition (OpIAdd) - |
*e1* : *T* *e2* : *T* *T* is *IntVec* | `e1 - e2` : *T* | Component-wise integer subtraction (OpISub) - |
*e1* : *T* *e2* : *T* *T* is *FloatVec* | `e1 - e2` : *T* | Component-wise floating point subtraction (OpISub) - |
*e1* : *T* *e2* : *T* *T* is *IntVec* | `e1 * e2` : *T* | Component-wise integer multiplication (OpIMul) - |
*e1* : *T* *e2* : *T* *T* is *FloatVec* | `e1 * e2` : *T* | Component-wise floating point multiplication (OpIMul) - |
*e1* : *T* *e2* : *T* *T* is *IntVec* with unsigned component | `e1 / e2` : *T* | Component-wise unsigned integer division (OpUDiv) - |
*e1* : *T* *e2* : *T* *T* is *IntVec* with signed component | `e1 / e2` : *T* | Component-wise signed integer division (OpSDiv) - |
*e1* : *T* *e2* : *T* *T* is *FloatVec* | `e1 / e2` : *T* | Component-wise floating point division (OpFDiv) - |
*e1* : *T* *e2* : *T* *T* is *IntVec* with unsigned component | `e1 % e2` : *T* | Component-wise unsigned integer modulus (OpUMod) - |
*e1* : *T* *e2* : *T* *T* is *IntVec* with signed component | `e1 % e2` : *T* | Component-wise signed integer remainder (OpSMod) - |
*e1* : *T* *e2* : *T* *T* is *FloatVec* | `e1 % e2` : *T* | Component-wise floating point modulus (OpFMod) - |
Preconditions | Conclusions | Semantics - | -
---|---|---|
|S| is one of f32, i32, u32 - |V| is vec|N|<|S|> - |es|: |S| - |ev|: |V| - | |ev| `+` |es|: |V| - | |ev| `+` |V|(|es|) - |
|es| `+` |ev|: |V| - | |V|(|es|) `+` |ev| - | |
|ev| `-` |es|: |V| - | |ev| `-` |V|(|es|) - | |
|es| `-` |ev|: |V| - | |V|(|es|) `-` |ev| - | |
|ev| `*` |es|: |V| - | |ev| `*` |V|(|es|) - | |
|es| `*` |ev|: |V| - | |V|(|es|) `*` |ev| - | |
|ev| `/` |es|: |V| - | |ev| `/` |V|(|es|) - | |
|es| `/` |ev|: |V| - | |V|(|es|) `/` |ev| - | |
|S| is one of i32, u32 - |V| is vec|N|<|S|> - |es|: |S| - |ev|: |V| - | |ev| `%` |es|: |V| - | |ev| `%` |V|(|es|) - |
|es| `%` |ev|: |V| - | |V|(|es|) `%` |ev| - |
Preconditions | Conclusions | Semantics - | -
---|---|---|
|e1|, |e2|: mat|M|x|N|<f32> - | |e1| `+` |e2|: mat|M|x|N|<f32> - | Matrix addition: column |i| of the result is |e1|[i] + |e2|[i] - |
|e1| `-` |e2|: mat|M|x|N|<f32> - | Matrix subtraction: column |i| of the result is |e1|[|i|] - |e2|[|i|] - | |
|m|: mat|M|x|N|<f32> - |s|: f32 - | |m| `*` |s| : mat|M|x|N|<f32> - | Component-wise scaling: (|m| `*` |s|)[i][j] is |m|[i][j] `*` |s| - |
|s| `*` |m| : mat|M|x|N|<f32> - | Component-wise scaling: (|s| `*` |m|)[i][j] is |m|[i][j] `*` |s| - | |
|m|: mat|M|x|N|<f32> - |v|: vec|M|<f32> - | |m| `*` |v| : vec|N|<f32> - | Linear algebra matrix-column-vector product:
- Component |i| of the result is `dot`(|m|[|i|],|v|)
- OpMatrixTimesVector - |
- |m|: mat|M|x|N|<f32> - |v|: vec|N|<f32> - | |v| `*` |m| : vec|M|<f32> - | Linear algebra row-vector-matrix product: - [=transpose=](transpose(|m|) `*` transpose(|v|)) - OpVectorTimesMatrix - |
|e1|: mat|K|x|N|<f32> - |e2|: mat|M|x|K|<f32> - | |e1| `*` |e2| : mat|M|x|N|<f32> - | Linear algebra matrix product. OpMatrixTimesMatrix - - |
Precondition | Conclusion | Notes - |
---|---|---|
*e1* : bool - *e2* : bool - | `e1 == e2` : bool - | Equality (OpLogicalEqual) - |
*e1* : bool - *e2* : bool - | `e1 != e2` : bool - | Inequality (OpLogicalNotEqual) - |
*e1* : i32 - *e2* : i32 - | `e1 == e2` : bool - | Equality (OpIEqual) - |
*e1* : i32 - *e2* : i32 - | `e1 != e2` : bool - | Inequality (OpINotEqual) - |
*e1* : i32 - *e2* : i32 - | `e1 < e2` : bool - | Less than (OpSLessThan) - |
*e1* : i32 - *e2* : i32 - | `e1 <= e2` : bool - | Less than or equal (OpSLessThanEqual) - |
*e1* : i32 - *e2* : i32 - | `e1 >= e2` : bool - | Greater than or equal (OpSGreaterThanEqual) - |
*e1* : i32 - *e2* : i32 - | `e1 > e2` : bool - | Greater than (OpSGreaterThan) - |
*e1* : u32 - *e2* : u32 - | `e1 == e2` : bool - | Equality (OpIEqual) - |
*e1* : u32 - *e2* : u32 - | `e1 != e2` : bool - | Inequality (OpINotEqual) - |
*e1* : u32 - *e2* : u32 - | `e1 < e2` : bool - | Less than (OpULessThan) - |
*e1* : u32 - *e2* : u32 - | `e1 <= e2` : bool - | Less than or equal (OpULessThanEqual) - |
*e1* : u32 - *e2* : u32 - | `e1 >= e2` : bool - | Greater than or equal (OpUGreaterThanEqual) - |
*e1* : u32 - *e2* : u32 - | `e1 > e2` : bool - | Greater than (OpUGreaterThan) - |
*e1* : f32 - *e2* : f32 - | `e1 == e2` : bool - | Equality (OpFOrdEqual) - |
*e1* : f32 - *e2* : f32 - | `e1 != e2` : bool - | Equality (OpFOrdNotEqual) - |
*e1* : f32 - *e2* : f32 - | `e1 < e2` : bool - | Less than (OpFOrdLessThan) - |
*e1* : f32 - *e2* : f32 - | `e1 <= e2` : bool - | Less than or equal (OpFOrdLessThanEqual) - |
*e1* : f32 - *e2* : f32 - | `e1 >= e2` : bool - | Greater than or equal (OpFOrdGreaterThanEqual) - |
*e1* : f32 - *e2* : f32 - | `e1 > e2` : bool - | Greater than (OpFOrdGreaterThan) - |
Precondition | Conclusion | Notes - |
---|---|---|
*e1* : *T* - *e2* : *T* - *T* is vec*N*<bool> - | `e1 == e2` : vec*N*<bool> - | Component-wise equality - Component |i| of the result is `(`|e1|`[`|i|`] == `|e2|`[`|i|`])` - (OpLogicalEqual) - |
*e1* : *T* - *e2* : *T* - *T* is vec*N*<bool> - | `e1 != e2` : vec*N*<bool> - | Component-wise inequality - Component |i| of the result is `(`|e1|`[`|i|`] != `|e2|`[`|i|`])` - (OpLogicalNotEqual) - |
*e1* : *T* - *e2* : *T* - *T* is vec*N*<i32> - | `e1 == e2` : vec*N*<bool> - | Component-wise equality (OpIEqual) - |
*e1* : *T* - *e2* : *T* - *T* is vec*N*<i32> - | `e1 != e2` : vec*N*<bool> - | Component-wise inequality (OpINotEqual) - |
*e1* : *T* - *e2* : *T* - *T* is vec*N*<i32> - | `e1 < e2` : vec*N*<bool> - | Component-wise less than (OpSLessThan) - |
*e1* : *T* - *e2* : *T* - *T* is vec*N*<i32> - | `e1 <= e2` : vec*N*<bool> - | Component-wise less than or equal (OpSLessThanEqual) - |
*e1* : *T* - *e2* : *T* - *T* is vec*N*<i32> - | `e1 >= e2` : vec*N*<bool> - | Component-wise greater than or equal (OpSGreaterThanEqual) - |
*e1* : *T* - *e2* : *T* - *T* is vec*N*<i32> - | `e1 > e2` : vec*N*<bool> - | Component-wise greater than (OpSGreaterThan) - |
*e1* : *T* - *e2* : *T* - *T* is vec*N*<u32> - | `e1 == e2` : vec*N*<bool> - | Component-wise equality (OpIEqual) - |
*e1* : *T* - *e2* : *T* - *T* is vec*N*<u32> - | `e1 != e2` : vec*N*<bool> - | Component-wise inequality (OpINotEqual) - |
*e1* : *T* - *e2* : *T* - *T* is vec*N*<u32> - | `e1 < e2` : vec*N*<bool> - | Component-wise less than (OpULessThan) - |
*e1* : *T* - *e2* : *T* - *T* is vec*N*<u32> - | `e1 <= e2` : vec*N*<bool> - | Component-wise less than or equal (OpULessThanEqual) - |
*e1* : *T* - *e2* : *T* - *T* is vec*N*<u32> - | `e1 >= e2` : vec*N*<bool> - | Component-wise greater than or equal (OpUGreaterThanEqual) - |
*e1* : *T* - *e2* : *T* - *T* is vec*N*<u32> - | `e1 > e2` : vec*N*<bool> - | Component-wise greater than (OpUGreaterThan) - *T* is vec*N*<u32> - |
*e1* : *T* - *e2* : *T* - *T* is vec*N*<f32> - | `e1 == e2` : vec*N*<bool> - | Component-wise equality (OpFOrdEqual) - |
*e1* : *T* - *e2* : *T* - *T* is vec*N*<f32> - | `e1 != e2` : vec*N*<bool> - | Component-wise inequality (OpFOrdNotEqual) - |
*e1* : *T* - *e2* : *T* - *T* is vec*N*<f32> - | `e1 < e2` : vec*N*<bool> - | Component-wise less than (OpFOrdLessThan) - |
*e1* : *T* - *e2* : *T* - *T* is vec*N*<f32> - | `e1 <= e2` : vec*N*<bool> - | Component-wise less than or equal (OpFOrdLessThanEqual) - |
*e1* : *T* - *e2* : *T* - *T* is vec*N*<f32> - | `e1 >= e2` : vec*N*<bool> - | Component-wise greater than or equal (OpFOrdGreaterThanEqual) - |
*e1* : *T* - *e2* : *T* - *T* is vec*N*<f32> - | `e1 > e2` : vec*N*<bool> - | Component-wise greater than (OpFOrdGreaterThan) - |
Precondition | Conclusion | Notes - |
---|---|---|
|e| : u32 - | `~`|e| : u32 - | Bitwise complement on unsigned integers. Result is the mathematical value (232 - 1 - |e|).
- OpNot - |
|e| : vec|N|<u32> - | `~`|e| : vec|N|<u32> - | Component-wise unsigned complement. Component |i| of the result is `~(`|e|`[`|i|`])`.
- OpNot - |
|e| : i32 - | `~`|e| : i32 - | Bitwise complement on signed integers. Result is i32(~u32(|e|)).
- OpNot - |
|e| : vec|N|<i32> - | `~`|e| : vec|N|<i32> - | Component-wise signed complement. Component |i| of the result is `~(`|e|`[`|i|`])`.
- OpNot - |
Precondition | Conclusion | Notes - |
---|---|---|
*e1* : *T* - *e2* : *T* - *T* is *Integral* - | `e1 | e2` : *T* - | Bitwise-or - |
*e1* : *T* - *e2* : *T* - *T* is *Integral* - | `e1 & e2` : *T* - | Bitwise-and - |
*e1* : *T* - *e2* : *T* - *T* is *Integral* - | `e1 ^ e2` : *T* - | Bitwise-exclusive-or - |
Precondition | Conclusion | Notes - |
---|---|---|
|e1| : |T| - |e2| : u32 - |T| is *Int* - | |e1| `<<` |e2| : |T| - | Shift left: - Shift |e1| left, inserting zero bits at the least significant positions, - and discarding the most significant bits. - The number of bits to shift is the value of |e2| modulo the bit width of |e1|. - (OpShiftLeftLogical) - |
|e1| : vec|N|<|T|> - |e2| : vec|N|<u32> - |T| is *Int* - | |e1| `<<` |e2| : vec|N|<|T|> - | Component-wise shift left: - Component |i| of the result is `(`|e1|`[`|i|`] << `|e2|`[`|i|`])` - (OpShiftLeftLogical) - |
|e1| : u32 - |e2| : u32 - | |e1| `>>` |e2| `: u32` - | Logical shift right: - Shift |e1| right, inserting zero bits at the most significant positions, - and discarding the least significant bits. - The number of bits to shift is the value of |e2| modulo the bit width of |e1|. - (OpShiftRightLogical) - |
|e1| : vec|N|<u32> - |e2| : vec|N|<u32> - | |e1| `>>` |e2| : vec|N|<u32> - | Component-wise logical shift right: - Component |i| of the result is `(`|e1|`[`|i|`] >> `|e2|`[`|i|`])` - (OpShiftRightLogical) - |
|e1| : i32 - |e2| : u32 - | |e1| `>>` |e2| : i32 - | Arithmetic shift right: - Shift |e1| right, copying the sign bit of |e1| into the most significant positions, - and discarding the least significant bits. - The number of bits to shift is the value of |e2| modulo the bit width of |e1|. - (OpShiftRightArithmetic) - |
|e1| : vec|N|<i32> - |e2| : vec|N|<u32> - | |e1| `>>` |e2| : vec|N|<i32> - | Component-wise arithmetic shift right: - Component |i| of the result is `(`|e1|`[`|i|`] >> `|e2|`[`|i|`])` - (OpShiftRightArithmetic) - |
Precondition | Conclusion | Description - |
---|---|---|
- |v| is an identifier [=resolves|resolving=] to - an [=in scope|in-scope=] variable declared in [=storage class=] |SC| - with [=store type=] |T| - | - |v| : ref<|SC|,|T|> - | Result is a reference to the storage for the named variable |v|. - |
Precondition | Conclusion | Description - |
---|---|---|
- |a| is an identifier [=resolves|resolving=] to - an [=in scope|in-scope=] formal paramter declaration with type |T| - | - |a| : |T| - | Result is the value supplied for the corresponding function call operand at the call site - invoking this instance of the function. - |
Precondition | Conclusion | Description - |
---|---|---|
- |r| : ref<|SC|,|T|> - | - `&`|r| : ptr<|SC|,|T|> - | Result is the pointer value corresponding to the - same [=memory view=] as the reference value |r|. - |
Precondition | Conclusion | Description - |
---|---|---|
- |p| : ptr<|SC|,|T|> - | - `*`|p| : ref<|SC|,|T|> - | Result is the reference value corresponding to the - same [=memory view=] as the pointer value |p|. - |
Precondition | Conclusion | Description - |
---|---|---|
- |c| is an identifier [=resolves|resolving=] to - an [=in scope|in-scope=] [=pipeline-overridable=] `let` declaration with type |T| - | - |c| : |T| - | If pipeline creation specified a value for the [=pipeline constant ID|constant ID=],
- then the result is that value.
- This value may be different for different pipeline instances. - Otherwise, the result is the value computed for the initializer expression. - Pipeline-overridable constants appear at module-scope, so evaluation occurs - before the shader begins execution. - Note: Pipeline creation fails if no initial value was specified in the API call - and the `let`-declaration has no intializer expression. - |
- |c| is an identifier [=resolves|resolving=] to - an [=in scope|in-scope=] `let` declaration with type |T|, - and is not pipeline-overridable - | - |c| : |T| - | Result is the value computed for the initializer expression. - For a `let` declaration at module scope, evaluation occurs before the shader begins execution. - For a `let` declaration inside a function, evaluation occurs each time control reaches - the declaration. - |
-primary_expression - : IDENT argument_expression_list? - | type_decl argument_expression_list - | const_literal - | paren_rhs_statement - | BITCAST LESS_THAN type_decl GREATER_THAN paren_rhs_statement - OpBitcast - -argument_expression_list - : PAREN_LEFT ((short_circuit_or_expression COMMA)* short_circuit_or_expression)? PAREN_RIGHT - -postfix_expression - : - | BRACKET_LEFT short_circuit_or_expression BRACKET_RIGHT postfix_expression - | PERIOD IDENT postfix_expression - -unary_expression - : singular_expression - | MINUS unary_expression - OpSNegate - OpFNegate - | BANG unary_expression - OpLogicalNot - | TILDE unary_expression - OpNot - | STAR unary_expression - | AND unary_expression - -singular_expression - : primary_expression postfix_expression - -multiplicative_expression - : unary_expression - | multiplicative_expression STAR unary_expression - OpVectorTimesScalar - OpMatrixTimesScalar - OpVectorTimesMatrix - OpMatrixTimesVector - OpMatrixTimesMatrix - OpIMul - OpFMul - | multiplicative_expression FORWARD_SLASH unary_expression - OpUDiv - OpSDiv - OpFDiv - | multiplicative_expression MODULO unary_expression - OpUMOd - OpSMod - OpFMod - -additive_expression - : multiplicative_expression - | additive_expression PLUS multiplicative_expression - OpIAdd - OpFAdd - | additive_expression MINUS multiplicative_expression - OpFSub - OpISub - -shift_expression - : additive_expression - | shift_expression SHIFT_LEFT additive_expression - OpShiftLeftLogical - | shift_expression SHIFT_RIGHT additive_expression - OpShiftRightLogical or OpShiftRightArithmetic - -relational_expression - : shift_expression - | relational_expression LESS_THAN shift_expression - OpULessThan - OpFOrdLessThan - | relational_expression GREATER_THAN shift_expression - OpUGreaterThan - OpFOrdGreaterThan - | relational_expression LESS_THAN_EQUAL shift_expression - OpULessThanEqual - OpFOrdLessThanEqual - | relational_expression GREATER_THAN_EQUAL shift_expression - OpUGreaterThanEqual - OpFOrdGreaterThanEqual - -equality_expression - : relational_expression - | relational_expression EQUAL_EQUAL relational_expression - OpIEqual - OpFOrdEqual - | relational_expression NOT_EQUAL relational_expression - OpINotEqual - OpFOrdNotEqual - -and_expression - : equality_expression - | and_expression AND equality_expression - -exclusive_or_expression - : and_expression - | exclusive_or_expression XOR and_expression - -inclusive_or_expression - : exclusive_or_expression - | inclusive_or_expression OR exclusive_or_expression - -short_circuit_and_expression - : inclusive_or_expression - | short_circuit_and_expression AND_AND inclusive_or_expression - -short_circuit_or_expression - : short_circuit_and_expression - | short_circuit_or_expression OR_OR short_circuit_and_expression -- - -# Statements TODO # {#statements} - -## Compound Statement ## {#compound-statement} - -A compound statement is a brace-enclosed group of zero or more statements. -When a declaration is one of those statements, its identifier is [=in scope=] -from the start of the next statement until the end of the compound statement. - -
-compound_statement - : BRACE_LEFT statements BRACE_RIGHT -- -## Assignment Statement ## {#assignment} - -An assignment statement replaces the contents of a variable, -or a portion of a variable, with a new value. - -The -expression to the left of the equals token is the left-hand side, -and the -expression to the right of the equals token is the right-hand side. - -
Precondition | Statement | Description - |
---|---|---|
|r| : ref<|SC|,|T|>, - |e| : |T|, - |T| is [=storable=], - |SC| is a writable [=storage class=] - | |r| = |e|; - | Evaluates |e|, evaluates |r|, then writes the value computed for |e| into
- the [=memory locations=] referenced by |r|. - (OpStore) - |
-assignment_statement - : singular_expression EQUAL short_circuit_or_expression - If singular_expression is a variable, this maps to OpStore to the variable. - Otherwise, singular expression is a pointer expression in an Assigning (L-value) context - which maps to OpAccessChain followed by OpStore -- -## Control flow TODO ## {#control-flow} - -### Sequence TODO ### {#sequence-statement} - -### If/elseif/else Statement TODO ### {#if-statement} - -
-if_statement - : IF paren_rhs_statement compound_statement elseif_statement? else_statement? - -elseif_statement - : ELSE_IF paren_rhs_statement compound_statement elseif_statement? - -else_statement - : ELSE compound_statement -- - -### Switch Statement ### {#switch-statement} - -
-switch_statement - : SWITCH paren_rhs_statement BRACE_LEFT switch_body+ BRACE_RIGHT - -switch_body - : CASE case_selectors COLON BRACE_LEFT case_body BRACE_RIGHT - | DEFAULT COLON BRACE_LEFT case_body BRACE_RIGHT - -case_selectors - : const_literal (COMMA const_literal)* - -case_body - : - | statement case_body - | FALLTHROUGH SEMICOLON -- -A switch statement transfers control to one of a set of case clauses, or to the `default` clause, -depending on the evaluation of a selector expression. - -The selector expression must be of a scalar integer type. -If the selector value equals a value in a case selector list, then control is transferred to -the body of that case clause. -If the selector value does not equal any of the case selector values, then control is -transferred to the `default` clause. - -Each switch statement must have exactly one default clause. - -The case selector values must have the same type as the selector expression. - -A literal value must not appear more than once in the case selectors for a switch statement. - -Note: The value of the literal is what matters, not the spelling. -For example `0`, `00`, and `0x0000` all denote the zero value. - -When control reaches the end of a case body, control normally transfers to the first statement -after the switch statement. -Alternately, executing a `fallthrough` statement transfers control to the body of the next case clause or -default clause, whichever appears next in the switch body. -A `fallthrough` statement must not appear as the last statement in the last clause of a switch. -When a declaration appears in a case body, its identifier is [=in scope=] from -the start of the next statement until the end of the case body. - -Note: Identifiers declared in a case body are not [=in scope=] of case bodies -which are reachable via a `fallthrough` statement. - - -### Loop Statement ### {#loop-statement} - -
-loop_statement - : LOOP BRACE_LEFT statements continuing_statement? BRACE_RIGHT -- -The loop body is special form [compound -statement](#compound-statement) that executes repeatedly. -Each execution of the loop body is called an iteration. - -The identifier of a declaration in a loop is [=in scope=] from the start of the -next statement until the end of the loop body. -The declaration is executed each time it is reached, so each new iteration -creates a new instance of the variable or constant, and re-initializes it. - -This repetition can be interrupted by a [[#break-statement]], `return`, or -`discard`. - -Optionally, the last statement in the loop body may be a -[[#continuing-statement]]. - -Note: The loop statement is one of the biggest differences from other shader -languages. - -This design directly expresses loop idioms commonly found in compiled code. -In particular, placing the loop update statements at the end of the loop body -allows them to naturally use values defined in the loop body. - -
-for_statement - : FOR PAREN_LEFT for_header PAREN_RIGHT compound_statement - -for_header - : (variable_statement | assignment_statement | func_call_statement)? SEMICOLON - short_circuit_or_expression? SEMICOLON - (assignment_statement | func_call_statement)? -- -The `for(initializer; condition; continuing) { body }` statement is syntactic sugar on top of a [[#loop-statement]] with the same `body`. Additionally: -* If `initializer` is non-empty, it is executed inside an additional scope before the first iteration. -* If `condition` is non-empty, it is checked at the beginning of the loop body and if unsatisfied then a [[#break-statement]] is executed. -* If `continuing` is non-empty, it becomes a [[#continuing-statement]] at the end of the loop body. - -The `initializer` of a for loop is executed once prior to executing the loop. -When a declaration appears in the initializer, its identifier is [=in scope=] until the end of the `body`. -Unlike declarations in the `body`, the declaration is not re-initialized each iteration. - -The `condition`, `body` and `continuing` execute in that order to form a loop [=iteration=]. -The `body` is a special form of [compound statement](#compound-statement). -The identifier of a declaration in the `body` is [=in scope=] from the start of -the next statement until the end of the `body`. -The declaration is executed each time it is reached, so each new iteration -creates a new instance of the variable or constant, and re-intializes it. - -
-break_statement - : BREAK -- -Use a `break` statement to transfer control to the first statement -after the body of the nearest-enclosing [[#loop-statement]] -or [[#switch-statement]]. - -When a `break` statement is placed such that it would exit from a loop's [[#continuing-statement]], -then: - -* The `break` statement must appear as either: - * The only statement in the true-branch clause of an `if` that has: - * no `else` clause or an empty `else` clause - * no `elseif` clauses - * The only statement in the `else` clause of an `if` that has an empty true-branch clause and no `elseif` clauses. -* That `if` statement must appear last in the `continuing` clause. - -
-continue_statement - : CONTINUE -- -Use a `continue` statement to transfer control in the nearest-enclosing [[#loop-statement]]: - -* forward to the [[#continuing-statement]] at the end of the body of that loop, if it exists. -* otherwise backward to the first statement in the loop body, starting the next iteration - -A `continue` statement must not be placed such that it would transfer -control to an enclosing [[#continuing-statement]]. -(It is a *forward* branch when branching to a `continuing` statement.) - -A `continue` statement must not be placed such that it would transfer -control past a declaration used in the targeted continuing construct. - -
-continuing_statement - : CONTINUING compound_statement -- -A *continuing* construct is a block of statements to be executed at the end of a loop iteration. -The construct is optional. - -The block of statements must not contain a return or discard statement. - -### Return Statement ### {#return-statement} - -
-return_statement - : RETURN short_circuit_or_expression? -- -A return statement ends execution of the current function. -If the function is an [=entry point=], then the current shader invocation -is terminated. -Otherwise, evaluation continues with the next expression or statement after -the evaluation of the call site of the current function invocation. - -If the function doesn't have a [=return type=], then the return statement is -optional. If the return statement is provided for such a function, it must not -supply a value. -Otherwise the expression must be present, and is called the *return value*. -In this case the call site of this function invocation evaluates to the return value. -The type of the return value must match the return type of the function. - - -### Discard Statement ### {#discard-statement} - -The `discard` statement must only be used in a [=fragment=] shader stage. -Executing a `discard` statement will: - -* immediately terminate the current invocation, and -* prevent evaluation and generation of a return value for the [=entry point=], and -* prevent the current fragment from being processed downstream in the [=GPURenderPipeline=]. - -Only statements -executed prior to the `discard` statement will have observable effects. - -Note: A `discard` statement may be executed by any -[=functions in a shader stage|function in a fragment stage=] and the effect is the same: -immediate termination of the invocation. - -After a `discard` statement is executed, control flow is non-uniform for the -duration of the entry point. - -Issue: [[#uniform-control-flow]] needs to state whether all invocations being discarded maintains uniform control flow. - -
-func_call_statement - : IDENT argument_expression_list -- -## Statements Grammar Summary ## {#statements-summary} - -
-compound_statement - : BRACE_LEFT statements BRACE_RIGHT - -paren_rhs_statement - : PAREN_LEFT short_circuit_or_expression PAREN_RIGHT - -statements - : statement* - -statement - : SEMICOLON - | return_statement SEMICOLON - | if_statement - | switch_statement - | loop_statement - | for_statement - | func_call_statement SEMICOLON - | variable_statement SEMICOLON - | break_statement SEMICOLON - | continue_statement SEMICOLON - | DISCARD SEMICOLON - | assignment_statement SEMICOLON - | compound_statement -- - -# Functions # {#functions} - -A function performs computational work when invoked. - -A function is invoked in one of the following ways: -* By evaluating a function call expression. See [[#function-call-expr]]. -* By executing a function call statement. See [[#function-call-statement]]. -* An [=entry point=] function is invoked by the WebGPU implementation to perform - the work of a [=shader stage=] in a [=pipeline=]. See [[#entry-points]] - -There are two kinds of functions: -* A [=built-in function=] is provided by the [SHORTNAME] implementation, - and is always available to a [SHORTNAME] program. - See [[#builtin-functions]]. -* A user-defined function is declared in a [SHORTNAME] program. - -## Declaring a user-defined function ## {#function-declaration-sec} - -A function declaration creates a user-defined function, by specifying: -* An optional set of attributes. -* The name of the function. -* The formal parameter list: an ordered sequence of zero - or more [=formal parameter=] declarations, - separated by commas, and - surrounded by parentheses. -* An optional, possibly decorated, return type. -* The function body. - -A function declaration must only occur at [=module scope=]. -The function name is [=in scope=] from the start of the formal parameter list -until the end of the program. - -A formal parameter declaration specifies an identifier name and a type for a value that must be -provided when invoking the function. -A formal parameter may have attributes. -See [[#function-calls]]. -The identifier is [=in scope=] until the end of the function. -Two formal parameters for a given function must not have the same name. - -If the return type is specified, then: -* The return type must be a [=plain type=]. -* The last statement in the function body must be a [=return=] statement. - -
-function_decl - : attribute_list* function_header compound_statement - -function_header - : FN IDENT PAREN_LEFT param_list PAREN_RIGHT function_return_type_decl_optional - -function_return_type_decl_optional - : - | ARROW attribute_list* type_decl - -param_list - : - | (param COMMA)* param - -param - : attribute_list* variable_ident_decl -- -[SHORTNAME] defines the following attributes that can be applied to function declarations: - * [=attribute/stage=] - * [=attribute/workgroup_size=] - -[SHORTNAME] defines the following attributes that can be applied to function -parameters and return types: - * [=attribute/builtin=] - * [=attribute/location=] - -
[SHORTNAME] resource - | WebGPU [[WebGPU#enumdef-gpubindingtype|GPUBindingType]] - |
---|---|
[=uniform buffer=] - | [[WebGPU#dom-gpubindingtype-uniform-buffer|uniform-buffer]] - |
read-write [=storage buffer=] - | [[WebGPU#dom-gpubindingtype-storage-buffer|storage-buffer]] - |
read-only [=storage buffer=] - | [[WebGPU#dom-gpubindingtype-readonly-storage-buffer|readonly-storage-buffer]] - |
sampler - | [[WebGPU#dom-gpubindingtype-sampler|sampler]] - |
sampler_comparison - | [[WebGPU#dom-gpubindingtype-comparison-sampler|comparison-sampler]] - |
sampled texture - | [[WebGPU#dom-gpubindingtype-sampled-texture|sampled-texture]] or - [[WebGPU#dom-gpubindingtype-multisampled-texture|multisampled-texture]] - |
[=read-only storage texture=] - | [[WebGPU#dom-gpubindingtype-readonly-storage-texture|readonly-storage-texture]] - |
[=write-only storage texture=] - | [[WebGPU#dom-gpubindingtype-writeonly-storage-texture|writeonly-storage-texture]] - |
-enable_directive - : ENABLE IDENT SEMICOLON -- -Note: The grammar rule includes the terminating semicolon token, -ensuring the additional functionality is usable only after that semicolon. -Therefore any [SHORTNAME] implementation can parse the entire `enable` directive. -When an implementation encounters an enable directive for an unsupported extension, -the implementation can issue a clear diagnostic. - -
-translation_unit - : global_decl_or_directive* EOF -- -
-global_decl_or_directive - : SEMICOLON - | global_variable_decl SEMICOLON - | global_constant_decl SEMICOLON - | type_alias SEMICOLON - | struct_decl SEMICOLON - | function_decl - | enable_directive -- -# Execution TODO # {#execution} - -## Invocation of an entry point TODO ## {#invocation-of-an-entry-point} - -### Before an entry point begins TODO ### {#before-entry-point-begins} - -TODO: *Stub* - -* Setting values of builtin variables -* External-interface variables have initialized backing storage -* Internal module-scope variables have backing storage - * Initializers evaluated in textual order -* No two variables have overlapping storage (might already be covered earlier?) - -### Program order (within an invocation) TODO ### {#program-order} - -#### Function-scope variable lifetime and initialization TODO #### {#function-scope-variable-lifetime} - -#### Statement order TODO #### {#statement-order} - -#### Intra-statement order (or lack) TODO #### {#intra-statement-order} - -TODO: *Stub*: Expression evaluation - -## Uniformity TODO ## {#uniformity} - -### Uniform control flow TODO ### {#uniform-control-flow} - -### Divergence and reconvergence TODO ### {#divergence-reconvergence} - -### Uniformity restrictions TODO ### {#uniformity-restrictions} - -## Compute Shaders and Workgroups ## {#compute-shader-workgroups} - -A workgroup is a set of invocations which -concurrently execute a [=compute shader stage=] [=entry point=], -and share access to shader variables in the [=storage classes/workgroup=] storage class. - -The workgroup grid for a compute shader is the set of points -with integer coordinates *(i,j,k)* with: - -* 0 ≤ i < workgroup_size_x -* 0 ≤ j < workgroup_size_y -* 0 ≤ k < workgroup_size_z - -where *(workgroup_size_x, workgroup_size_y, workgroup_size_z)* is -the value specified for the [=workgroup_size=] attribute of the -entry point, or (1,1,1) if the entry point has no such attribute. - -There is exactly one invocation in a workgroup for each point in the workgroup grid. - -An invocation's local invocation ID is the coordinate -triple for the invocation's corresponding workgroup grid point. - -When an invocation has [=local invocation ID=] (i,j,k), then its -local invocation index is - - i + - (j * workgroup_size_x) + - (k * workgroup_size_x * workgroup_size_y) - -
Note that if a workgroup has |W| invocations, -then each invocation |I| the workgroup has a unique local invocation index |L|(|I|) -such that 0 ≤ |L|(|I|) < |W|, -and that entire range is covered.
- -A compute shader begins execution when a WebGPU implementation -removes a dispatch command from a queue and begins the specified work on the GPU. -The dispatch command specifies a dispatch size, -which is an integer triple *(group_count_x, group_count_y, group_count_z)* -indicating the number of workgroups to be executed, as described in the following. - -The compute shader grid for a particular dispatch -is the set of points with integer coordinates *(CSi,CSj,CSk)* with: - -* 0 ≤ CSi ≤ workgroup_size_x × group_count_x -* 0 ≤ CSj ≤ workgroup_size_y × group_count_y -* 0 ≤ CSk ≤ workgroup_size_z × group_count_z - -where *workgroup_size_x*, -*workgroup_size_y*, and -*workgroup_size_z* are as above for the compute shader entry point. - -The work to be performed by a compute shader dispatch is to execute exactly one -invocation of the entry point for each point in the compute shader grid. - -An invocation's global invocation ID is the coordinate -triple for the invocation's corresponding compute shader grid point. - -The invocations are organized into workgroups, so that each invocation -*(CSi, CSj, CSk)* is identified with the workgroup grid point - - ( *CSi* mod workgroup_size_x , - *CSj* mod workgroup_size_y , - *CSk* mod workgroup_size_z ) - -in workgroup ID - - ( ⌊ *CSi* ÷ workgroup_size_x ⌋, - ⌊ *CSj* ÷ workgroup_size_y ⌋, - ⌊ *CSk* ÷ workgroup_size_z ⌋). - -WebGPU provides no guarantees about: - -* Whether invocations from different workgroups execute concurrently. - That is, you cannot assume more than one workgroup executes at a time. -* Whether, once invocations from a workgroup begin executing, that other workgroups - are blocked from execution. - That is, you cannot assume that only one workgroup executes at a time. - While a workgroup is executing, the implementation may choose to - concurrently execute other workgroups as well, or other queued but unblocked work. -* Whether invocations from one particular workgroup begin executing before - the invocations of another workgroup. - That is, you cannot assume that workgroups are launched in a particular order. - -Issue: [WebGPU issue 1045](https://github.com/gpuweb/gpuweb/issues/1045): -Dispatch group counts must be positive. -However, how do we handle an indirect dispatch that specifies a group count of zero. - -## Collective operations TODO ## {#collective-operations} - -### Barrier TODO ### {#barrier} - -### Image Operations Requiring Uniformity TODO ### {#image-operations-requiring-uniformity} - -### Derivatives TODO ### {#derivatives} - -### Arrayed resource access TODO ### {#arrayed-resource-access} - -## Floating Point Evaluation TODO ## {#floating-point-evaluation} - -TODO: *Stub* - -* Infinities, NaNs, negative zeros -* Denorms, flushing -* fast-math rules: e.g. reassociation, fusing -* Invariance (or is this more general than floating point) -* Rounding -* Error bounds on basic operations - -### Floating point conversion ### {#floating-point-conversion} - -When converting a floating point scalar value to an integral type: -* If the original value is exactly representable in the destination type, then the result is that value. -* If the original value has a fractional component, then it cannot be represented exactly in the destination type, and the result is TODO -* If the original value is out of range of the destination type, then TODO. - -When converting a value to a floating point type: -* If the original value is exactly representable in the destination type, then the result is that value. - * If the original value is zero and of integral type, then the resulting value has a zero sign bit. -* Otherwise, the original value is not exactly representable. - * If the original value is different from but lies between two adjacent values representable in the destination type, - then the result is one of those two values. - [SHORTNAME] does not specify whether the larger or smaller representable - value is chosen, and different instances of such a conversion may choose differently. - * Otherwise, if the original value lies outside the range of the destination type. - * This does not occur when the original types is one of [=i32=] or [=u32=] and the destination type is [=f32=]. - * This does not occur when the source type is a floating point type with fewer exponent and mantissa bits. - * If the source type is a floating point type with more mantissa bits than the destination type, then: - * The extra mantissa bits of the source value may be discarded (treated as if they are 0). - * If the resulting value is the maximum normal value of the destination type, then that is the result. - * Otherwise the result is the infinity value with the same sign as the source value. - * Otherwise, if the original value is a NaN for the source type, then the result is a NaN in the destination type. - -NOTE: An integer value may lie between two adjacent representable floating point values. -In particular, the [=f32=] type uses 23 explicit fractional bits. -Additionally, when the floating point value is in the normal range (the exponent is neither extreme value), then the mantissa is -the set of fractional bits together with an extra 1-bit at the most significant position at bit position 23. -Then, for example, integers 228 and 1+228 both map to the same floating point value: the difference in the -least significant 1 bit is not representable by the floating point format. -This kind of collision occurs for pairs of adjacent integers with a magnitude of at least 225. - -Issue: (dneto) Default rounding mode is an implementation choice. Is that what we want? - -Issue: Check behaviour of the f32 to f16 conversion for numbers just beyond the max normal f16 values. -I've written what an NVIDIA GPU does. See https://github.com/google/amber/pull/918 for an executable test case. - -# Memory Model TODO # {#memory-model} - -# Keyword and Token Summary # {#grammar} - -## Keyword Summary ## {#keyword-summary} - -Token | Definition - |
---|---|
`ARRAY` | array - |
`BOOL` | bool - |
`FLOAT32` | f32 - |
`INT32` | i32 - |
`MAT2x2` | mat2x2 // 2 column x 2 row - |
`MAT2x3` | mat2x3 // 2 column x 3 row - |
`MAT2x4` | mat2x4 // 2 column x 4 row - |
`MAT3x2` | mat3x2 // 3 column x 2 row - |
`MAT3x3` | mat3x3 // 3 column x 3 row - |
`MAT3x4` | mat3x4 // 3 column x 4 row - |
`MAT4x2` | mat4x2 // 4 column x 2 row - |
`MAT4x3` | mat4x3 // 4 column x 3 row - |
`MAT4x4` | mat4x4 // 4 column x 4 row - |
`POINTER` | ptr - |
`SAMPLER` | sampler - |
`SAMPLER_COMPARISON` | sampler_comparison - |
`STRUCT` | struct - |
`TEXTURE_1D` | texture_1d - |
`TEXTURE_2D` | texture_2d - |
`TEXTURE_2D_ARRAY` | texture_2d_array - |
`TEXTURE_3D` | texture_3d - |
`TEXTURE_CUBE` | texture_cube - |
`TEXTURE_CUBE_ARRAY` | texture_cube_array - |
`TEXTURE_MULTISAMPLED_2D` | texture_multisampled_2d - |
`TEXTURE_STORAGE_1D` | texture_storage_1d - |
`TEXTURE_STORAGE_2D` | texture_storage_2d - |
`TEXTURE_STORAGE_2D_ARRAY` | texture_storage_2d_array - |
`TEXTURE_STORAGE_3D` | texture_storage_3d - |
`TEXTURE_DEPTH_2D` | texture_depth_2d - |
`TEXTURE_DEPTH_2D_ARRAY` | texture_depth_2d_array - |
`TEXTURE_DEPTH_CUBE` | texture_depth_cube - |
`TEXTURE_DEPTH_CUBE_ARRAY` | texture_depth_cube_array - |
`UINT32` | u32 - |
`VEC2` | vec2 - |
`VEC3` | vec3 - |
`VEC4` | vec4 - |
Token | Definition - |
`BITCAST` | bitcast - |
`BLOCK` | block - |
`BREAK` | break - |
`CASE` | case - |
`CONTINUE` | continue - |
`CONTINUING` | continuing - |
`DEFAULT` | default - |
`DISCARD` | discard - |
`ELSE` | else - |
`ELSE_IF` | elseif - |
`ENABLE` | enable - |
`FALLTHROUGH` | fallthrough - |
`FALSE` | false - |
`FN` | fn - |
`FOR` | for - |
`FUNCTION` | function - |
`IF` | if - |
`LET` | let - |
`LOOP` | loop - |
`PRIVATE` | private - |
`RETURN` | return - |
`STORAGE` | storage - |
`SWITCH` | switch - |
`TRUE` | true - |
`TYPE` | type - |
`UNIFORM` | uniform - |
`VAR` | var - |
`WORKGROUP` | workgroup - |
Token | Definition - |
`R8UNORM` | r8unorm - |
`R8SNORM` | r8snorm - |
`R8UINT` | r8uint - |
`R8SINT` | r8sint - |
`R16UINT` | r16uint - |
`R16SINT` | r16sint - |
`R16FLOAT` | r16float - |
`RG8UNORM` | rg8unorm - |
`RG8SNORM` | rg8snorm - |
`RG8UINT` | rg8uint - |
`RG8SINT` | rg8sint - |
`R32UINT` | r32uint - |
`R32SINT` | r32sint - |
`R32FLOAT` | r32float - |
`RG16UINT` | rg16uint - |
`RG16SINT` | rg16sint - |
`RG16FLOAT` | rg16float - |
`RGBA8UNORM` | rgba8unorm - |
`RGBA8UNORM-SRGB` | rgba8unorm_srgb - |
`RGBA8SNORM` | rgba8snorm - |
`RGBA8UINT` | rgba8uint - |
`RGBA8SINT` | rgba8sint - |
`BGRA8UNORM` | bgra8unorm - |
`BGRA8UNORM-SRGB` | bgra8unorm_srgb - |
`RGB10A2UNORM` | rgb10a2unorm - |
`RG11B10FLOAT` | rg11b10float - |
`RG32UINT` | rg32uint - |
`RG32SINT` | rg32sint - |
`RG32FLOAT` | rg32float - |
`RGBA16UINT` | rgba16uint - |
`RGBA16SINT` | rgba16sint - |
`RGBA16FLOAT` | rgba16float - |
`RGBA32UINT` | rgba32uint - |
`RGBA32SINT` | rgba32sint - |
`RGBA32FLOAT` | rgba32float - |
asm - | bf16 - | do - | enum - | f16 - |
f64 - | i8 - | i16 - | i64 - | const - |
typedef - | u8 - | u16 - | u64 - | unless - |
using - | while - | regardless - | premerge - | handle - |
`AND` | `&` - |
`AND_AND` | `&&` - |
`ARROW` | `->` - |
`ATTR_LEFT` | `[[` - |
`ATTR_RIGHT` | `]]` - |
`FORWARD_SLASH` | `/` - |
`BANG` | `!` - |
`BRACKET_LEFT` | `[` - |
`BRACKET_RIGHT` | `]` - |
`BRACE_LEFT` | `{` - |
`BRACE_RIGHT` | `}` - |
`COLON` | `:` - |
`COMMA` | `,` - |
`EQUAL` | `=` - |
`EQUAL_EQUAL` | `==` - |
`NOT_EQUAL` | `!=` - |
`GREATER_THAN` | `>` - |
`GREATER_THAN_EQUAL` | `>=` - |
`SHIFT_RIGHT` | `>>` - |
`LESS_THAN` | `<` - |
`LESS_THAN_EQUAL` | `<=` - |
`SHIFT_LEFT` | `<<` - |
`MODULO` | `%` - |
`MINUS` | `-` - |
`MINUS_MINUS` | `--` - |
`PERIOD` | `.` - |
`PLUS` | `+` - |
`PLUS_PLUS` | `++` - |
`OR` | `|` - |
`OR_OR` | `||` - |
`PAREN_LEFT` | `(` - |
`PAREN_RIGHT` | `)` - |
`SEMICOLON` | `;` - |
`STAR` | `*` - |
`TILDE` | `~` - |
`XOR` | `^` - |
Built-in | Stage | Input or Output | Store type | Description - |
---|---|---|---|---|
`vertex_index` - | vertex - | in - | u32 - | Index of the current vertex within the current API-level draw command, - independent of draw instancing. - - For a non-indexed draw, the first vertex has an index equal to the `firstIndex` argument - of the draw, whether provided directly or indirectly. - The index is incremented by one for each additional vertex in the draw instance. - - For an indexed draw, the index is equal to the index buffer entry for - vertex, plus the `baseVertex` argument of the draw, whether provided directly or indirectly. - - |
`instance_index` - | vertex - | in - | u32 - | Instance index of the current vertex within the current API-level draw command. - - The first instance has an index equal to the `firstInstance` argument of the draw, - whether provided directly or indirectly. - The index is incremented by one for each additional instance in the draw. - - |
`position` - | vertex - | out - | vec4<f32> - | Output position of the current vertex, using homogeneous coordinates. - After homogeneous normalization (where each of the *x*, *y*, and *z* components - are divided by the *w* component), the position is in the WebGPU normalized device - coordinate space. - See [[WebGPU#coordinate-systems|WebGPU § Coordinate Systems]]. - - |
`position` - | fragment - | in - | vec4<f32> - | Framebuffer position of the current fragment, using normalized homogeneous - coordinates. - (The *x*, *y*, and *z* components have already been scaled such that *w* is now 1.) - See [[WebGPU#coordinate-systems|WebGPU § Coordinate Systems]]. - - |
`front_facing` - | fragment - | in - | bool - | True when the current fragment is on a front-facing primitive. - False otherwise. - See [[WebGPU#dom-gpurasterizationstatedescriptor-frontface|WebGPU § Rasterization State]]. - - |
`frag_depth` - | fragment - | out - | f32 - | Updated depth of the fragment, in the viewport depth range. - See [[WebGPU#coordinate-systems|WebGPU § Coordinate Systems]]. - - |
`local_invocation_id` - | compute - | in - | vec3<u32> - | The current invocation's [=local invocation ID=], - i.e. its position in the [=workgroup grid=]. - - |
`local_invocation_index` - | compute - | in - | u32 - | The current invocation's [=local invocation index=], a linearized index of - the invocation's position within the [=workgroup grid=]. - - |
`global_invocation_id` - | compute - | in - | vec3<u32> - | The current invocation's [=global invocation ID=], - i.e. its position in the [=compute shader grid=]. - - |
`workgroup_id` - | compute - | in - | vec3<u32> - | The current invocation's [=workgroup ID=], - i.e. the position of the workgroup in the [=workgroup grid=]. - - |
`workgroup_size` - | compute - | in - | vec3<u32> - | The [=workgroup_size=] of the current entry point. - - |
`sample_index` - | fragment - | in - | u32 - | Sample index for the current fragment.
- The value is least 0 and at most `sampleCount`-1, where
- [[WebGPU#dom-gpurenderpipelinedescriptor-samplecount|sampleCount]]
- is the number of MSAA samples specified for the GPU render pipeline.
- See [[WebGPU#gpurenderpipe|WebGPU § GPURenderPipeline]]. - - |
`sample_mask` - | fragment - | in - | u32 - | Sample coverage mask for the current fragment.
- It contains a bitmask indicating which samples in this fragment are covered
- by the primitive being rendered.
- See [[WebGPU#sample-masking|WebGPU § Sample Masking]]. - - |
`sample_mask` - | fragment - | out - | u32 - | Sample coverage mask control for the current fragment.
- The last value written to this variable becomes the
- [[WebGPU#shader-output-mask|shader-output mask]].
- Zero bits in the written value will cause corresponding samples in
- the color attachments to be discarded.
- See [[WebGPU#sample-masking|WebGPU § Sample Masking]]. - |
Logical built-in functions | SPIR-V - |
---|---|
all(BoolVec) -> bool | OpAll - |
any(BoolVec) -> bool | OpAny - |
select(*T*,*T*,bool) -> *T* |
- For scalar or vector type *T*.
- `select(a,b,c)` evaluates to *a* when *c* is true, and *b* otherwise. - OpSelect - |
select(vec*N*<*T*>,vec*N*<*T*>,vec*N*<bool>) -> vec*N*<*T*> |
- For scalar type *T*.
- `select(a,b,c)` evaluates to a vector with component *i* being `select(a[i], b[i], c[i])`. - OpSelect - |
Precondition | Conclusion | Notes - |
---|---|---|
|e| : f32 | `isNan(`|e|`)` : bool - | Returns true if |e| is NaN according to IEEE. (OpIsNan) - |
|e| : |T|, |T| is *FloatVec* - | `isNan(`|e|`)` : vec|N|<bool>, where |N| = *Arity(*|T|*)* | Component-wise test for NaN. Component *i* of the result is *isNan(e[i])*. (OpIsNan) - |
|e| : f32 | `isInf(`|e|`)` : bool - | Returns true if |e| is an infinity according to IEEE. (OpIsInf) - |
|e| : |T|, |T| is *FloatVec* - | `isInf(`|e|`)` : vec|N|<bool>, where |N| = *Arity(*|T|*)* | Component-wise test for inifinity. Component *i* of the result is *isInf(e[i])*. (OpIsInf) - |
|e| : f32 | `isFinite(`|e|`)` : bool - | Returns true if |e| is finite according to IEEE. (emulated) - |
|e| : |T|, |T| is *FloatVec* - | `isFinite(`|e|`)` : vec|N|<bool>, where |N| = *Arity(*|T|*)* | Component-wise finite value test. Component *i* of the result is *isFinite(e[i])*. (emulated) - |
|e| : f32 | `isNormal(`|e|`)` : bool - | Returns true if |e| is a normal number according to IEEE. (emulated) - |
|e| : |T|, |T| is *FloatVec* - | `isNormal(`|e|`)` : vec|N|<bool>, where |N| = *Arity(*|T|*)* | Component-wise test for normal number. Component *i* of the result is *isNormal(e[i])*. (emulated) - |
|e| : ptr<storage,array<|T|>> - | `arrayLength(`|e|`)` : u32 | Returns the number of elements in the runtime array. - (OpArrayLength, but you have to trace back to get the pointer to the enclosing struct.) - |
Precondition | Built-in | Description - |
---|---|---|
|T| is f32 - | `abs(`|e|`:` |T| `) -> ` |T| - | (GLSLstd450FAbs) - |
|T| is f32 - | `abs(`|e|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | (GLSLstd450FAbs) - |
|T| is f32 - | `acos(`|e|`:` |T| `) -> ` |T| - | (GLSLstd450Acos) - |
|T| is f32 - | `acos(`|e|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | (GLSLstd450Acos) - |
|T| is f32 - | `asin(`|e|`:` |T| `) -> ` |T| - | (GLSLstd450Asin) - |
|T| is f32 - | `asin(`|e|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | (GLSLstd450Asin) - |
|T| is f32 - | `atan(`|e|`:` |T| `) -> ` |T| - | (GLSLstd450Atan) - |
|T| is f32 - | `atan(`|e|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | (GLSLstd450Atan) - |
|T| is f32 - | `atan2(`|e1|`:` |T| `, `|e2|`:` |T| `) -> ` |T| - | (GLSLstd450Atan2) - |
|T| is f32 - | `atan2(`|e1|`:` vec|N|<|T|> `, `|e2|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | (GLSLstd450Atan2) - |
|T| is f32 - | `ceil(`|e|`:` |T| `) -> ` |T| - | (GLSLstd450Ceil) - |
|T| is f32 - | `ceil(`|e|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | (GLSLstd450Ceil) - |
|T| is f32 - | `clamp(`|e1|`:` |T| `, `|e2|`:` |T| `, `|e3|`:` |T|`) -> ` |T| - | (GLSLstd450NClamp) - |
|T| is f32 - | `clamp(`|e1|`:` vec|N|<|T|> `, `|e2|`:` vec|N|<|T|>`, `|e3|`:` vec|N|<|T|>`) -> ` vec|N|<|T|> - | (GLSLstd450NClamp) - |
|T| is f32 - | `cos(`|e|`:` |T| `) -> ` |T| - | (GLSLstd450Cos) - |
|T| is f32 - | `cos(`|e|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | (GLSLstd450Cos) - |
|T| is f32 - | `cosh(`|e|`:` |T| `) -> ` |T| - | (GLSLstd450Cosh) - |
|T| is f32 - | `cosh(`|e|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | (GLSLstd450Cosh) - |
|T| is f32 - | `cross(`|e1|`:` vec3<|T|> `, `|e2|`:` vec3<|T|>`) -> ` vec3<|T|> - | (GLSLstd450Cross) - |
|T| is f32 - | `distance(`|e1|`:` |T| `, `|e2|`:` |T| `) -> ` |T| - | (GLSLstd450Distance) - |
|T| is f32 - | `distance(`|e1|`:` vec|N|<|T|> `, `|e2|`:` vec|N|<|T|>`) -> ` |T| - | (GLSLstd450Distance) - |
|T| is f32 - | `exp(`|e|`:` |T| `) -> ` |T| - | (GLSLstd450Exp) - |
|T| is f32 - | `exp(`|e|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | (GLSLstd450Exp) - |
|T| is f32 - | `exp2(`|e|`:` |T| `) -> ` |T| - | (GLSLstd450Exp2) - |
|T| is f32 - | `exp2(`|e|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | (GLSLstd450Exp2) - |
|T| is f32 - | `faceForward(`|e1|`:` |T| `, `|e2|`:` |T| `, `|e3|`:` |T| `) -> ` |T| - | (GLSLstd450FaceForward) - |
|T| is f32 - | `faceForward(`|e1|`:` vec|N|<|T|> `, `|e2|`:` vec|N|<|T|>`, `|e3|`:` vec|N|<|T|>`) -> ` vec|N|<|T|> - | (GLSLstd450FaceForward) - |
|T| is f32 - | `floor(`|e|`:` |T| `) -> ` |T| - | (GLSLstd450Floor) - |
|T| is f32 - | `floor(`|e|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | (GLSLstd450Floor) - |
|T| is f32 - | `fma(`|e1|`:` |T| `, `|e2|`:` |T| `, `|e3|`:` |T| `) -> ` |T| - | (GLSLstd450Fma) - |
|T| is f32 - | `fma(`|e1|`:` vec|N|<|T|> `, `|e2|`:` vec|N|<|T|>`, `|e3|`:` vec|N|<|T|>`) -> ` vec|N|<|T|> - | (GLSLstd450Fma) - |
|T| is f32 - | `fract(`|e|`:` |T| `) -> ` |T| - | (GLSLstd450Fract) - |
|T| is f32 - | `fract(`|e|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | (GLSLstd450Fract) - |
|T| is f32 - |I| is i32 or u32 - | `frexp(`|e1|`:` |T| `, `|e2|`:` ptr<|I|> `) -> ` |T| - | (GLSLstd450Frexp) - |
|T| is f32 - |I| is i32 or u32 - | `frexp(`|e1|`:` vec|N|<|T|> `, `|e2|`:` ptr<vec|N|<|I|>>`) -> ` vec|N|<|T|> - | (GLSLstd450Frexp) - |
|T| is f32 - | `inverseSqrt(`|e|`:` |T| `) -> ` |T| - | (GLSLstd450InverseSqrt) - |
|T| is f32 - | `inverseSqrt(`|e|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | (GLSLstd450InverseSqrt) - |
|T| is f32 - |I| is i32 or u32 - | `ldexp(`|e1|`:` |T| `, `|e2|`:` |I| `) -> ` |T| - | (GLSLstd450Ldexp) - |
|T| is f32 - |I| is i32 or u32 - | `ldexp(`|e1|`:` vec|N|<|T|> `, `|e2|`:` vec|N|<|I|>`) -> ` vec|N|<|T|> - | (GLSLstd450Ldexp) - |
|T| is f32 - | `length(`|e|`:` |T| `) -> ` |T| - | (GLSLstd450Length) - |
|T| is f32 - | `length(`|e|`:` vec|N|<|T|> `) -> ` |T| - | (GLSLstd450Length) - |
|T| is f32 - | `log(`|e|`:` |T| `) -> ` |T| - | (GLSLstd450Log) - |
|T| is f32 - | `log(`|e|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | (GLSLstd450Log) - |
|T| is f32 - | `log2(`|e|`:` |T| `) -> ` |T| - | (GLSLstd450Log2) - |
|T| is f32 - | `log2(`|e|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | (GLSLstd450Log2) - |
|T| is f32 - | `max(`|e1|`:` |T| `, `|e2|`:` |T| `) -> ` |T| - | (GLSLstd450NMax) - |
|T| is f32 - | `max(`|e1|`:` vec|N|<|T|> `, `|e2|`:` vec|N|<|T|>`) -> ` vec|N|<|T|> - | (GLSLstd450NMax) - |
|T| is f32 - | `min(`|e1|`:` |T| `, `|e2|`:` |T| `) -> ` |T| - | (GLSLstd450NMin) - |
|T| is f32 - | `min(`|e1|`:` vec|N|<|T|> `, `|e2|`:` vec|N|<|T|>`) -> ` vec|N|<|T|> - | (GLSLstd450NMin) - |
|T| is f32 - | `mix(`|e1|`:` |T| `, `|e2|`:` |T| `, `|e3|`:` |T|`) -> ` |T| - | (GLSLstd450FMix) - |
|T| is f32 - | `mix(`|e1|`:` vec|N|<|T|> `, `|e2|`:` vec|N|<|T|>`, `|e3|`:` vec|N|<|T|>`) -> ` vec|N|<|T|> - | (GLSLstd450FMix) - |
|T| is f32 - | `modf(`|e1|`:` |T| `, `|e2|`:` ptr<|T|> `) -> ` |T| - | (GLSLstd450Modf) - |
|T| is f32 - | `modf(`|e1|`:` vec|N|<|T|> `, `|e2|`:` ptr<vec|N|<|T|>>`) -> ` vec|N|<|T|> - | (GLSLstd450Modf) - |
|T| is f32 - | `normalize(`|e|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | (GLSLstd450Normalize) - |
|T| is f32 - | `pow(`|e1|`:` |T| `, `|e2|`:` |T| `) -> ` |T| - | (GLSLstd450Pow) - |
|T| is f32 - | `pow(`|e1|`:` vec|N|<|T|> `, `|e2|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | (GLSLstd450Pow) - |
|T| is f32 - | `reflect(`|e1|`:` |T| `, `|e2|`:` |T| `) -> ` |T| - | (GLSLstd450Reflect) - |
|T| is f32 - | `reflect(`|e1|`:` vec|N|<|T|> `, `|e2|`:` vec|N|<|T|>`) -> ` vec|N|<|T|> - | (GLSLstd450Reflect) - |
|T| is f32 - | `round(`|e|`:` |T| `) -> ` |T| - | Result is the integer |k| nearest to |e|, as a floating point value. - When |e| lies halfway between integers |k| and |k|+1, - the result is |k| when |k| is even, and |k|+1 when |k| is odd. - (GLSLstd450RoundEven) - |
|T| is f32 - | `round(`|e|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | Component-wise rounding. - Component |i| of the result is `round`(|e|[|i|]) - (GLSLstd450RoundEven) - |
|T| is f32 - | `sign(`|e|`:` |T| `) -> ` |T| - | (GLSLstd450FSign) - |
|T| is f32 - | `sign(`|e|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | (GLSLstd450FSign) - |
|T| is f32 - | `sin(`|e|`:` |T| `) -> ` |T| - | (GLSLstd450Sin) - |
|T| is f32 - | `sin(`|e|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | (GLSLstd450Sin) - |
|T| is f32 - | `sinh(`|e|`:` |T| `) -> ` |T| - | (GLSLstd450Sinh) - |
|T| is f32 - | `sinh(`|e|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | (GLSLstd450Sinh) - |
|T| is f32 - | `smoothStep(`|e1|`:` |T| `, `|e2|`:` |T| `, `|e3|`:` |T| `) -> ` |T| - | (GLSLstd450SmoothStep) - |
|T| is f32 - | `smoothStep(`|e1|`:` vec|N|<|T|> `, `|e2|`:` vec|N|<|T|>`, `|e3|`:` vec|N|<|T|>`) -> ` vec|N|<|T|> - | (GLSLstd450SmoothStep) - |
|T| is f32 - | `sqrt(`|e|`:` |T| `) -> ` |T| - | (GLSLstd450Sqrt) - |
|T| is f32 - | `sqrt(`|e|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | (GLSLstd450Sqrt) - |
|T| is f32 - | `step(`|e1|`:` |T| `, `|e2|`:` |T| `) -> ` |T| - | (GLSLstd450Step) - |
|T| is f32 - | `step(`|e1|`:` vec|N|<|T|> `, `|e2|`:` vec|N|<|T|>`) -> ` vec|N|<|T|> - | (GLSLstd450Step) - |
|T| is f32 - | `tan(`|e|`:` |T| `) -> ` |T| - | (GLSLstd450Tan) - |
|T| is f32 - | `tan(`|e|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | (GLSLstd450Tan) - |
|T| is f32 - | `tanh(`|e|`:` |T| `) -> ` |T| - | (GLSLstd450Tanh) - |
|T| is f32 - | `tanh(`|e|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | (GLSLstd450Tanh) - |
|T| is f32 - | `trunc(`|e|`:` |T| `) -> ` |T| - | (GLSLstd450Trunc) - |
|T| is f32 - | `trunc(`|e|`:` vec|N|<|T|> `) -> ` vec|N|<|T|> - | (GLSLstd450Trunc) - |
Precondition | Built-in | Description - |
---|---|---|
- | `abs`(|e|: i32 ) -> i32 - | The absolute value of |e|. - (GLSLstd450SAbs) - |
- | `abs`(|e| : vec|N|<i32> ) -> vec|N|<i32> - | Component-wise absolute value:
- Component |i| of the result is `abs(`|e|`[`|i|`])` - (GLSLstd450SAbs) - |
- | `abs`(|e| : u32 ) -> u32 - | Result is |e|. This is provided for symmetry with `abs` for signed integers. - |
- | `abs(`|e|`:` vec|N|<u32> `) ->` vec|N|<u32> - | Result is |e|. This is provided for symmetry with `abs` for signed integer vectors. - |
|T| is u32 - | `clamp(`|e1|`:` |T| `, `|e2|`:` |T|`, `|e3|`:` |T|`) ->` |T| - | (GLSLstd450UClamp) - |
|T| is u32 - | `clamp(`|e1|`:` vec|N|<|T|> `, `|e2|`:` vec|N|<|T|>`, `|e3|`:`vec|N|<|T|> `) ->` vec|N|<|T|> - | (GLSLstd450UClamp) - |
|T| is i32 - | `clamp(`|e1|`:` |T| `, `|e2|`:` |T|`, `|e3|`:` |T|`) ->` |T| - | (GLSLstd450SClamp) - |
|T| is i32 - | `clamp(`|e1|`:` vec|N|<|T|> `, `|e2|`:` vec|N|<|T|>`, `|e3|`:`vec|N|<|T|> `) ->` vec|N|<|T|> - | (GLSLstd450SClamp) - |
|T| is u32 or i32 - | `countOneBits(`|e|`:` |T| `) ->` |T| - | The number of 1 bits in the representation of |e|. - Also known as "population count". - (SPIR-V OpBitCount) - |
|T| is u32 or i32 - | `countOneBits(`|e|`:` vec|N|<|T|>`) ->` vec|N|<|T|> - | Component-wise population count:
- Component |i| of the result is `countOneBits(`|e|`[`|i|`])` - (SPIR-V OpBitCount) - |
|T| is u32 - | `max(`|e1|`:` |T| `, `|e2|`:` |T|`) ->` |T| - | (GLSLstd450UMax) - |
|T| is u32 - | `max(`|e1|`:` vec|N|<|T|> `, `|e2|`:` vec|N|<|T|>`) ->` vec|N|<|T|> - | (GLSLstd450UMax) - |
|T| is i32 - | `max(`|e1|`:` |T| `, `|e2|`:` |T|`) ->` |T| - | (GLSLstd450SMax) - |
|T| is i32 - | `max(`|e1|`:` vec|N|<|T|> `, `|e2|`:` vec|N|<|T|>`) ->` vec|N|<|T|> - | (GLSLstd450SMax) - |
|T| is u32 - | `min(`|e1|`:` |T| `, `|e2|`:` |T|`) ->` |T| - | (GLSLstd450UMin) - |
|T| is u32 - | `min(`|e1|`:` vec|N|<|T|> `, `|e2|`:` vec|N|<|T|>`) ->` vec|N|<|T|> - | (GLSLstd450UMin) - |
|T| is i32 - | `min(`|e1|`:` |T| `, `|e2|`:` |T|`) ->` |T| - | (GLSLstd450SMin) - |
|T| is i32 - | `min(`|e1|`:` vec|N|<|T|> `, `|e2|`:` vec|N|<|T|>`) ->` vec|N|<|T|> - | (GLSLstd450SMin) - |
|T| is u32 or i32 - | `reverseBits(`|e|`:` |T| `) ->` |T| - | Reverses the bits in |e|: The bit at position |k| of the result equals the
- bit at position 31-|k| of |e|. - (SPIR-V OpBitReverse) - |
|T| is u32 or i32 - | `reverseBits(`|e|`:` vec|N|<|T|> `) ->` vec|N|<|T|> - | Component-wise bit reversal:
- Component |i| of the result is `reverseBits(`|e|`[`|i|`])` - (SPIR-V OpBitReverse) - |
Precondition | Built-in | Description - |
---|---|---|
|T| is f32 - | `determinant(`|e|`:` mat|N|x|N|<|T|> `) -> ` |T| - | (GLSLstd450Determinant) - |
Vector built-in functions | SPIR-V - |
---|---|
dot(vecN<f32>, vecN<f32>) -> float | OpDot - |
Precondition | Derivative built-in functions | SPIR-V - |
---|---|---|
|T| is f32 or vecN<f32> | dpdx(T) -> T | OpDPdx - |
dpdxCoarse(T) -> T | OpDPdxCoarse - | |
dpdxFine(T) -> T | OpDPdxFine - | |
dpdy(T) -> T | OpDPdy - | |
dpdyCoarse(T) -> T | OpDPdyCoarse - | |
dpdyFine(T) -> T | OpDPdyFine - | |
fwidth(T) -> T | OpFwidth - | |
fwidthCoarse(T) -> T | OpFwidthCoarse - | |
fwidthFine(T) -> T | OpFwidthFine - |
`t` | - The [sampled](#sampled-texture-type), - [multisampled](#multisampled-texture-type), [depth](#texture-depth), or - [storage](#texture-storage) texture. - |
`level` |
- The mip level, with level 0 containing a full size version of the texture. - If omitted, the dimensions of level 0 are returned. - |
`t` | - The [sampled](#sampled-texture-type), - [multisampled](#multisampled-texture-type), [depth](#texture-depth) or - [read-only storage](#texture-storage) texture. - |
`coords` | - The 0-based texel coordinate. - |
`array_index` | - The 0-based texture array index. - |
`level` | - The mip level, with level 0 containing a full size version of the texture. - |
`sample_index` | - The 0-based sample index of the multisampled texture. - |
`t` | - The [sampled](#sampled-texture-type), - [multisampled](#multisampled-texture-type), [depth](#texture-depth) or - [storage](#texture-storage) array texture. - |
`t` | - The [sampled](#sampled-texture-type) or [depth](#texture-depth) texture. - |
`t` | - The [multisampled](#multisampled-texture-type) texture. - |
`t` | - The [sampled](#sampled-texture-type) or [depth](#texture-depth) texture to - sample. - |
`s` | - The [sampler type](#sampler-type). - |
`coords` | - The texture coordinates used for sampling. - |
`array_index` | - The 0-based texture array index to sample. - |
`offset` |
- The optional texel offset applied to the unnormalized texture coordinate
- before sampling the texture. This offset is applied before applying any
- texture wrapping modes. - `offset` must be compile time constant, and may only be provided as a - [literal](#literals) or `const_expr` expression (e.g. `vec2 - Each `offset` component must be at least `-8` and at most `7`. Values outside - of this range will be treated as a compile time error. - |