OpenCL for VEX users

Why OpenCL? ¶

VEX is ideal for most computation needs. But some types of algorithms have a structure that lend themselves to fast GPU execution. Since VEX can’t run on the GPU, a different language is needed to describe your algorithm.

OpenCL is a standard kernel language that supports both GPU and CPU implementations. This ensures your algorithm can run on machines lacking a GPU without the need for maintaining two versions.

A key requirement to gain the benefit of a GPU is to have the data stay on the GPU. If any non-OpenCL nodes process the geometry they will pull the data back to the CPU, often eliminating the performance benefit. Note in SOPs every SOP must bring the geometry back after it completes cooking to store in its local cache. This can be avoided by ensuring a chain of OpenCL nodes are in a Compiled Block.

Why Not OpenCL? ¶

OpenCL is an even less friendly language than VEX, though experience in VEX will make a lot of it familiar. Its use of pointers means it is easy to write unsafe code that might crash your video card driver or Houdini. While developing you may find yourself in a bad state where all kernels error on compile - restarting Houdini may be necessary to restore the driver. Rarely, you may need to restart your machine.

GPUs do not have virtual memory, so while normally allocating too much data will cause swapping and slowness, with OpenCL you instead get allocation failures and errors.

Syntax and Language ¶

OpenCL is a C-like language, as is VEX. Both share the same look - semicolons are required to terminate each line; braces are used to enclose blocks, etc.

OpenCL itself follows C99, so online resources for C99 can be referred to as well.

There is no foreach() in OpenCL.

Parameters to a function must have a type for each parameter and are separate by commas, not semicolons:

Vex:

int foo(float a; int b, c)
{
    return a + b + c;
}

OpenCL:

int foo(float a, int b, int c)
{
    return a + b + c;
}

Parameters in OpenCL are always passed by value. This means changing a parameter inside of the function will not affect the passed in value. But note that OpenCL has pointers, so changing what is pointed to will affect the original value.

While functions can be declared inside of functions in VEX, this is not possible in OpenCL.

Kernels ¶

Both VEX and OpenCL are kernel languages. The idea of a kernel language is to avoid the usual outer-loop for many algorithms. If one wants to add something to every point, in a traditional language you would have:

for (int i = 0; i < number_of_points; i++)
{
    process_point(i);
}

An explicit outer loop makes it extremely difficult to optimize, however, as process_point may depend on earlier calls. By instead only writing process_point, we both have simpler code but something that can be run in parallel easily. VEX uses this to swiftly process millions of points, and OpenCL uses it to ensure the GPU can efficiently run the program.

In a usual VEX Wrangle SOP the process_point() function is invisible. You write the code inside of process_point(), so you might have a snippet like:

@P += 1;

Behind the scenes the Wrangle nodes will build a wrapper function for the kernel, something like:

void process_point(vector _bound_P)
{
_bound_P += 1;
}

cvex dowrangle(export vector P)
{
    process_point(P);
}

The dowrangle kernel is what is passed to the VEX runtime to run across all the points.

In OpenCL nodes there isn’t the same auto-wrapper. Instead the equivalent of dowrangle must be written as a kernel function.

The equivalent OpenCL code would be:

kernel void dowrangle(
                 int P_length,
                 global float * restrict P
)
{
    int idx = get_global_id(0);
    if (idx >= P_length)
        return;

    float3 pos = vload3(idx, P);
    pos += 1;
    vstore3(pos, idx, P);
}

The OpenCL node must be explicitly configured to bind the P point attribute for writing.

To minimize the boilerplate and enable direct access to attributes without having to use vload or vstore, the OpenCL nodes also have an @-Binding mode. When this mode is enabled the kernel can be specified with @KERNEL and the appropriate boilerplate will be injected. This code then becomes:

#bind point &P float3

@KERNEL
{
    float3 pos = @P;
    pos += 1;
    @P.set(pos);
}

For SOPs where write back kernels are enabled, @WRITEBACK can be used to generate the write back kernel.

Run Over ¶

An important component of a VEX wrangle is the Run Over parameter. This controls what elements the kernel is executed on. For example, with Run Over set to points, kernel will run once for each point. The Volume Wrangle by its nature has an implicit run over of Volumes and VDBs - it will run the kernel for each voxel.

With OpenCL you also have to specify what the primary target of the kernel is. Note this doesn’t have to be what is written to, it merely defines how many workitems are given to the GPU.

The most straightforward run over modes are the First-Writable modes. These will find the first writable attribute, volume, or vdb, and use that as the template for the execution.

Unlike VEX that can stumble along even when forced to work serially with Run Over Detail; GPUs do not work well with serial execution and should always target many elements.

