WGSL 内存和内存视图

Night_Aurora

已于 2023-02-06 11:43:19 修改

阅读量924

点赞数

分类专栏： WebGPU 文章标签： webgpu wgsl

于 2023-02-03 17:58:07 首次发布

本文链接：https://blog.csdn.net/fatcat123/article/details/128873701

版权

WebGPU 专栏收录该内容

11 篇文章 2 订阅

订阅专栏

在 WGSL 中，可存储类型的值被存储到内存中，以便后续检索。本节讲述内存结构和 WGSL变量如何在内存中保存内容。

内存 Memory

内存位置 Memory Locations

内存是由一些列不同的memory location组成, 每个 memory location 8 比特大小。内存操作会影响一个或者多个 memory loacation。在结构体或数组上的内存操作不会访问内存填充区域。如果两组内存位置的交集是非空的，则它们是重叠的。

内存访问模式 Memory Access Mode

内存访问是发生在 memory location 上操作，有三种模式：

（1）读模式 read access

（2）写模式 write access

（3）读写模式 read_ write access

特定的内存位置可能只支持某些类型的访问，用内存的访问模式表示。

可存储类型 Storable Types

可存储类型有以下几种：

a scalar type
a vector type
a matrix type
an atomic type
an array type
a structure type
a texture type
a sampler type

也就是说，可存储类型是具体普通类型，纹理类型，采样器类型。

主机可共享类型 Host-shareable Types

主机可共享类型用于描述缓冲区的内容，这些内容在主机和GPU之间共享，或者在主机和GPU之间无需格式转换就可以复制。当用于此目的时，类型可能还具有内存布局属性。

主机可共享类型有以下几种：

a numeric scalar type 数值型标量
a numeric vector type 数值型向量
a matrix type 矩阵类型
an atomic type 原子类型
a fixed-size array type, if it has creation-fixed footprint and its element type is host-shareable
a runtime-sized array type, if its element type is host-shareable
a structure type, if all its members are host-shareable

固定大小的数组和非固定大小的数组，只要里面的元素是主机可共享类型，它也是主机可共享的。

许多类型可以在主机上共享，但不能在io上共享，包括原子类型、运行时大小的数组，以及包含它们的任何复合类型。

地址空间 Address Spaces

内存位置（memory location）被划分到不同的地址空间(address spaces)中，每个地址空间（address space）都有唯一的属性(unique properities)，用于确定可变性、可见性、可能包含的值以及如何在其中使用变量。

(1) function 函数地址空间

(2) private 私有地址空间

(3) workgroup 工作组地址空间

(4) uniform 统一变量缓冲地址空间, 相同的着色器阶段

(5) storage 存储环控地址空间，相同的着色器阶段

(6) handle 句柄地址空间 texture and sampler 相同的着色器阶段

内存布局 Memory Layout

Uniform buffer和storage buffer变量用于共享在内存中组织为字节序列的批量数据。缓冲区在CPU和GPU之间共享，或者在流水线中的不同着色阶段之间共享，或者在不同流水线之间共享。

因为缓冲区数据是在没有重新格式化或转换的情况下共享的，如果缓冲区的生产者和消费者在内存布局(即缓冲区中的字节如何组织成类型化的WGSL值的描述)上没有达成一致，就会出现动态错误。

8位字节是主机共享内存的最基本单位。本节定义的术语表示8位字节的计数。

内存视图 Memory Views

引用和指针类型 Reference and Pointer Types

WGSL有两种类型来表示内存视图:引用类型(reference types)和指针类型(pointer types)。

AS ： address space;

T ： storable type

AM : access mode

引用类型的表示：

ref<AS, T, AM> ， 引用类型不用来写 WGSL 程序代码，它只是用来分析 WGSL 程序。

指针类型的表示：

ptr<AS, T, AM> ， 指针类型被用来写 WGSL 程序。

如果两个指针类型的地址空间，存储类型和访问模式一致，我们就说它们是一样的。

指针类型示例代码：

fn my_function(
  /* 'ptr<function,i32,read_write>' is the type of a pointer value that references
     memory for keeping an 'i32' value, using memory locations in the 'function'
     address space.  Here 'i32' is the store type.
     The implied access mode is 'read_write'. See below for access mode defaults. */
  ptr_int: ptr<function,i32>,

  // 'ptr<private,array<f32,50>,read_write>' is the type of a pointer value that
  // refers to memory for keeping an array of 50 elements of type 'f32', using
  // memory locations in the 'private' address space.
  // Here the store type is 'array<f32,50>'.
  // The implied access mode is 'read_write'. See below for access mode defaults.
  ptr_array: ptr<private, array<f32, 50>>
) { }

