typedef struct
{
    int a;
    int b;
    int c;
    int d;
} MY_TYPE_T;

typedef INTERLEAVED_T MY_TYPE_T[1024]; 

typedef int ARRAY_T[1024];

typedef struct
{
    ARRAY_T a;
    ARRAY_T b;
    ARRAY_T c;
    ARRAY_T d;
}NON_INTERLEAVED_T;

__host__ void add_test_non_interleaved_cpu(NON_INTERLEAVED_T * const host_dest_ptr, const NON_INTERLEAVED_T * const host_src_ptr, const int iter, const num_elements)
{
    for(int index = 0; index < num_elements; index++)
    {
        for(int i = 0; i < iter; i++)
        {
            host_dest_ptr->a[index] += host_src_ptr->a[index];
            host_dest_ptr->b[index] += host_src_ptr->b[index];
            host_dest_ptr->c[index] += host_src_ptr->c[index];
            host_dest_ptr->d[index] += host_src_ptr->d[index];
        }
    }
}

__host__ void add_test_interleaved_cpu(INTERLEAVED_T * const host_dest_ptr, const INTERLEAVED_T * const host_src_ptr, const int iter, const num_elements)
{
    for(int index = 0; index < num_elements; index++)
    {
        for(int i = 0; i < iter; i++)
        {
            host_dest_ptr[index].a += host_src_ptr[index].a;
            host_dest_ptr[index].b += host_src_ptr[index].b;
            host_dest_ptr[index].c += host_src_ptr[index].c;
            host_dest_ptr[index].d += host_src_ptr[index].d;
        }
    }
}

這兩個加和函數(shù)明顯類似，每個函數(shù)都對列表中的所有元素迭代iter次，從源數(shù)據(jù)結(jié)構(gòu)中讀取一個值，然后加和到目標數(shù)據(jù)結(jié)構(gòu)中。利用CPU系統(tǒng)時間統(tǒng)計這兩個函數(shù)分別運行的時間可以發(fā)現(xiàn)
“非交錯內(nèi)存訪問方式的執(zhí)行時間比交錯訪問方式的時間多出3~4倍?！?br> 這是意料之中的，因為在交錯訪問的例子中，CPU訪問元素a的同時會將結(jié)構(gòu)體中元素b、c和d讀入緩存中，使他們在相同的緩存行中。然而非交錯版本則需要對4個獨立的物理內(nèi)存進行訪問，也就是說存儲事務(wù)的數(shù)目為交錯版本的4倍，并且CPU使用的預(yù)讀策略不會起作用。
我們再看一下GPU版本的代碼：

__global__ void add_non_test_interleaved_kernel(NON_INTERLEAVED_T * const gpu_dest_ptr, const NON_INTERLEAVED_T * const gpu_src_ptr, const int iter, const int num_elements)
{
    const int tid = (blockIdx.x * blockDim.x) + threadIdx.x;
    if(tid < num_elements)
    {
        for(int i = 0; i < iter; i++)
        {
            gpu_dest_ptr->a[tid] += gpu_src_ptr->a[tid];
            gpu_dest_ptr->b[tid] += gpu_src_ptr->b[tid];
            gpu_dest_ptr->c[tid] += gpu_src_ptr->c[tid];
            gpu_dest_ptr->d[tid] += gpu_src_ptr->d[tid];
        }
    }
} 

__global__ void add_test_interleaved_kernel(INTERLEAVED_T * const gpu_dest_ptr, const INTERLEAVED_T * const gpu_src_ptr, const int iter, const num_elements)
{
    const int tid = (blockIdx.x * blockDim.x) + threadIdx.x;
    if(tid < num_elements)
    {
        for(int i = 0; i < iter; i++)
        {
            gpu_dest_ptr[tid].a += gpu_src_ptr[tid].a;
            gpu_dest_ptr[tid].b += gpu_src_ptr[tid].b;
            gpu_dest_ptr[tid].c += gpu_src_ptr[tid].c;
            gpu_dest_ptr[tid].d += gpu_src_ptr[tid].d;
        }
    }
}

   這兩個函數(shù)與CPU版本的類似，不過在GPU上，每個線程迭代iter計算一個元素。利用GPU系統(tǒng)統(tǒng)計分別統(tǒng)計這兩個函數(shù)運行的時間，可以發(fā)現(xiàn)與CPU版本不同，在GPU上
    “交錯內(nèi)存訪問方式的執(zhí)行時間比非交錯內(nèi)存訪問方式的時間多出3~4倍?！?

因為在GPU上，相比于交錯的訪問方式，非交錯訪問使我們得到了4個合并的訪問（所有線程訪問連續(xù)的對齊的內(nèi)存塊），保持全局內(nèi)存帶寬最優(yōu)。因此，在使用GPU全局內(nèi)存時，我們要注意連續(xù)合并的內(nèi)存訪問方式，從而擁有全局內(nèi)存帶寬最優(yōu)化。