The equivalent of For Each Number is the Detail Attribute of Worksets. Instead of a fixed number, however, it takes a pair of detail array attributes on the incoming geometry to create the work items. This allows for very precise control of the GPU’s execution at the cost of a lot of complexity.

Dot operator ¶

In VEX you can access components of a vector using the dot operator (.). x,y,z, and w can be used, as can r, g, b, a or u, v. They can be used in conjunction to swizzle a vector, so v.zyx would create a three-vector in reverse order.

OpenCL has a similar notation. It also has .lo and .hi that will extract the top or lower half of a vector, and .odd and .even to pull out corresponding slices. Because it can have up to float16, there is also .s notation; .s0 will be the same as .x, and .s01 the same as .xy. To go past 9, a through f or A through F can be used for 10 through 15. Note OpenCL does not support the .u or .v variants.

Built-in Functions ¶

Unlike VEX, OpenCL has a very sparse set of built-in functions. VEX has a very rich set of operations to do everything from computing the error function to querying points on geometry. OpenCL’s internal functions are largely purely mathematical in nature.

Sadly, there is little in the way of built-in matrix support. matrix.h is provided (and included by default in snippets) which provides some methods for doing matrix multiplications.

Note that mat3 is actually 12 floats in size.

Note

select is in both OpenCL and VEX, but tragically they have different argument orders. VEX has the conditional first and OpenCL has the conditional last.

Matrix Functions ¶

The matrix.h functions of note are:

Function	Description
mat2 transpose2(const mat2 a)	Return `a` transposed.
mat3 transpose3(const mat3 a)	Return `a` transposed.
mat2 mat2mul(const mat2 a, const mat2 b)	Return `a * b`
fpreal2 mat2vecmul(const mat2 a, const fpreal2 b)	Return `b * A`
fpreal squaredNorm2(const mat2 a)	Return the square of the L2-norm of a 2×2 matrix.
fpreal squaredNorm3(const mat3 a)	Return the square of the L2-norm of a 3×3 matrix.
void mat3add(const mat3 a, const mat3 b, mat3 c)	`c = a + b`
void mat3sub(const mat3 a, const mat3 b, mat3 c)	`c = a - b`
mat3zero(mat3 a)	`a = 0`
mat3identity(mat3 a)	`a = identity`
mat3copy(const mat3 a, mat3 b)	`b = a`
mat3load(size_t idx, const global float *a, mat3 m)	Load into `m` the 9 floats stored starting at `a[idx*9]` This is loaded row-major.
mat3store(mat3 in, size_t idx, const global float *data)	Store 9 floats sequentially into `data[idx*9]` in a row-major manor.
void mat3fromcols(const fpreal3 c0, const fpreal3 c1, const fpreal3 c2, mat3 m)	Create a 3×3 matrix with columns of the specified vectors.
mat3 mat3mul(const mat3 a, const mat3 b)	Return `a * b`
fpreal3 mat3vecmul(const mat3 a, const fpreal3 b)	Return `b * A`
fpreal3 mat3Tvecmul(const mat3 a, const fpreal3 b)	Return `b * A^T`
fpreal2 mat3vecmul(const mat3 a, const fpreal3 b)	Return `b * A`. Discards third component.
fpreal2 mat3Tvecmul(const mat3 a, const fpreal3 b)	Return `b * A^T`. Discards third component.
void outerprod3(const fpreal3 a, const fpreal3 b, mat3 c)	Set c to the outerproduct of a and b.
void mat3lcombine(const fpreal s, const mat3 a, const fpreal t, const mat3 b, mat3 c)	`c = s * a + t * b`
fpreal det3(const mat3 a)	Return the determinant of `a`
fpreal3 diag3(const mat3 a)	Return diagonal vector of `a`
fpreal2 mat4vec2mul(const mat4 a, const fpreal2 b)	Returns `b * A`.
fpreal3 mat43vec3mul(const mat4 a, const fpreal3 b)	Returns b * A, assuming fourth component of b is 0.
fpreal3 mat4vec3mul(const mat4 a, const fpreal3 b)	Returns b * A, assuming fourth component of b is 1.
fpreal4 mat4vecmul(const mat4 a, const fpreal4 b)	Returns b * A.
int mat4invert(fpreal16 *m)	Update `m` with its inverse. Returns 0 on success, 1 if singular.
fpreal mat2det(const mat2 m)	Determinant of `m`
fpreal mat2inv(const mat2 m, mat2 *minvvout)	Sets minvout to the inverse of m. Returns the determinant.
fpreal mat3inv(const mat3 m, mat3 minvout)	Sets minvout to the inverse of m. Returns the determinant.
void mat3scale(mat3 mout, const mat3 m, fpreal scale)	`mout = m * scale`
void mat3lincomb2(mat3 mout, const mat3 m1, fpreal scale1, const mat3 m2, fpreal scale2)	`mout = m1 * scale1 + m2 * scale2`