指针类型默认访问模式是 read_write !

ptr<AS,T,AM>类型的每个指针值p对应于ref<AS,T,AM>类型的唯一引用值r，反之亦然，其中p和r描述的是同一个内存视图。

默认内存访问模式 Access Mode Defaults

内存视图的默认访问模式是上下文相关的， storage address spaces 既支持读和也支持写模式。

Access Mode Defaults.

Address Space	Default Access Mode
function	read_write
private	read_write
workgroup	read_write
uniform	read
storage	read
handle	read

起始变量 Originating Variable

在WGSL中，引用值 总是对应于某个变量的某些或所有内存位置（memory locations）的 内存视图（memory view），这定义了 引用值（reference value）的 起始变量。

指针值（pointer value）总是对应于引用值（reference value）, 所以指针值的起始值与引用值的起始值是相同的。

注意：起始变量（originating variable）是一个动态的概念，函数形参的起始变量依赖于函数的调用方，不同的调用方会指向不同的起始变量（Originalting variable）。

一句话：指针值和引用值关联到相同的内存视图，它们是等效的。

无效的内存引用 Invalid Memory Reference

使用指针或引用访问越界会产生一个无效的内存引用错误。从无效的引用被加载时会产生下列值：

当原始变量是统一缓冲区或存储缓冲区时，绑定到原始变量的WebGPU GPUBuffer的任何内存位置的值；
当原始变量不是统一缓冲区或存储缓冲区时，原始变量中任何存储位置的值；
引用的存储类型的零值；

存储到无效的内存引用行为：

当原始变量是统一缓冲区或存储缓冲区时，将值存储到绑定到原始变量的WebGPU GPUBuffer的任何内存位置；
当原始变量不是统一缓冲区或存储缓冲区时，将值存储到原始变量中的任何存储位置；
不执行；

如果读-修改-写原子操作在无效的内存引用加载上，并且在访问内存时从不同的内存位置进行存储，则这是一个动态错误。

引用和指针用例 Use Cases for References and Pointers

引用（references）和指针(pointers) 是按使用方式来区分的：

变量的类型是引用类型 reference type; （理所当然的事）
使用一元取址操作符（&）将引用值转换为指针值；（类似于 c/c++）
使用一元解引用操作符（*）将指针转换为引用值； (类似于 c/c++)
let 声明类型可以是指针类型，但不能是引用类型；
形参可以是指针类型，但不能是引用类型；
简单的赋值语句可以通过引用可以更新内存中的内容；
Load Rule: 在函数内部，引用会自动解引用以满足类型规则:

引用简化了变量的使用：

@compute @workgroup_size(1)
fn main() {
  // 'i' has reference type ref<function,i32,read_write>
  // The memory locations for 'i' store the i32 value 0.
  var i: i32 = 0;

  // 'i + 1' can only match a type rule where the 'i' subexpression is of type i32.
  // So the expression 'i + 1' has type i32, and at evaluation, the 'i' subexpression
  // evaluates to the i32 value stored in the memory locations for 'i' at the time
  // of evaluation.
  let one: i32 = i + 1;

  // Update the value in the locations referenced by 'i' so they hold the value 2.
  i = one + 1;

  // Update the value in the locations referenced by 'i' so they hold the value 5.
  // The evaluation of the right-hand-side occurs before the assignment takes effect.
  i = i + 3;
}

var<private> age: i32;
fn get_age() -> i32 {
  // The type of the expression in the return statement must be 'i32' since it
  // must match the declared return type of the function.
  // The 'age' expression is of type ref<private,i32,read_write>.
  // Apply the Load Rule, since the store type of the reference matches the
  // required type of the expression, and no other type rule applies.
  // The evaluation of 'age' in this context is the i32 value loaded from the
  // memory locations referenced by 'age' at the time the return statement is
  // executed.
  return age;
}

fn caller() {
  age = 21;
  // The copy_age constant will get the i32 value 21.
  let copy_age: i32 = get_age();
}

指针的用方法有两个关键用例：

