LambdaReduction

Lambda Reduction

Reduction with the binary operator passed as a Func<float, float, float>. The lambda is captured at codegen time and inlined into the kernel — no runtime function-pointer indirection.

Source: src/6.Advanced/LambdaReduction/Program.cs

[Kernel]
public static void InnerReduce([Out] float[] result, [In] float[] input, int N,
                               float neutral, Func<float, float, float> reductor)
{
    var cache = new SharedMemoryAllocator<float>().allocate(blockDim.x);
    int tid = threadIdx.x + blockDim.x * blockIdx.x;
    int cacheIndex = threadIdx.x;

    float tmp = neutral;
    while (tid < N) { tmp = reductor(tmp, input[tid]); tid += blockDim.x * gridDim.x; }
    cache[cacheIndex] = tmp;

    CUDAIntrinsics.__syncthreads();
    int i = blockDim.x / 2;
    while (i != 0)
    {
        if (cacheIndex < i) cache[cacheIndex] = reductor(cache[cacheIndex], cache[cacheIndex + i]);
        CUDAIntrinsics.__syncthreads();
        i >>= 1;
    }
    // … atomic combine across blocks …
}

// Call sites:
InnerReduce(result, input, N, 0.0f,                    (a, b) => a + b);
InnerReduce(result, input, N, float.NegativeInfinity, (a, b) => Math.Max(a, b));

The lambda must:

• Capture no state beyond its parameters (no closure over outer locals — capture is unsupported in kernels).
• Be expressible at the call site so codegen can see the body. A Func<…> field set at runtime won't work.

Compare with InterfacesReduction (interface-based) and GenericReduction (intrinsic-based). The lambda form is the lightest at the call site but the strictest about composition.