Precision ¶

VEX has a single global precision it executes in. A float is 32-bits in 32-bit mode and 64-bit in 64-bit mode. Thus there is no distinction of double vs float.

OpenCL has fixed precision. However it is often the case that one wants to write code that will switch precision on demand, as 32-bit is much faster to compute; but some problems will need 64-bit.

The built-in OpenCL types have a specific precision. We have added two new types, fpreal and exint which have a precision that can depend on the node or incoming geometry, allowing code to be written once but work in several precisions.

Name	Precision	Type
half	16	Float
float	32	Float
double	64	Float
fpreal	Variable	Float
short	16	Integer
int	32	Integer
long	64	Integer
exint	Variable	Integer

Types ¶

OpenCL and VEX have similar basic types but with different names. The mat2/3/4 types are defined in matrix.h. float below could be replaced by any of the float precisions (half, float, double, and fpreal), as can int.

VEX	OpenCL	Type
int	int	Integer value
`--`	int2/3/4/8/16	Integer vector of the given number of components.
float	float	Float value
vector2	float2	Two component vector
vector	float3	Three component vector. Note that in OpenCL this is laid out in memory with four floats, so care has to be taken if used in a structure.
vector4	float4	Four component vector.
`--`	float8/16	OpenCL also has native 8 and 16 component vectors.
matrix2	float4, mat2	A float4 is the same size as a matrix2, but we've defined a mat2 to make the semantics more clear and provide some utility functions.
matrix3	mat3	There is no native 9-element type. The mat3 we have defined is actually a 12-element type, so care has to be done to move it to or from memory.
matrix	float16, mat4	A float16 is the same size as a matrix, but we've defined a mat4 to make the semantics clear and provide some utility functions.
string	`--`	You can define string literals for things like `printf()` in OpenCL, but as this is on the GPU there are many oddities, so these are not the same as the `const char *` you may expect from other C languages.
dict	`--`	Dictionaries aren’t supported in OpenCL.
`int[]`, `float[]`	`int `, `float `	VEX has an array type which can grow and shrink dynamically. In OpenCL, pointers can be used to a similar effect, but cannot change dynamically.

Binding and Prototypes ¶

VEX wrangles and snippets will usually auto-bind attributes when the @ syntax is used to refer to them. For complicated snippets, the Enforce Prototypes can be turned on that requires the binding to be explicitly listed before it is used.

Because of the auto-binding, VEX’s @ syntax has a prefix to specify the type where it isn’t implied.

OpenCL has no auto-binding, and also no type prefixes. Instead any @ binding must be explicitly provided either as a parameter in the Bindings tab of the node; or as a #bind directive in the snippet.

The #bind directive provides a way to specify the same information as in the Bindings parameters in a concise line. Where the same @ binding is controlled both by the #bind and a binding parameter, it is the binding parameter that is used.

The basic format of the #bind directive is

//    type name   flags...
#bind parm foobar read write

The type is what to bind to. The name is the token used in the @ statements to refer to this binding. By default it will match the attribute name, but a different attribute can be selected using the flags. (This is required if you have two attributes of the same name bound, you will need different names for them for the @ to disambiguate)

Flags are a space separated list of flags.

Binding Name Decorations ¶

The name can be decorated. Decorations can go before or after the name.

Decoration	Implied Flag
`&`	write
`?`	opt
`!`	noread

Binding Types ¶

Valid bindings words are listed here. Note the corresponding binding in the multiparm can provide more details of how these work.

Binding	Type
parm	Matches integer, float, float vec style bindings. A flag of `float`, `float3`, `int`, etc, will specify which type. This is used to control the snippet with a parameter constant across all work items.
point/prim/primitive/global/detail/vertex	Binds a geometry attribute of this class.
ramp	Binds a ramp. The flag of `float` or `float3` controls if it is a scalar or vector ramp.
volume	Binds a volume primitive in a geometry.
vdb	Binds a VDB primitive in a geometry.
option	Binds a DOP Option in DOPs.
sfield/vfield/mfield/scalarfield/vectorfield/matrixfield	Binds to a DOP field of the provided type.
layer	Binds to a COP layer.

Type Flag ¶

One type of flag is to use an OpenCL type as the flag, such as float or float2. This then specifies the vector size, precision, etc, of the binding. float[] can be used to specify a binding to an array attribute. In addition, float? and int? can be used to bind to attributes of arbitrary size. The tuplesize method can be used to see what the actual bound size was.

Access Flags ¶

Access flags control how the source attribute is bound to the kernel.