（1）用 let 声明指针类型，作为变量部分内容的简短名称：

struct Particle {
  position: vec3<f32>,
  velocity: vec3<f32>
}
struct System {
  active_index: i32,
  timestep: f32,
  particles: array<Particle,100>
}
@group(0) @binding(0) var<storage,read_write> system: System;

@compute @workgroup_size(1)
fn main() {
  // Form a pointer to a specific Particle in storage memory.
  let active_particle: ptr<storage,Particle> =
      &system.particles[system.active_index];

  let delta_position: vec3<f32> = (*active_particle).velocity * system.timestep;
  let current_position: vec3<f32>  = (*active_particle).position;
  (*active_particle).position = delta_position + current_position;
}

（2）作为函数的形参访问和修改内存值：

fn add_one(x: ptr<function,i32>) {
  /* Update the locations for 'x' to contain the next higher integer value,
     (or to wrap around to the largest negative i32 value).
     On the left-hand side, unary '*' converts the pointer to a reference that
     can then be assigned to. It has a read_write access mode, by default.
     /* On the right-hand side:
        - Unary '*' converts the pointer to a reference, with a read_write
          access mode.
        - The only matching type rule is for addition (+) and requires '*x' to
          have type i32, which is the store type for '*x'.  So the Load Rule
          applies and '*x' evaluates to the value stored in the memory for '*x'
          at the time of evaluation, which is the i32 value for 0.
        - Add 1 to 0, to produce a final value of 1 for the right-hand side. */
     Store 1 into the memory for '*x'. */
  *x = *x + 1;
}

@compute @workgroup_size(1)
fn main() {
  var i: i32 = 0;

  // Modify the contents of 'i' so it will contain 1.
  // Use unary '&' to get a pointer value for 'i'.
  // This is a clear signal that the called function has access to the memory
  // for 'i', and may modify it.
  add_one(&i);
  let one: i32 = i;  // 'one' has value 1.
}

形成引用和指针值 Forming Reference and Pointer Values

形成引用的方法有：

（1）标识符解析在作用域内解析为变量内存的引用；

（2）对指针类型使用解引用标识符（*）转换成引用；

（3）使用复合引用组件表达式（composite reference component expression）:

在每种情况下，结果的起始变量被定义为原始参考的起始变量。

给定一个向量存储类型的引用，附加一个单个字母的向量访问短语会得到一个指向向量中命名成分的引用；
给定一个向量存储类型的引用，添加数组索引访问短语会得到一个指向向量中被索引的部分的引用；

（例如 var a:vec3<f32>; a.x 或 a[0] 的形式都会返回组件引用）

