Basic Ripple Effect Kernel

A straightforward example demonstrating fundamental OpenCL kernel operations. This kernel reads point positions, applies a sine-wave displacement based on distance from origin, and writes the modified positions back.

kernel void rippleEffect(
    global float *positionData,
    int positionCount,
    float waveAmplitude,
    float wavePeriod,
    float wavePhase
) {
    int pointId = get_global_id(0);
    if (pointId >= positionCount)
        return;
    
    float3 currentPos = vload3(pointId, positionData);
    float distanceFromOrigin = length(currentPos);
    currentPos.y += waveAmplitude * sin(distanceFromOrigin / wavePeriod + wavePhase);
    
    vstore3(currentPos, pointId, positionData);
}

This produces a convincing ripple animation effect suitable for water surface simulations or abstract motion graphics.

3D Volume Blur Implementation

When processing volumetric data in Houdini's OpenCL context, a multi-pass approach often proves more effective than attempting all operations in a single kernel. The following example demonstrates a separable blur operation using scratch buffer intermediate storage.

kernel void blur3D(
    int strideX,
    int strideY,
    int strideZ,
    int baseOffset,
    int resolutionX,
    int resolutionY,
    int resolutionZ,
    global float *inputField,
    global float *scratchField
#ifdef ENABLE_MASK
    , global float *maskField
#endif
) {
    int x = get_global_id(0);
    int y = get_global_id(1);
    int z = get_global_id(2);
    int voxelIndex = baseOffset + strideX * x + strideY * y + strideZ * z;

#ifdef ENABLE_MASK
    if (maskField[voxelIndex] < 0.5) {
        scratchField[voxelIndex] = inputField[voxelIndex];
        return;
    }
#endif

    float accumulatedValue = 0.0;
    for (int offsetX = -1; offsetX <= 1; offsetX++) {
        for (int offsetY = -1; offsetY <= 1; offsetY++) {
            for (int offsetZ = -1; offsetZ <= 1; offsetZ++) {
                int clampedX = clamp(x + offsetX, 0, resolutionX - 1);
                int clampedY = clamp(y + offsetY, 0, resolutionY - 1);
                int clampedZ = clamp(z + offsetZ, 0, resolutionZ - 1);
                
                int neighborIndex = baseOffset 
                    + strideX * clampedX 
                    + strideY * clampedY 
                    + strideZ * clampedZ;
                
                accumulatedValue += inputField[neighborIndex];
            }
        }
    }
    
    scratchField[voxelIndex] = accumulatedValue / 27.0;
}

kernel void copyBack(
    int strideX,
    int strideY,
    int strideZ,
    int baseOffset,
    int resolutionX,
    int resolutionY,
    int resolutionZ,
    global float *outputField,
    global float *scratchField
) {
    int x = get_global_id(0);
    int y = get_global_id(1);
    int z = get_global_id(2);
    int voxelIndex = baseOffset + strideX * x + strideY * y + strideZ * z;
    
    outputField[voxelIndex] = scratchField[voxelIndex];
}

The blur kernel iterates through all 27 neighboring voxels (3×3×3 kernel), accumulates their values, and computes the mean. This box blur approach provides smooth results while maintaining computational efficiency.

The second kernel handles result transfer from scratch buffer back to the original field, completing the two-pass workflow.

OpenCL Type Strictness Pitfalls

OpenCL enforces stricter type rules than many high-level languages expect. Implicit type conversion can introduce subtle bugs that produce unexpected results without triggering compilation errors.

// PROBLEMATIC: Integer literal causes type inference issue
float signedDistance = abs(worldPos.x) + abs(worldPos.y) + abs(worldPos.z) - 1;

// CORRECTED: Explicit floating-point literal ensures float result
float signedDistance = abs(worldPos.x) + abs(worldPos.y) + abs(worldPos.z) - 1.0;

The problematic line appears syntactically correct but fails becuase the integer literal 1 causes the entire expression to evaluate as an integer before assignment to a float variable. Modern compilers may handle this more gracefully, but explicitly using floating-point literals remains the safest practice.

Volume Color Modification with Spatial Transformations

This advanced example demonstrates mapping point positions into volumetric data space and modifying point colors based on density values. It showcases essential techniques for hybrid geometry-volume workflows.