Flag	Meaning
read / noread	Defaults to read. If a binding is readable it will have its OpenCL buffer initialized from its original value. If you do not mark it as readable, it is important you overwrite it as otherwise arbitrary data will be present.
write / nowrite	Defaults to nowrite. A binding not marked writable will not be written back to the CPU as it will be assumed to be unchanged. Note you may still actually write to it if you try hard enough. Doing so will result in undefined behavior.
opt / noopt	Defaults to noopt. An optional binding doesn’t need to exist on the original geometry. The macro define HAS_name will be defined if the binding does exist, allowing one to have different code paths. Alternatively, if the def flag is present, the missing binding will be replaced by a single default value.
def / nodef	Defaults to def. If an optional binding isn’t present, will instead bind a parameter in the same place. For example, if one binds the point attribute `mass` optionally and mass isn’t present, a float parameter `mass` will be bound in its place. This lets code use @mass to refer to the value and switch between the default value or the geometry value without any code changes or use of HAS_mass #ifdefs.

Target Flags ¶

Binding target flags control exactly what is used for the source of the binding. The name used for the @ is always the binding name, which often forms the default for what is bound, but these allow finer control.

Flag	Meaning
val=#	Defaults to 0. Set the default value for the parameter, volume, or vdb. If it is a vector binding and a single scalar is provided, it is repeated, so `val=1` on a vector will use a default value of `(1, 1, 1)`. Vector values can be specified by using braces, so `val={1, 2, 3}` will use the corresponding default value.
name=SSS	Defaults to the binding name. The name of the attribute, volume, or VDB to bind to.
input=#	Defaults to 0. In SOPs, the input to bind to. Only the first input can be written to.
geo=SSS or geometry=SSS	Defaults to Geometry. In DOPs, the simulation geometry to bind to, relative to the root of the current simulation object.
port=SSS or portname=SSS	Defaults to the binding name. The name of the port, ie: input or output, in COPs to bind to.

Feature Flags ¶

In addition there are general feature flags that are often specific to certain binding types. These may be needed to enable certain methods on the binding as well. These correspond to toggles in the Bindings multiparm of either the OpenCL SOP or DOP, so more details can be found there as well.

Flag	Meaning
fieldoffsets / nofieldoffsets	Default fieldoffsets. For Field Bindings in DOPs, enable the offset parameters for this field.
forcealign / noforcealign	Default forcealign. Requires all the fields or volumes to have the same resolution. This lets the code be simpler and allow assumptions about samples lining up. Note that the first writeable field or volume (the one being iterated over…) should not be forcealigned, so its default is noforcealign.
resolution / noresolution	Default noresolution. Include the resolution of the field or volume. This is required for bounds checking, but only usually required for the first volume if the others are aligned. The first writeable volume has a default of resolution.
voxelsize / novoxelsize	Default novoxelsize. Include the size of the voxels of the volume.
xformtoworld / noxformtoworld	Default noxformtoworld. Include the transform from index space to geometry (SOP) space.
xformtovoxel / noxformtovoxel	Default noxformtovoxel. Include the transform from geometry (SOP) space to index space.

Spare Parameter Creation ¶

In VEX you can instrument your function by including ch("myparm") directly in the code. Then the generate-parameter button can be pressed and a spare parameter myparm will be created that the VEX code will read at execution time.

OpenCL, running on the GPU, cannot have the equivalent of ch("myparm") as everything must be explicitly bound. Thus the equivalent is to have a #bind directive such as #bind parm myparm float.

To actually control the value that is bound, you can manually add myparm as a binding in the Bindings tab. By setting it to float type, the float value can be animated there to control what is passed to OpenCL.

To make things easier, the generate-parameter button on the OpenCL node will parse all the #bind directives and create a corresponding entry in the Bindings page, but also create a spare parameter next to the code that drives the multiparm value. It is important to note the resulting spare parameter is not directly read by the OpenCL code - it is driving the value of the multiparm, and that mulitparm is then bound to the kernel before execution.

Topology and Group Binding Names ¶

Vex has several magic names that affect the binding process. @group_foo will bind to the group foo, for example, and @opinput1_P will bind to P on the second input.

OpenCL also has similar options. Binding to the attribute group:foo will bind to the group foo. Note the binding has to be of integer type.

Several topology attributes are also supported using the topo: prefix. These must be bound to the correct class and type. These are all read-only. Where possible they closely match the corresponding VEX functions.

