groupshared memory

groupsharedmem.hlsl

RWStructuredBuffer<float> g_data;
#define groupDim_x 128
groupshared float sdata[groupDim_x];
[numthreads(groupDim_x, 1, 1)]
void reduce1(uint3 threadIdx : SV_GroupThreadID,
uint3 groupIdx : SV_GroupID)
{
    // each thread loads one element from global to shared mem
    unsigned int tid = threadIdx.x;
    unsigned int i = groupIdx.x * groupDim_x + threadIdx.x;
    sdata[tid] = g_data[i];
    GroupMemoryBarrierWithGroupSync();
    // do reduction in shared mem
    for (unsigned int s = 1; s < groupDim_x; s *= 2)
    {
        if (tid % (2 * s) == 0)
        {
            sdata[tid] += sdata[tid + s];
        }
        GroupMemoryBarrierWithGroupSync();
    }
    // write result for this block to global mem
    if (tid == 0) g_data[groupIdx.x] = sdata[0];
}