Types and Values

Goal. Give a node a number, vector or matrix authored in USD, and read it in the GPU code.

Before you start. A working node, from Your First Node.

The Shape

Declare a typed attribute on the shader definition. Anything not opaque is a value-input:

float inputs:strength = 1.0

How It Works

A value-input is declared once, alongside the opaque ports:

InvertKernel.usda, from the runnable pipeline example

# Reusable SPG shader definition with a typed value-input (strength) alongside
# the opaque AOV ports. SPG finds the launch script by appending ".lua" to the
# source asset path (InvertKernel.cu -> InvertKernel.cu.lua), so keep them
# co-located.
def Shader "InvertKernel"
{
    uniform token info:implementationSource = "sourceAsset"
    uniform asset info:spg:sourceAsset = @InvertKernel.cu@
    uniform token info:spg:sourceAsset:subIdentifier = "invert"

    float inputs:strength = 1.0

    opaque inputs:Image
    opaque outputs:Inverted
}

From there it reaches the GPU by a different route in each language.

CUDA

It becomes a kernel parameter, declared in the signature:

InvertKernel.cu, from the runnable pipeline example

// Second pass of the pipeline: invert an image. The input is whatever upstream
// shader is connected to inputs:Image; here, the grayscale output.
extern "C" __global__ void invert(
    int width,
    int height,
    float strength,
    cudaTextureObject_t inputImage,
    cudaSurfaceObject_t outputInverted)
{
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    if (x < width && y < height)
    {
        uchar4 pixel = tex2D<uchar4>(inputImage, x, y);

        // lerp(original, 255 - original, strength) per RGB channel.
        unsigned char r = (unsigned char)(pixel.x + strength * (255 - 2 * pixel.x));
        unsigned char g = (unsigned char)(pixel.y + strength * (255 - 2 * pixel.y));
        unsigned char b = (unsigned char)(pixel.z + strength * (255 - 2 * pixel.z));

        uchar4 out = { r, g, b, pixel.w };
        surf2Dwrite<uchar4>(out, outputInverted, x * sizeof(uchar4), y);
    }
}

The launch script places it in args at the position the signature expects:

InvertKernel.cu.lua, from the runnable pipeline example