Topology Name	Class	Type	Meaning
topo:primpoints	prim	`int[]`	Indexed by primitive number, stores all the points indices belonging to that primitive.
topo:primvertices	prim	`int[]`	Indexed by primitive number, stores all the vertex indices belonging to that primitive.
topo:primvertexcount	prim	`int`	Indexed by primitive number, stores the number of vertices on that primitive.
topo:pointprims	point	`int[]`	Indexed by point number, stores the primitive indices that refer to that point.
topo:pointvertices	point	`int[]`	Indexed by point number, stores the vertex indices that refer to that point.
topo:pointneighbours	point	`int[]`	Indexed by point number, stores the point indices of the ring neighbors of the point; as per the `neighbours` VEX function.
topo:vertexpoint	vertex	`int`	Indexed by vertex number, stores the point index that the vertex refers to.
topo:vertexprim	vertex	`int`	Indexed by vertex number, stores the primitive index that the vertex is owned by.
topo:vertexprimindex	vertex	`int`	Indexed by vertex number, stores where in the primitive this vertex belongs.

Global Bindings ¶

VEX has various default bindings that are presented in many wrangles. OpenCL likewise has some, but while VEX will often bind these just-in-time if they show up in your code, OpenCL often requires an explicit feature flag to be set on the node to enable the binding.

Binding	Type	Validity	Meaning
`@ix`, `@iy`, `@iz`, `@ixy`	`int`, `int2`	Running over volumes, VDBs, layers, or fields.	The currently processed voxel/buffer integer index.
`@xres`, `@yres`, `@res`	`int`, `int`, `int2`	Running over layers.	The output buffer resolution.
`@tilesize`	`int2`	Running over layers.	The tile size. The kernel will be run once for each tile of this size. `0` means the tile is implied to be the full size of the image in that dimension. The created tiles may extend past the actual image size.
`@elemnum`	`int`	Running over volumes, attributes, or VDBs. Running over fields where the fields are aligned.	An integer index representing the currently processed element. Note for VDBs this will vary with the topology.
`@P`	`float2/3`	Running over layers or fields.	The current Image/DOP/Model space of the processed voxel. In layer mode, .image, .texture, .pixel, and .world specify a space with .image the default.
`@dPdx`, `@dPdy`	`float2`	Running over layers.	The change in image space when the buffer coordinate changes by one in either x or y. While this is a float2, only one component will be non-zero. .image, .texture and .pixel specify a space with .image the default.
`@dPdxy`	`float2`	Running over layers.	The size of the current output buffer element in image space. Equivalent to (@dPdx.x, @dPdy.y) .image, .texture and .pixel specify a space with .image the default.
`@Time`	`float`	Include Time specified on node	Current time in seconds. Makes the node time dependent.
`@TimeInc`	`float`	Include Timestep specified on node	The current timestep, usually `1/$FPS` in SOPs, but the current simulation step size if in DOPs.
`@SimFrame`	`int`	Include Simulation Frame specified on node	The current simulation frame, `$SF`.

@-Binding Methods ¶

In VEX an @ binding is the actual value. @P will be the vector of the current position and be of type vector. Writing to the current position is done by simply assigning to @P. To work with position on points other than the current point, the point and setpointattrib functions are used.

In VEX the @ binding is effectively a local variable - writing or reading to it is the same as any variable. In OpenCL, the @ binding may represent more computation, so it is often useful to copy it into a local variable if it is to be used extensively.

With OpenCL reading and writing to arbitrary attributes such as done in the point and setpointattrib functions is not possible as all data has to be explicitly bound. But this also means that when the position attribute is bound with @P, all of the position data is available for reading. Access to the full data is provided by adding various methods to the @ binding. @P.getAt(53) will read the point with index 53, for example. These methods include bounds checking to ensure that your kernel does not fail catastrophically due to programming errors, much like the point and setpointattrib methods.

The setpointattrib functions in VEX are designed to be thread-safe and always operate the same. The corresponding setAt functions in OpenCL have no such guarantee, it is important you are careful not to have two work items write to the same point.

What methods are available depend on the type of the binding, the type of data it is bound to, whether the data is readable or writable, and whether the data is aligned to the current work item. In particular, the last point is that you can only refer to @P as a raw value if the position attribute, or a matching attribute, is being run over. Otherwise which point you want must be specified.

Parameter Binding Methods ¶

A parm binding is merely the value itself. Thus there are no methods as @name will refer to the actual value.

Ramp Binding Methods ¶

A ramp in OpenCL is always a 1d array of floats that is interpolated to get the ramp values. The size of the ramp is specified in the binding. The provided methods ensure one does not have to worry about the nuances of the data layout, however.

Method	Meaning
`@name.data`	The float array of the actual ramp data.
`@name.len`	The size of the ramp.
`@name(float)`, `@name.getAt(float)`	Look up the value of the ramp at the 0..1 position requested, interpolating between the samples.

Field Binding Methods ¶

The vector and matrix fields have _x and _xx suffixes to refer to the different components. So if the binding name is name, a vector field will have @name.stride_x_y for the y stride of the x field of the vector field.