给定一个矩阵存储类型的引用，添加数组索引访问短语会得到一个指向矩阵中已索引列向量的引用；(例如 mat[0] 形式）；
给定一个具有数组存储类型的引用，添加数组索引访问短语将得到一个对被索引的数组元素的引用；
给定具有结构存储类型的引用，添加成员访问短语将得到对结构的已命名成员的引用；

在所有情况下，结果的访问权限都与原始引用的访问权限相同。

复合引用的组件引用代码示例：

struct S {
    age: i32,
    weight: f32
}

var<private> person: S;
// Elsewhere, 'person' denotes the reference to the memory underlying the variable,
// and will have type ref<private,S,read_write>.

fn f() {
    var uv: vec2<f32>;
    // For the remainder of this function body, 'uv' denotes the reference
    // to the memory underlying the variable, and will have type
    // ref<function,vec2<f32>,read_write>.

    // Evaluate the left-hand side of the assignment:
    //   Evaluate 'uv.x' to yield a reference:
    //   1. First evaluate 'uv', yielding a reference to the memory for
    //      the 'uv' variable. The result has type ref<function,vec2<f32>,read_write>.
    //   2. Then apply the '.x' vector access phrase, yielding a reference to
    //      the memory for the first component of the vector pointed at by the
    //      reference value from the previous step.
    //      The result has type ref<function,f32,read_write>.
    // Evaluating the right-hand side of the assignment yields the f32 value 1.0.
    // Store the f32 value 1.0 into the storage memory locations referenced by uv.x.
    uv.x = 1.0;

    // Evaluate the left-hand side of the assignment:
    //   Evaluate 'uv[1]' to yield a reference:
    //   1. First evaluate 'uv', yielding a reference to the memory for
    //      the 'uv' variable. The result has type ref<function,vec2<f32>,read_write>.
    //   2. Then apply the '[1]' array index phrase, yielding a reference to
    //      the memory for second component of the vector referenced from
    //      the previous step.  The result has type ref<function,f32,read_write>.
    // Evaluating the right-hand side of the assignment yields the f32 value 2.0.
    // Store the f32 value 2.0 into the storage memory locations referenced by uv[1].
    uv[1] = 2.0;

    var m: mat3x2<f32>;
    // When evaluating 'm[2]':
    // 1. First evaluate 'm', yielding a reference to the memory for
    //    the 'm' variable. The result has type ref<function,mat3x2<f32>,read_write>.
    // 2. Then apply the '[2]' array index phrase, yielding a reference to
    //    the memory for the third column vector pointed at by the reference
    //    value from the previous step.
    //    Therefore the 'm[2]' expression has type ref<function,vec2<f32>,read_write>.
    // The 'let' declaration is for type vec2<f32>, so the declaration
    // statement requires the initializer to be of type vec2<f32>.
    // The Load Rule applies (because no other type rule can apply), and
    // the evaluation of the initializer yields the vec2<f32> value loaded
    // from the memory locations referenced by 'm[2]' at the time the declaration
    // is executed.
    let p_m_col2: vec2<f32> = m[2];

    var A: array<i32,5>;
    // When evaluating 'A[4]'
    // 1. First evaluate 'A', yielding a reference to the memory for
    //    the 'A' variable. The result has type ref<function,array<i32,5>,read_write>.
    // 2. Then apply the '[4]' array index phrase, yielding a reference to
    //    the memory for the fifth element of the array referenced by
    //    the reference value from the previous step.
    //    The result value has type ref<function,i32,read_write>.
    // The let-declaration requires the right-hand-side to be of type i32.
    // The Load Rule applies (because no other type rule can apply), and
    // the evaluation of the initializer yields the i32 value loaded from
    // the memory locations referenced by 'A[4]' at the time the declaration
    // is executed.
    let A_4_value: i32 = A[4];

    // When evaluating 'person.weight'
    // 1. First evaluate 'person', yielding a reference to the memory for
    //    the 'person' variable declared at module scope.
    //    The result has type ref<private,S,read_write>.
    // 2. Then apply the '.weight' member access phrase, yielding a reference to
    //    the memory for the second member of the memory referenced by
    //    the reference value from the previous step.
    //    The result has type ref<private,f32,read_write>.
    // The let-declaration requires the right-hand-side to be of type f32.
    // The Load Rule applies (because no other type rule can apply), and
    // the evaluation of the initializer yields the f32 value loaded from
    // the memory locations referenced by 'person.weight' at the time the
    // declaration is executed.
    let person_weight: f32 = person.weight;
}

形成指针的方法有：

（1）对引用使用操作符（&）取址；

（2）如果函数形参是指针类型，那么当函数在运行时被调用时，形参的使用表示的是在调用函数的调用位置提供给相应操作数的指针值。

// Declare a variable in the private address space, for storing an f32 value.
var<private> x: f32;

fn f() {
    // Declare a variable in the function address space, for storing an i32 value.
    var y: i32;

    // The name 'x' resolves to the module-scope variable 'x',
    // and has reference type ref<private,f32,read_write>.
    // Applying the unary '&' operator converts the reference to a pointer.
    // The access mode is the same as the access mode of the original variable, so
    // the fully specified type is ptr<private,f32,read_write>.  But read_write
    // is the default access mode for function address space, so read_write does not
    // have to be spelled in this case
    let x_ptr: ptr<private,f32> = &x;

    // The name 'y' resolves to the function-scope variable 'y',
    // and has reference type ref<private,i32,read_write>.
    // Applying the unary '&' operator converts the reference to a pointer.
    // The access mode defaults to 'read_write'.
    let y_ptr: ptr<function,i32> = &y;

    // A new variable, distinct from the variable declared at module scope.
    var x: u32;

    // Here, the name 'x' resolves to the function-scope variable 'x' declared in
    // the previous statement, and has type ref<function,u32,read_write>.
    // Applying the unary '&' operator converts the reference to a pointer.
    // The access mode defaults to 'read_write'.
    let inner_x_ptr: ptr<function,u32> = &x;
}

与其他语言中的引用和指针的比较: