 template <
         typename DataType,
         unsigned int ItemsPerThread,
-        unsigned int ThreadsInBlock,
-        template<typename, unsigned int, unsigned int> class TestType>
+        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, ThreadsInBlock> test;
+    TestType<DataType, ItemsPerThread, BlockDimX, BlockDimY, BlockDimZ> test;
 
     Iota(num_items, d_data);
-    BlockShuffleTest<DataType, ItemsPerThread, ThreadsInBlock>(d_data, test);
+    BlockShuffleTest<DataType, ItemsPerThread, BlockDimX, BlockDimY, BlockDimZ>(d_data, test);
     AssertEquals(0, CheckResult(num_items, d_data, h_data, test));
 }
 

};

template <typename DataType,
          unsigned int ThreadsInBlock>
struct OffsetUpTest : public OffsetTestBase<DataType, ThreadsInBlock, -1 /* offset */>
{
    static __host__ bool check(const DataType *data, int i)
    {
        return UpTest<DataType, 1 /* ItemsPerThread */, ThreadsInBlock>::check (data, i);
    }
};

template<typename DataType,
         unsigned int ThreadsInBlock>
struct OffsetDownTest : public OffsetTestBase<DataType, ThreadsInBlock, 1 /* offset */>
{
    static __host__ bool check(const DataType *data, int i)
    {
        return DownTest<DataType, 1 /* ItemsPerThread */, ThreadsInBlock>::check (data, i);
    }
};

template<typename DataType,
         unsigned int ThreadsInBlock,
         unsigned int offset>
struct RotateTestBase
{
    static constexpr unsigned int ItemsPerThread = 1;

    __device__ void operator()(
            BlockShuffle<DataType, ThreadsInBlock> &block_shuffle,
            DataType (&thread_data)[ItemsPerThread]) const
    {
        block_shuffle.Rotate(thread_data[0], thread_data[0], offset);

template <
        typename DataType,
        unsigned int ItemsPerThread,
        unsigned int BlockDimX,
        unsigned int BlockDimY,
        unsigned int BlockDimZ>
struct DownTest
{
    __device__ void operator()(
            BlockShuffle<DataType, BlockDimX, BlockDimY, BlockDimZ> &block_shuffle,
            DataType (&thread_data)[ItemsPerThread]) const
    {
        block_shuffle.Down(thread_data, thread_data);