A field is aligned if it has matching resolution and transform to the field being run over.

Method	Meaning
`@name.data`	The raw float data of the field.
`@name.stride_x/y/z/offset`	The stride of a step in the x, y, and z axes of the data. The offset is an offset from the data to the 0,0,0 voxel.
`@name.x/y/zres`	The x, y, or z resolution of the field.
`@name`, `@name.get`	The current processed value. Valid if aligned.
`@name.getAt(x,y,z)`	The the voxel at index location x, y, z.
`@name.set(val)`	Set the currently processed voxel to val. Valid if aligned.
`@name.setAt(x, y, z, val)`	Set the voxel at (x, y, z) to val.
`@name.sample(xyz)`	Evaluate the field in index space at `float3 xyz`.

Attribute Binding Methods ¶

Attributes bind to geometry attributes, so may be point, primitive, vertex, or detail. These include groups, which have a group: prefix and topology that have a topo: prefix.

An attribute has a matching index if it the kernel is running over attributes and the attribute is the same class and size as the attribute that is being run over.

Method	Meaning
`@name.data`	The raw attribute data. All the element data is flattened into a single array with no holes or pages. For array attributes, the array data is stored contiguously. Usually one will not directly access this but use the other methods that provide bounds checked methods.
`@name.len`	The length of the attribute. This is the number of points, primitives, or vertices it has.
`@name`, `@name.get`	The attribute value of the current index. Valid if readable. Valid if not an array. Valid if matching index.
`@name(idx)`, `@name.getAt(idx)`	Get the attribute at a specific index. Valid if readable. Valid if not an array.
`@name.set(val)`	Set the attribute value at the current index. Vaild if writable. Valid if not an array. Valid if matching index.
`@name.setAt(idx, val)`	Set the attribute at the given index to the value. Valid if writable. Valid if not an array.
`@name.tuplesize`	The size of each element of the attribute. A float attribute has size `1`, a vector attribute would have size `3`.
`@name.tuple`	A pointer to an array of the currently processed item. This is helpful if working with attribute sizes that do not map to OpenCL native types. To access the elements one would use `[]`, as in `@name.tuple[2]` to get the third part of the current attribute. Note in this case you must do bounds checking to ensure you do not read past the attribute’s tuplesize. Valid if matching index.
`@name.tupleAt(idx)`	A pointer to an array of a specific item at index `idx`. Like `tuple`, care must be taken in using this.
`@name.bound`	Whether the attribute is bound or not. For optional attributes this can be 1 if bound and 0 if not. Non-optional attributes are always bound as the kernel will not be run if they aren’t present.
`@name.comp(compidx)`	For an array attribute, extracts a the specific component of the current index’s data. Valid if an array. Valid if matching index.
`@name.entries`	The length of the current index’s array. Valid if an array. Valid if matching index.
`@name.compAt(idx, compidx)`	For an array attribute, extracts the specific component of a specific index’s data. `idx` is the element number in the attribute, ie, point number or primitive number. `compidx` is the where in that element’s array to access. Valid if an array.
`@name.entriesAt(idx)`	The length of the index `idx`'s array. Valid if an array.

Volume Binding Methods ¶

Volumes on geometry can bound.

A volume has a matching index if the kernel is running over a volume and the bound volume matches its resolution and transform.

Volumes have a fair number of flags to add bind extra data that can be read by the methods. Some methods may require extra flags to work, especially if the volume isn’t aligned with the volume being run over.

Method	Meaning
`@name.data`	The raw data of the volume. This is a flat array of the volume’s values indexed using the `stride_` values. It is recommended to use `getAt` or similar methods rather than directly use the `data` as it will have bounds checking.
`@name.bound`	Whether the volume is bound or not. For optional volume this can be 1 if bound and 0 if not. Non-optional volumes are always bound as the kernel will not be run if they aren’t present.
`@name.stride_x/y/z/offset`	The step size of the x, y, and z axes, along with the offset of the 0,0,0 voxel from the start of the `data` array.
`@name.x/y/zres`	The resolution of the volume. Note it is often padded so the stride must be used to index.
`@name.voxelsize_x/y/z`	The size of the 0,0,0 voxel along each of the axes. Requires `voxelsize` flag.
`@name.xformtoworld`	A `fpreal16` that maps a voxel index into the model/SOP space. Requires `xformtoworld` flag.
`@name.xformtovoxel`	A `fpreal16` that maps a model/SOP space into index space. Requires `xformtovoxel` flag.
`@name`, `@name.get`	The value of the volume at the current voxel index. Valid if readable. Valid if matching index.
`@name.valid(x,y,z)`	True if voxel x, y, z is within the bounds of the volume.
`@name.getAt(x,y,z)`	Returns the value at voxel (x, y, z). If out of the range of the volume, the index is clamped to the volume’s range. Valid if readable.
`@name.set(val)`	Sets the currently processed voxel to `val`. Valid if writable. Valid if matching index.
`@name.setAt(x, y, z, val)`	Sets the voxel at (x, y, z) to `val`. This is bounds checked and a no-op if it is out of range. Valid if writable.
`@name.sample(xyz)`	Trilinearly interpolates the volume at the index-space `fpreal3` location `xyz`. Clamps to within the valid range. Valid if readable.
`@name.sample2d(xy)`	Bilinearly interpolates the volume at the index-space `fpreal2` location `xy`. Clamps to within the valid range. Valid if readable.

