NOTE: To decide which address space (
device
or
constant
) a read-only buffer passed
to a graphics or kernel function uses, look at how the buffer is accessed inside the
graphics or kernel function. The
constant
address space is optimized for multiple
instances executing a graphics or kernel function accessing the same location in the
buffer. Some examples of this access pattern are accessing light or material properties
for lighting / shading, matrix of a matrix array used for skinning, filter weight accessed
from a filter weight array for convolution. If multiple executing instances of a graphics
or kernel function are accessing the buffer using an index such as the vertex ID,
fragment coordinate, or the thread position in grid, the buffer must be allocated in the
device
address space.
4.2.5
thread Address Space
The
thread
address space refers to the per-thread memory address space. Variables allocated
in this address space are not visible to other threads. Variables declared inside a graphics or
kernel function are allocated in the
thread
address space.
kernel void
my_kernel(…)
{
// A float allocated in the per-thread address space
float x;
// A pointer to variable x in per-thread address space
thread float p = &x;
…
}
4.3 Function Arguments and Variables
Most inputs and outputs to a graphics and kernel functions are passed as arguments.
(Exceptions are initialized variables in the constant address space and samplers declared in
program scope.) Arguments to graphics (vertex and fragment) and kernel functions can be one
of the following:
• device buffer – a pointer or reference to any data type in the
device
address space (see
section 2.6)
• constant buffer – a pointer or reference to any data type in the
constant
address space
(see section 2.6)
•
texture
object (see section 2.8) or an array of textures
•
sampler
object (see section 2.9) or an array of samplers
2017-9-12 | Copyright © 2017 Apple Inc. All Rights Reserved.
Page of
58
174
• a buffer shared between threads in a threadgroup – a pointer to a type in the
threadgroup
address space that can only be used as arguments for kernel functions.
• imageblock (see section 2.10)
• argument buffer (see section 2.12)
• A struct with elements that are either buffers or textures.
Buffers (device and constant) specified as argument values to a graphics or kernel function
cannot alias; i.e., a buffer passed as an argument value cannot overlap another buffer passed to
a separate argument of the same graphics or kernel function.
Arguments to graphics and kernel functions cannot be declared to be of type
size_t
,
ptrdiff_t
, or a struct and/or union that contains members declared to be one of these built-in
scalar types.
The arguments to these functions are often specified with attributes to provide further guidance
on their use. Attributes are used to specify:
• the resource location for the argument (see section 4.3.1),
• built-in variables that support communicating data between fixed-function and
programmable pipeline stages (see section 4.3.3),
• which data is sent down the pipeline from vertex function to fragment function (see
section 4.3.5).
4.3.1
Attributes to Locate Buffers, Textures and Samplers
For each argument, an attribute can be optionally specified to identify the location of a buffer,
texture or sampler to use for this argument type. The Metal framework API uses this attribute to
identify the location for these argument types.
• device and constant buffers –
[[buffer(index)]]
• texture and an array of textures –
[[texture(index)]]
• sampler –
[[sampler(index)]]
• threadgroup buffer –
[[threadgroup(index)]]
The
index
value is an unsigned integer that identifies the location of a buffer, texture or sampler
argument that is being assigned. The proper syntax is for the attribute to follow the argument/
variable name.
The example below is a simple kernel function,
add_vectors
, that adds an array of two buffers
in the device address space,
inA
and
inB
, and returns the result in the buffer out. The attributes
(
buffer(index)
) specify the buffer locations for the function arguments.
kernel void
add_vectors(const device float4 *inA [[buffer(0)]],
const device float4 *inB [[buffer(1)]],
device float4 *out [[buffer(2)]],
2017-9-12 | Copyright © 2017 Apple Inc. All Rights Reserved.
Page of
59
174
uint id [[thread_position_in_grid]])
{
out[id] = inA[id] + inB[id];
}
The example below shows attributes used for function arguments of several different types (a
buffer, a texture, and a sampler):
kernel void
my_kernel(device float4 *p [[buffer(0)]],
texture2d img [[texture(0)]],
sampler sam [[sampler(1)]])
{…}
If the location indices are not specified the Metal compiler assigns them using the first available
location index. In the following example,
src
is assigned texture index 0,
dst
texture index 1,
s
sampler index 0 and
u
buffer index 0.:
kernel void
my_kernel(texture2d src,
texture2d dst,
sampler s,
device myUserInfo *u)
{…}
In the following example, some kernel arguments have explicitly assigned location indices and
some do not.
src
is explicitly assigned texture index 0, and
f
is explicitly assigned buffer index
10. The other arguments are assigned the first available location index:
dst
texture index 1,
s
sampler index 0, and
u
buffer index 0.
kernel void
my_kernel(texture2d src [[texture(0)]],
texture2d dst,
sampler s,
device myUserInfo *u,
device float *f [[buffer(10)]])
{…}
2017-9-12 | Copyright © 2017 Apple Inc. All Rights Reserved.
Page of
60
174
Dostları ilə paylaş: |