     Test<2> (g_allocator);
     Test<15> (g_allocator);
 
+    Test<int32_t, 1, 64, 2, 2>(g_allocator);
+
     return 0;
 }


template <
        typename DataType,
        unsigned int ThreadsInBlock,
        unsigned int ItemsPerThread,
        typename ActionType>
__global__ void BlockShuffleTestKernel(
        DataType *data,
        ActionType action)
{
    typedef cub::BlockShuffle<DataType, ThreadsInBlock> BlockShuffle;

    __shared__ typename BlockShuffle::TempStorage temp_storage_shuffle;

    DataType thread_data[ItemsPerThread];

    data += threadIdx.x * ItemsPerThread;
    for (unsigned int item = 0; item < ItemsPerThread; item++)
    {
        thread_data[item] = data[item];

template <
        typename DataType,
        unsigned int ItemsPerThread,
        unsigned int BlockDimX, 
        unsigned int BlockDimY, 
        unsigned int BlockDimZ, 
        template<typename, unsigned int, unsigned int, unsigned int, unsigned int> class TestType>
void Test(unsigned int num_items,
          DataType *d_data,
          DataType *h_data)
{
    TestType<DataType, ItemsPerThread, BlockDimX, BlockDimY, BlockDimZ> test;

    Iota(num_items, d_data);
    BlockShuffleTest<DataType, ItemsPerThread, BlockDimX, BlockDimY, BlockDimZ>(d_data, test);
    AssertEquals(0, CheckResult(num_items, d_data, h_data, test));
}

