     static_array<int [3][3], memory_space::local, layout::rmo, align<1024, 2>> A;
 
     printf("Host\n");
-    test(A);
+    test_single(A);
 
     printf("Device\n");
-    test_kernel<<<1, 1>>>();
+    test_kernel_single<<<1, 1>>>();
     cudaDeviceSynchronize();
 
     printf("Device __shared__\n");
-    test_kernel_shared<<<1, dim3(3, 3)>>>();
+    test_kernel_shared<<<1, dim3(BLOCK_X, BLOCK_Y)>>>();
     cudaDeviceSynchronize();
 
     return 0;

}

__global__
void test_kernel()
{
    static_array<int [3][3]> A;

    test(A);
}

__global__
void test_kernel_shared()
{
    static_array<int [3][3], memory_space::shared, layout::rmo, align<4>> A;

    bool first = threadIdx.y == 0 && threadIdx.x == 0;

    if (threadIdx.y == 1 && threadIdx.x == 1)
        A(1, 1) = 3;
    else
        A(threadIdx.y, threadIdx.x) = 1;

    __syncthreads();



template <typename T>
__host__ __device__
void test_single(T &A)
{
    for (unsigned i = 0; i < A.template dim<0>(); ++i) {
        for (unsigned j = 0; j < A.template dim<1>(); ++j) {