VDB Binding Methods ¶

VDBs on a geometry can bound. VDBs in the OpenCL world have been transformed into a NanoVDB layout. While CNanoVDB and similar headers can be used to directly access them, the methods provided will simplify the access patterns so should be preferred.

VDBs have a matching index if it is running over VDBs and the VDB has matching transforms.

While VDBs can be written to, doing so cannot change their topology. You can only write to voxels that were already activated outside of OpenCL. Note that constant regions of a VDB are often compressed into constant tiles, so cannot be written to. These need to be explicitly densified - a volume wrangle adding a random number per voxel will do this.

VDBs have a fair number of flags to add bind extra data that can be read by the methods. Some methods may require extra flags to work, especially if the VDB isn’t aligned with the VDB being run over.

VDBs are usually bound explicitly to a type using the float or float3 type tags. If no tag is specified, they are bound as any and fewer methods are available.

Method	Meaning
`@name.data`	The raw NanoVDB data. The `CNanoVDB` header can be used to inspect this directly, but using methods is the preferred approach as CNanoVDB is very verbose.
`@name.bound`	Whether the VDB is bound or not. For optional VDBs this can be 1 if bound and 0 if not. Non-optional VDBs are always bound as the kernel will not be run if they aren’t present.
`@name.voxelsize_x/y/z`	The size of the VDB (0,0,0) voxel in the x, y, and z axes. Requires `voxelsize` flag.
`@name.xformtoworld`	An `fpreal16` transform from index space into model/SOP space. Requires `xformtoworld` flag.
`@name.xformtovoxel`	An `fpreal16` transform from model/SOP space to index space. Requires `xformtovoxel` flag.
`@name`, `@name.get`	Returns the value at the current processed index. Valid if readable. Valid if matching index.
`@name.pos`	Returns the model/SOP location of the currently processed index. Valid if matching index.
`@name.active`	Whether the currently processed VDB index is active on this VDB. Requires `vector` or `float` binding. Requires matching index.
`@name.activeAt(x, y, z)`	Whether the provided x, y, z index coordinates are active. Valid if `vector` or `float` binding.
`@name.valid(x, y, z)`	Whether the provided x, y, z index coordinates are active. Valid if `vector` or `float` binding. Valid if matching index.
`@name.getAt(x, y, z)`	Returns the value at the index (x, y, z). VDBs possess background values in inactive areas so this can be run on any index. Valid if `vector` or `float` binding. Valid if readable.
`@name.set(val)`	Sets the currently processed voxel to the value `val`. Valid if `vector` or `float` binding. Valid if writable. Valid if matching index.
`@name.setAt(x, y, z, val)`	Sets the voxel at index (x, y, z) to the value `val`. If the index isn’t in the VDBs topology the write will do nothing. Valid if `vector` or `float` binding. Valid if writable.
`@name.sample(xyz)`	Trilinearly interpolates the VDB at the provided index space location. Valid if `vector` or `float` binding. Valid if readable.
`@name.worldSample(xyz)`	Trilinearly interpolates the VDB at the provided model/SOP space location. Valid if `vector` or `float` binding. Valid if readable.
`@name.worldGradient(xyz)`	Computes the gradient of the VDB at the provided model/SOP space and returns the gradient in the model/SOP space. Valid if `float` binding. Valid if readable.

Option Binding Methods ¶

Options on simulation data can be bound by the DOP OpenCL nodes.

Which data to bind to is controlled by the geo flag. Which option is controlled by the name flag.

Method	Meaning
`@name`, `@name.get`	The value of the option on the specified data.
`@name.bound`	Whether the option is bound or not. For optional options this can be 1 if bound and 0 if not. Non-optional options are always bound.

Layer Binding Methods ¶

Layers in COPs can be bound. Layers within OpenCL are dense 2d arrays of buffer elements. To support cooking at variable resolutions, the actual data elements of the array are not necessarily the same as the pixels of the final image.

