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;
}
}
}
這兩個(gè)加和函數(shù)明顯類似纵装，每個(gè)函數(shù)都對(duì)列表中的所有元素迭代iter次，從源數(shù)據(jù)結(jié)構(gòu)中讀取一個(gè)值据某，然后加和到目標(biāo)數(shù)據(jù)結(jié)構(gòu)中橡娄。利用CPU系統(tǒng)時(shí)間統(tǒng)計(jì)這兩個(gè)函數(shù)分別運(yùn)行的時(shí)間可以發(fā)現(xiàn)
“非交錯(cuò)內(nèi)存訪問(wèn)方式的執(zhí)行時(shí)間比交錯(cuò)訪問(wèn)方式的時(shí)間多出3~4倍⊙⒆眩”
這是意料之中的挽唉，因?yàn)樵诮诲e(cuò)訪問(wèn)的例子中滤祖，CPU訪問(wèn)元素a的同時(shí)會(huì)將結(jié)構(gòu)體中元素b、c和d讀入緩存中瓶籽，使他們?cè)谙嗤木彺嫘兄薪惩Ｈ欢墙诲e(cuò)版本則需要對(duì)4個(gè)獨(dú)立的物理內(nèi)存進(jìn)行訪問(wèn)，也就是說(shuō)存儲(chǔ)事務(wù)的數(shù)目為交錯(cuò)版本的4倍塑顺，并且CPU使用的預(yù)讀策略不會(huì)起作用汤求。
我們?cè)倏匆幌翯PU版本的代碼：
__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;
}
}
}
這兩個(gè)函數(shù)與CPU版本的類似，不過(guò)在GPU上严拒，每個(gè)線程迭代iter計(jì)算一個(gè)元素扬绪。利用GPU系統(tǒng)統(tǒng)計(jì)分別統(tǒng)計(jì)這兩個(gè)函數(shù)運(yùn)行的時(shí)間，可以發(fā)現(xiàn)與CPU版本不同裤唠，在GPU上
“交錯(cuò)內(nèi)存訪問(wèn)方式的執(zhí)行時(shí)間比非交錯(cuò)內(nèi)存訪問(wèn)方式的時(shí)間多出3~4倍挤牛。”
因?yàn)樵贕PU上种蘸，相比于交錯(cuò)的訪問(wèn)方式墓赴，非交錯(cuò)訪問(wèn)使我們得到了4個(gè)合并的訪問(wèn)（所有線程訪問(wèn)連續(xù)的對(duì)齊的內(nèi)存塊），保持全局內(nèi)存帶寬最優(yōu)航瞭。因此诫硕，在使用GPU全局內(nèi)存時(shí)，我們要注意連續(xù)合并的內(nèi)存訪問(wèn)方式沧奴，從而擁有全局內(nèi)存帶寬最優(yōu)化。
