CUDA版は？　1フレーム連続バッファ前提で，CUDAで高速化しやすい形に書き直してもよい．
__global__ void unpack_v210_to_planar_i16_kernel(
    const uint4* __restrict__ src,   // 16-byte aligned v210 blocks
    int blockCount,                  // number of 6-pixel blocks
    int16_t* __restrict__ dstY,
    int16_t* __restrict__ dstU,
    int16_t* __restrict__ dstV)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= blockCount) {
        return;
    }

    // 1 thread = 1 v210 block = 16 bytes = 4 words = 6 pixels
    uint4 p = src[i];

    uint32_t w0 = p.x;
    uint32_t w1 = p.y;
    uint32_t w2 = p.z;
    uint32_t w3 = p.w;

    // v210 unpack
    // w0: U0  Y0  V0
    // w1: Y1  U2  Y2
    // w2: V2  Y3  U4
    // w3: Y4  V4  Y5

    int16_t u0 = static_cast<int16_t>((w0 >>  0) & 0x3FF);
    int16_t y0 = static_cast<int16_t>((w0 >> 10) & 0x3FF);
    int16_t v0 = static_cast<int16_t>((w0 >> 20) & 0x3FF);

    int16_t y1 = static_cast<int16_t>((w1 >>  0) & 0x3FF);
    int16_t u2 = static_cast<int16_t>((w1 >> 10) & 0x3FF);
    int16_t y2 = static_cast<int16_t>((w1 >> 20) & 0x3FF);

    int16_t v2 = static_cast<int16_t>((w2 >>  0) & 0x3FF);
    int16_t y3 = static_cast<int16_t>((w2 >> 10) & 0x3FF);
    int16_t u4 = static_cast<int16_t>((w2 >> 20) & 0x3FF);

    int16_t y4 = static_cast<int16_t>((w3 >>  0) & 0x3FF);
    int16_t v4 = static_cast<int16_t>((w3 >> 10) & 0x3FF);
    int16_t y5 = static_cast<int16_t>((w3 >> 20) & 0x3FF);

    // flat planar index
    // 1 block = 6 Y samples, 3 U samples, 3 V samples
    int yBase  = i * 6;
    int uvBase = i * 3;

    dstY[yBase + 0] = y0;
    dstY[yBase + 1] = y1;
    dstY[yBase + 2] = y2;
    dstY[yBase + 3] = y3;
    dstY[yBase + 4] = y4;
    dstY[yBase + 5] = y5;

    dstU[uvBase + 0] = u0;
    dstU[uvBase + 1] = u2;
    dstU[uvBase + 2] = u4;

    dstV[uvBase + 0] = v0;
    dstV[uvBase + 1] = v2;
    dstV[uvBase + 2] = v4;
}