function invert(inputs, outputs)
    assert(#inputs["Image"].shape == 2, "Input must be a 2D image")
    assert(inputs["Image"].dtype == cuda.uchar4, "Input must be uchar4")

    local height = inputs["Image"].shape[1]
    local width  = inputs["Image"].shape[2]

    outputs["Inverted"] = cuda.image(width, height, cuda.uchar4)

    return cuda.kernel({
        -- void invert(int, int, float, cudaTextureObject_t, cudaSurfaceObject_t)
        args = {
            cuda.int(width),
            cuda.int(height),
            cuda.float(inputs["strength"]),
            cuda.TextureObject(inputs["Image"]),
            cuda.SurfaceObject(outputs["Inverted"]),
        },
        block = { 32, 32 },
        grid  = { math.ceil(width / 32), math.ceil(height / 32) },
    })
end

Slang

Values do not arrive individually. They are packed into one constant buffer, which the shader declares as a struct and reads through a ParameterBlock:

InvertKernel.slang, from the runnable pipeline example

// Second pass of the pipeline: invert an image. Slang counterpart of
// InvertKernel.cu. The input is whatever upstream shader is connected to
// inputs:Image; here, the grayscale output.
//
// On Slang, value-inputs are packed into
// a constant buffer described by a struct, which the shader reads through a
// ParameterBlock. The struct field order must match the order the launch
// script lists the values in slang.ParameterBlock(...).
struct Params
{
    float strength;
};

[[vk::binding(0, 1)]] ParameterBlock<Params> g_Params;
[[vk::binding(1, 1)]] Texture2D<float4> g_InImage;
[[vk::binding(2, 1)]] RWTexture2D<float4> g_OutInverted;

[shader("compute")]
[numthreads(32, 32, 1)]
void invert(uint3 tid : SV_DispatchThreadID)
{
    uint width = 0, height = 0;
    g_InImage.GetDimensions(width, height);

    if (tid.x >= width || tid.y >= height)
        return;

    float4 pixel = g_InImage.Load(int3(tid.xy, 0));

    // lerp(original, inverted, strength) per RGB channel. The CUDA kernel does
    // the same blend on 0..255 integers; here the channels are normalized
    // floats, so full inversion is 1.0 - value.
    float3 inverted = lerp(pixel.rgb, 1.0 - pixel.rgb, g_Params.strength);

    g_OutInverted[tid.xy] = float4(inverted, pixel.a);
}

The launch script fills that block, in the order the struct declares the fields:

InvertKernel.slang.lua, from the runnable pipeline example

-- Launch script for InvertKernel.slang. The function name must equal the
-- subIdentifier ("invert").
function invert(inputs, outputs)
    assert(inputs["Image"].rank == 2, "Input must be a 2D image")

    outputs["Inverted"] = slang.image(inputs["Image"].shape, inputs["Image"].dtype)

    return slang.compute({
        -- The bind list is positional. The parameter block comes first and
        -- carries the value-inputs in the order the shader's Params struct
        -- declares them; the resources follow in declaration order.
        bind = {
            slang.ParameterBlock(
                slang.float(inputs["strength"])      -- -> float strength
            ),
            slang.Texture2D(inputs["Image"]),        -- -> Texture2D<float4> g_InImage
            slang.RWTexture2D(outputs["Inverted"]),  -- -> RWTexture2D<float4> g_OutInverted
        },
        -- The shader's [numthreads(32, 32, 1)] places the dispatch; the grid
        -- follows from the output's shape.
    })
end

A value-input is a wrapper, not a number. inputs["strength"] cannot be used in arithmetic; that raises an error. Pass it through a dtype constructor to hand it to the GPU, or read .value to compute with it in Lua.

A scene can override the default. The scene states the value on the shader instance; a value stated there overrides the definition’s default.

pipeline_scene.usda, from the runnable pipeline example

def Scope "Render"
{
    def RenderProduct "PipelineDemo"
    {
        uniform int2 resolution = (1280, 720)
        rel camera = </World/Camera>
        rel orderedVars = [ <LdrColor>, <LdrInverted> ]

        def RenderVar "LdrColor"
        {
            uniform string sourceName = "LdrColor"
            opaque omni:rtx:aov
        }

        def RenderVar "LdrInverted"
        {
            uniform string sourceName = "LdrInverted"
            opaque omni:rtx:aov.connect = <../InvertKernel.outputs:Inverted>
        }

        def Shader "GrayscaleKernel" (
            references = @GrayscaleKernel.usda@
        )
        {
            opaque inputs:LdrColor.connect = <../LdrColor.omni:rtx:aov>
        }

        def Shader "InvertKernel" (
            references = @InvertKernel.usda@
        )
        {
            # Chain: consume the grayscale output directly (no intermediate RenderVar).
            opaque inputs:Image.connect = <../GrayscaleKernel.outputs:LdrGrayscale>
            float inputs:strength = 1.0
        }
    }
}

Every type you can author is listed in Value-Input Types, together with what each one becomes on the GPU.

Two of them will catch you out. A quatf or quatd is written (w, x, y, z) in USDA and arrives as (x, y, z, w), so the identity quaternion authored (1, 0, 0, 0) reaches the GPU as (0, 0, 0, 1). And 64-bit integers bind on CUDA but not on Slang, where the Vulkan shaderInt64 feature is off.

On CUDA a numeric vector can pass by value or by pointer. A dtype constructor such as cuda.float3 passes every component by value to a matching float3 kernel parameter. cuda.array passes a pointer to the components instead; use it for matrices or when the kernel expects a pointer. A quaternion can use a matching four-component dtype constructor or cuda.array. On Slang the value goes into the parameter block and the shader declares it as a field of its struct.

Slang integer fields require matching element widths. A narrower integer value is rejected with an error naming the value and both widths, for scalar and vector fields as well as arrays and matrices. Use slang.int or slang.uint for a 32-bit integer field.

A token and an asset arrive as arrays. A token arrives as a null-terminated char array; CUDA can pass it directly with array. To pack token bytes into a Slang uint4 field, construct the four 32-bit words explicitly, as below; a byte array does not match that field’s element width. An asset arrives as the file’s raw bytes and is read as whichever dtype you name. The bytes must make whole elements of that dtype: an asset with bytes left over after its last whole element, or too short for one element, stops the node’s graph.

CUDA

cuda.array(inputs["mode"])              -- token -> const char*
cuda.array(inputs["lut"], cuda.uint)    -- asset -> const unsigned int*

Slang

local token = inputs["mode"].value
assert(#token < 16) -- reserve a byte for the null terminator
local words = { 0, 0, 0, 0 }
for i = 1, #token do
    local word = (i - 1) // 4 + 1
    words[word] = words[word] | (string.byte(token, i) << (((i - 1) % 4) * 8))
end
-- Inside the parameter block, matching a uint4 field:
slang.uint4(words[1], words[2], words[3], words[4])
slang.StructuredBuffer(slang.array(inputs["lut"], slang.uint))  -- asset -> StructuredBuffer<uint>

A USD array attribute, float[] say, is a buffer already and reaches a Slang shader as a read-only StructuredBuffer<float>. Refer to Upload Your Own Data.

A USD array attribute handed to cuda.array or slang.array:

• Without a dtype, it reaches the GPU as authored: a float3[] arrives as float3 elements.

• With a dtype of another number of components per element, its components are regrouped: a float[] passed as float2 arrives as pairs. The components must make whole elements of the dtype, or the node’s graph stops.

• A uchar[] array given a dtype is read the way an asset is: its bytes are taken as elements of that dtype.

• Any other array given a dtype of another element type has each value converted, then grouped as above: a double[] passed as float2 arrives as pairs of floats. The conversion runs in the launch script on every frame and counts every value against the script’s instruction and memory limits, so a large array stops the node’s graph.

A value converted to an integer element type:

• An integer that an unsigned type does not hold wraps: the element keeps the integer modulo 2 to the power of the type’s bits, so -1 from an int[] arrives as 4294967295 in a uint. A uint64 keeps all 64 bits of any integer. SPG logs the first integer that wraps at info level, once per node and input.

• An integer that a signed type does not hold stops the node’s graph.

• Any other value is truncated toward zero, and the node’s graph stops when the result does not fit the type, such as -1.0 from a double[] converted to uint, or when the value is not a finite number.

Author the array in the element type the GPU code reads.

An array may hold no elements. An array authored as [], written empty while the stage runs, or connected to an attribute that holds none reaches the node with no elements: .value is an empty list and sourceShape[1] is 0.

• cuda.array hands the kernel a null pointer, and a Slang buffer binds a placeholder the shader must not read, so pass the count to the GPU code and read no element past it.

• A parameter-block array field, or a place the launch script states, takes no elements and stays zero. A scalar, vector or matrix field cannot take an empty array, and the node’s graph stops.

• In a shader that does not describe itself, such as a .spv binary, SPG sizes a parameter block from the offsets the launch script states and the sizes of the values placed there. A block whose size comes out as zero, such as one that holds only an empty array at offset 0, stops the node’s graph with an error naming the first value placed there that cannot be laid out, else the empty array.

• cuda.array({}, dtype) and slang.array({}, dtype) make an array of no elements as well, and so does a uchar[] of no bytes given a dtype.

Passing an array on is free in Lua. An array, token or asset input handed to array without a dtype, with a dtype of its own element type, or, for a uchar[] array or an asset, with any dtype, reaches the GPU without its elements passing through the launch script. Reading .value builds a Lua table of all its elements, every frame; reading it alone does not change what the kernel receives.

Changes the script makes to that table before it hands the input to array reach the kernel, which still receives as many elements as the input has, taken in order from the table:

• An element the table lacks arrives as 0, and elements past the input’s length are not passed on.

• A fraction in an integer array arrives truncated toward zero, and true and false arrive as 1 and 0.

• A string, even one that holds a number, or a table where a number belongs stops the node’s graph.

• An integer written into an unsigned element that does not hold it wraps, so 300 in a uchar array arrives as 44. A fraction whose truncated value does not fit, such as 300.5 in a uchar array, stops the graph, and so does an integer that a signed element type does not hold.

• A uint64 value is the Lua integer with the same 64 bits, as Lua’s own string.unpack gives it. The largest Lua integer is 9223372036854775807, so a value above it reads as a negative number and converts to another element type as one. Any Lua integer written into a uint64 keeps its bits, so 0xFFFFFFFFFFFFFFFF and -1 both set every bit.

Verify It Worked

Change the value in the scene and re-run. If the output does not change, the value never reached the GPU. In the pipeline example, strength at 0.0 passes the image through unchanged and 1.0 fully inverts it, so the two ends of the range are unmistakable.

When It Goes Wrong

• Wrong value, or garbage: the argument order does not match. Refer to Wrong Pixels Rather Than No Pixels.

• On CUDA, check that the dtype constructor matches the kernel parameter’s element type and component count, or use cuda.array when the parameter is a pointer.

• On Slang, a wrong field count is an error and a wrong order only a warning. Refer to Bind Resources and Descriptor Spaces.

Related

The Lua Launch Script, The cuda Lua Table, The slang Lua Table