kernel void modifyColors(
    int pointCount,
    global float *positionData,
    int colorCount,
    global float *colorData,
    int densityStrideX,
    int densityStrideY,
    int densityStrideZ,
    int densityBaseOffset,
    int densityResolutionX,
    int densityResolutionY,
    int densityResolutionZ,
    float16 worldToVoxelMatrix,
    global float *densityField
) {
    int pointId = get_global_id(0);
    if (pointId >= colorCount)
        return;
    
    float3 worldPosition = vload3(pointId, positionData);
    
    // Apply 4x4 transformation matrix (treating position as homogeneous coordinate)
    float4 voxelPosition = worldPosition.x * worldToVoxelMatrix.lo.lo +
                          worldPosition.y * worldToVoxelMatrix.lo.hi +
                          worldPosition.z * worldToVoxelMatrix.hi.lo +
                          worldToVoxelMatrix.hi.hi;
    
    // Convert to discrete voxel indices
    int3 voxelCoords;
    voxelCoords.x = clamp((int)(floor(voxelPosition.x)), 0, densityResolutionX - 1);
    voxelCoords.y = clamp((int)(floor(voxelPosition.y)), 0, densityResolutionY - 1);
    voxelCoords.z = clamp((int)(floor(voxelPosition.z)), 0, densityResolutionZ - 1);
    
    float3 pointColor = {1.0, 1.0, 1.0};
    
    // Sample density at corresponding voxel
    float voxelDensity = densityField[densityBaseOffset +
                                      densityStrideX * voxelCoords.x +
                                      densityStrideY * voxelCoords.y +
                                      densityStrideZ * voxelCoords.z];
    
    if (voxelDensity > 0.5)
        pointColor.y = 0.0;
    
    vstore3(pointColor, pointId, colorData);
}

Matrix Decomposition for OpenCL Compatibility

OpenCL handles 4×4 matrices as composite types. When multiplying a homogeneous coordinate vector by a transformation matrix, decompose the 16-element matrix into four 4-element components:

float16 transformMatrix;
// Matrix layout: 4 rows × 4 elements each
float4 row0 = transformMatrix.lo.lo;  // First row
float4 row1 = transformMatrix.lo.hi;  // Second row
float4 row2 = transformMatrix.hi.lo;  // Third row
float4 row3 = transformMatrix.hi.hi;  // Fourth row (translation + w component)

// Vector-matrix multiplication
float4 result;
result.x = dot((float4)(pos.x, pos.y, pos.z, 1.0), row0);
result.y = dot((float4)(pos.x, pos.y, pos.z, 1.0), row1);
result.z = dot((float4)(pos.x, pos.y, pos.z, 1.0), row2);
result.w = dot((float4)(pos.x, pos.y, pos.z, 1.0), row3);

This decomposition approach provides explicit control over matrix operations without relying on built-in matrix multiplication that may not be available in all OpenCL implementations.

Understanding Stride Parameters and Memory Padding

When working with volumetric data in Houdini's OpenCL environment, the stride parameters (density_stride_x, density_stride_y, density_stride_z) encode critical information about memory layout. These values account for both voxel spacing and boundary padding.

For a volume with resolution [100, 95, 110], examining the stride values reveals the padding mechanism:

Houdini adds a one-voxel border around volumetric data to prevent boundary access violations during filtering operations. This padding appears on all six faces of the volume, effectively transforming dimensions according to:

• padded_x = original_x + 2
• padded_y = original_y + 2
• padded_z = original_z + 2

The stride values enable direct memory access by encoding the offset required to advance one voxel in each dimension within the flattened 1D buffer representation.

Force Alignment Considerations

When configuring OpenCL kernels for volumetric processing, the Force Alignment option significantly affects kernel interface complexity:

Enabled (default): Houdini automatically infers resolution and transformation parameters, providing a simplified kernel signature with only the essential data pointers.

Disabled: Explicit stride and resolution parameters become necessary, offering finer control over memory access patterns at the cost of additional parameter complexity.

This trade-off becomes particularly relevant when processing multiple volumes with different resolutions or transformations, where manual parameter management enables optimized memory access patterns that the automatic alignment system cannot achieve.