A layer binding of a specific type will have .set and .get methods of that type. However, that may not be the actual data layout of the bound data - implicit conversion is done with the accessors. Thus a layer binding of float4 can take a 1-channel grey scale buffer and it will see the grey channel replicated across the xyzw values.

Note

While VEX supports .rgba access, and some OpenCL drivers support it, this is not part of the standard. Use .xyzw.

Method	Meaning
`@name.data`	The raw buffer data. It is recommended to use the proper accessors to get at this data. Note that the actual layout may not match the bound layer type, as just-in-time conversion is often done.
`@name.stat`	A pointer to the IMX_Stat structure describing various metadata of this layer. Usually one will use an accessor that ties directly to the layer values.
`@name.bound`	For optional layers, whether they have been bound. Mandatory layers will always be 1, as otherwise the kernel won’t run.
`@name.xres`, @name.yres`,` @name.res`	Buffer resolution of the layer.
`@name.border`	The border type of the layer. Of type BorderType. Can be CONSTANT, CLAMP, MIRROR, or WRAP.
`@name.storage`	The storage type of the layer. Of type StorageType. Can be INT8, INT16, INT32, FLOAT16, FLOAT32. There are additional CONSTANT_ prefixed versions for constant buffers.
`@name.channels`, `@name.tuplesize`	The number of channels in the layer. This is in the actual data itself, so may differ from what is specified on the #bind.
`@name.imageToBuffer(xy)`	Returns the image coordinate transformed to buffer coordinates.
`@name.bufferToImage(xy)`	Returns the buffer coordinate transformed to image coordinates.
`@name.pixelToBuffer(xy)`	Returns the pixel coordinate transformed to buffer coordinates.
`@name.bufferToPixel(xy)`	Returns the buffer coordinate transformed to pixel coordinates.
`@name.imageToWorld(xy)`	Returns the image coordinate transformed to world coordinates.
`@name.image3ToWorld(xyz)`	Returns the image float3 coordinate transformed to world coordinates.
`@name.worldToImage(xyz)`	Returns the world coordinate transformed to image coordinates.
`@name.worldToImage3(xyz)`	Returns the world coordinate transformed to image float3 coordinates.
`@name.bufferIndex(ixy)`	Returns value of buffer element at integer coordinates.
`@name.bufferSample(xy)`	Bilinearly interpolates buffer elements at floating point buffer coordinates.
`@name.imageSample(xy)`	Bilinearly interpolates at location in image space.
`@name.imageNearest(xy)`	Closest buffer element to location in image space.
`@name.textureSample(xy)`	Bilinearly interpolates at location in texture space.
`@name.textureNearest(xy)`	Closest buffer element to location in texture space.
`@name`	Bilinear interpolation of this buffer at currently processed output. Equivalent to `@name.imageSample(@P)`
`@name.set(v)`	Requires `name` is writable and aligned to output buffer. Sets currently processed buffer element to v.
`@name.setIndex(ixy, v)`	Requires `name` is writable. Sets buffer element at integer ixy to v.
`@name.dCdx`, `@name.dCdy`, `@name.dCdx(xy)`, `@name.dCdy(xy)`	The change in the values of `name`. This is measured along either the x or y axes of the output buffer. The location is either @P if no parameter, or the given position in image space. Only for non-integer layers.
`@name.bufferSampleRect(centerxy, sizewh)`, `@name.bufferSampleRectClip(centerxy, sizewh)`, `@name.imageSampleRect(centerxy, sizewh)`, `@name.imageSampleRectClip(centerxy, sizewh)`, `@name.textureSampleRect(centerxy, sizewh)`, `@name.textureSampleRectClip(centerxy, sizewh)`	Sample a rectangular region using the node’s filter settings. The region is specified in either buffer or image coordinates. The clip versions clip the sample region at the source image boundary rather than using border conditions to sample outside.

OpenCL for VEX users

Why OpenCL? ¶

Why Not OpenCL? ¶

Syntax and Language ¶

Kernels ¶

Run Over ¶

Dot operator ¶

Built-in Functions ¶

Matrix Functions ¶

Precision ¶

Types ¶

Binding and Prototypes ¶

Binding Name Decorations ¶

Binding Types ¶

Type Flag ¶

Access Flags ¶

Target Flags ¶

Feature Flags ¶

Spare Parameter Creation ¶

Topology and Group Binding Names ¶

Global Bindings ¶

@-Binding Methods ¶

Parameter Binding Methods ¶

Ramp Binding Methods ¶

Field Binding Methods ¶

Attribute Binding Methods ¶

Volume Binding Methods ¶

VDB Binding Methods ¶

Option Binding Methods ¶

Layer Binding Methods ¶

VEX

Language ¶

Next steps ¶

Reference ¶