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.

It becomes an ordinary kernel parameter, declared in the signature like any other:

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

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.
//
// Value-inputs are not passed as arguments like in CUDA. They 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.dispatch({
        -- 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 value on the shader definition is a default; a shader instance in a scene may state its own.

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 vector does not pass by value. A float3, a matrix or a quaternion is bound with cuda.array, and the kernel parameter is a pointer to its components. Passed with a dtype constructor instead, it reaches the kernel with only its first component intact. On Slang the value goes into the parameter block and the shader declares it as an ordinary field.

A token and an asset are arrays, not scalars, so both are bound with array rather than with a dtype constructor. A token arrives as a null-terminated char array and needs no dtype; an asset arrives as the file’s raw bytes and is read as whichever dtype you name.

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

slang.array(inputs["mode"])                             -- token, inside the parameter block
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.

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, the first component of a vector is right and the rest is noise: it was passed with a dtype constructor rather than with cuda.array.

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