9. CUDA shared memory使用------GPU的革命
序言:明年就毕业了,下半年就要为以后的生活做打算。这半年,或许就是一个抉择的时候,又是到了一个要做选择的时候。或许是自己的危机意识比较强,一直都觉得自己做得不够好,还需要积累和学习。或许是知足常乐吧,从小山沟,能到香港,一步一步,自己都比较满足,只是心中一直抱着一个理想,坚持做一件事情,坚持想做点事情,踏踏实实,曾经失败过,曾经迷茫过,才学会了坚持,学会了坚毅,才体会了淡定和从容。人生路上,一路走,一路看,一路学,抱着感恩的心,帮助别人,就是帮助自己,未来的路才会更宽……
正文:书接上文《8. CUDA 内存使用 global 二------GPU的革命》 讲了global内存访问的时候,需要对齐的问题,只有在对齐的情况下才能保证global内存的高效访问。这一章节准备写一下shared memory的访问的问题,首先是讲一下shared的memory的两种使用方法,然后讲解一下shared memory的bank conflict的问题,这个是shared memory访问能否高效的问题所在;
Shared memory的常规使用:
1. 使用固定大小的数组:
/************************************************************************/
/* Example */
/************************************************************************/
__global__ void shared_memory_1(float* result, int num, float* table_1)
{
__shared__ float sh_data[THREAD_SIZE];
int idx = threadIdx.x;
float ret = 0.0f;
sh_data[idx] = table_1[idx];
for (int i = 0; i
{
ret += sh_data[idx %BANK_CONFLICT];
}
result[idx] = ret;
}
这里的sh_data就是固定大小的数组;
2. 使用动态分配的数组:
extern __shared__ char array[];
__global__ void shared_memory_1(float* result, int num, float* table_1, int shared_size)
{
float* sh_data = (float*)array; // 这里就让sh_data指向了shared memory的第一个地址,就可以动态分配空间
float* sh_data2 = (float*)&sh_data[shared_size]; // 这里的shared_size的大小为sh_data的大小;
int idx = threadIdx.x;
float ret = 0.0f;
sh_data[idx] = table_1[idx];
for (int i = 0; i
{
ret += sh_data[idx %BANK_CONFLICT];
}
result[idx] = ret;
}
这里是动态分配的空间,extern __shared__ char array[];指定了shared的第一个变量的地址,这里其实是指向shared memory空间地址;后面的动态分配float* sh_data = (float*)array;让sh_data指向array其实就是指向shared memory上的第一个地址;
后面的float* sh_data2 = (float*)&sh_data[shared_size];这里的sh_data2是指向的第一个sh_data的shared_size的地址,就是sh_data就是有了shared_size的动态分配的空间;
入下图:
3. 下面是讲解bank conflict
我们知道有每一个half-warp是16个thread,然后shared memory有16个bank,怎么分配这16个thread,分别到各自的bank去取shared memory,如果大家都到同一个bank取款,就会排队,这就造成了bank conflict,上面的代码可以用来验证一下bank conflict对代码性能造成的影响:
/************************************************************************/
/* Example */
/************************************************************************/
__global__ void shared_memory_1(float* result, int num, float* table_1)
{
__shared__ float sh_data[THREAD_SIZE];
int idx = threadIdx.x;
float ret = 0.0f;
sh_data[idx] = table_1[idx];
for (int i = 0; i
{
ret += sh_data[idx %BANK_CONFLICT];
}
result[idx] = ret;
}
// 1,2,3,4,5,6,7.....16
#define BANK_CONFLICT 16
这里的BANK_CONFLICT 定义为从1到16的大小,可以自己修改,来看看bank conflict对性能的影响;当BANK_CONFLICT为2的时候,就会通用有8个thread同时访问同一个bank,因为idx%2的取值只有2个0和1,所以16个都会访问bank0和bank1,以此类推,就可以测试整个的性能;
下面为示意图:
当然我们还可以利用16bank conflict,大家都访问同一个bank的同一个数据的时候,就可以形成一个broadcast,那样就会把数据同时广播给16个thread,这样就可以合理利用shared memory的broadcast的机会。
下面贴出代码,最好自己测试一下;
/********************************************************************
* shared_memory_test.cu
* This is a example of the CUDA program.
* Author: zhao.kaiyong(at)gmail.com
* http://blog.csdn.net/openhero
* http://www.comp.hkbu.edu.hk/~kyzhao/
*********************************************************************/
#include <stdio.h></stdio.h>
#include <stdlib.h></stdlib.h>
#include <cutil.h></cutil.h>
#include <cutil_inline.h></cutil_inline.h>
// 1,2,3,4,5,6,7.....16
#define BANK_CONFLICT 16
#define THREAD_SIZE 16
/************************************************************************/
/* static */
/************************************************************************/
__global__ void shared_memory_static(float* result, int num, float* table_1)
{
__shared__ float sh_data[THREAD_SIZE];
int idx = threadIdx.x;
float ret = 0.0f;
sh_data[idx] = table_1[idx];
for (int i = 0; i
{
ret += sh_data[idx%BANK_CONFLICT];
}
result[idx] = ret;
}
/************************************************************************/
/* dynamic */
/************************************************************************/
extern __shared__ char array[];
__global__ void shared_memory_dynamic(float* result, int num, float* table_1, int shared_size)
{
float* sh_data = (float*)array; // 这里就让sh_data指向了shared memory的第一个地址,就可以动态分配空间
float* sh_data2 = (float*)&sh_data[shared_size]; // 这里的shared_size的大小为sh_data的大小;
int idx = threadIdx.x;
float ret = 0.0f;
sh_data[idx] = table_1[idx];
for (int i = 0; i
{
ret += sh_data[idx%BANK_CONFLICT];
}
result[idx] = ret;
}
/************************************************************************/
/* Bank conflict */
/************************************************************************/
__global__ void shared_memory_bankconflict(float* result, int num, float* table_1)
{
__shared__ float sh_data[THREAD_SIZE];
int idx = threadIdx.x;
float ret = 0.0f;
sh_data[idx] = table_1[idx];
for (int i = 0; i
{
ret += sh_data[idx % BANK_CONFLICT];
}
result[idx] = ret;
}
/************************************************************************/
/* HelloCUDA */
/************************************************************************/
int main(int argc, char* argv[])
{
if ( cutCheckCmdLineFlag(argc, (const char**) argv, "device"))
{
cutilDeviceInit(argc, argv);
}else
{
int id = cutGetMaxGflopsDeviceId();
cudaSetDevice(id);
}
float *device_result = NULL;
float host_result[THREAD_SIZE] ={0};
CUDA_SAFE_CALL( cudaMalloc((void**) &device_result, sizeof(float) * THREAD_SIZE));
float *device_table_1 = NULL;
float host_table1[THREAD_SIZE] = {0};
for (int i = 0; i
{
host_table1[i] = rand()%RAND_MAX;
}
CUDA_SAFE_CALL( cudaMalloc((void**) &device_table_1, sizeof(float) * THREAD_SIZE));
CUDA_SAFE_CALL( cudaMemcpy(device_table_1, host_table1, sizeof(float) * THREAD_SIZE, cudaMemcpyHostToDevice));
unsigned int timer = 0;
CUT_SAFE_CALL( cutCreateTimer( &timer));
CUT_SAFE_CALL( cutStartTimer( timer));
shared_memory_static>>(device_result, 1000, device_table_1);
//shared_memory_dynamic>>(device_result, 1000, device_table_1, 16);
//shared_memory_bankconflict>>(device_result, 1000, device_table_1);
CUT_CHECK_ERROR("Kernel execution failed/n");
CUDA_SAFE_CALL( cudaMemcpy(host_result, device_result, sizeof(float) * THREAD_SIZE, cudaMemcpyDeviceToHost));
CUT_SAFE_CALL( cutStopTimer( timer));
printf("Processing time: %f (ms)/n", cutGetTimerValue( timer));
CUT_SAFE_CALL( cutDeleteTimer( timer));
for (int i = 0; i
{
printf("%f ", host_result[i]);
}
CUDA_SAFE_CALL( cudaFree(device_result));
CUDA_SAFE_CALL( cudaFree(device_table_1));
cutilExit(argc, argv);
}
这里只是一个简单的demo,大家可以测试一下。下一章节会将一些shared memory的更多的特性,更深入的讲解shared memory的一些隐藏的性质;
再在接下来的章节会讲一些constant和texture的使用;
写的内容一直都是文字比较多,代码比较少,其实学习的过程更重要的思想,实践的代码,最好是自己写,唯一可以学习的是思想,学习更重要的也是思想的交流,知识的传播,最好的是思想的传播,代码,方法,都是只是一些工具而已。但是工具的熟练层度,就得靠自己下来多练习。
分享到:
相关推荐
GPU求矩阵中最大值 sharedmemory共享内存
Total amount of shared memory per block: 49152 bytes Total number of registers available per block: 65536 Warp size: 32 Maximum number of threads per multiprocessor: 2048 Maximum number of ...
It provides several key abstractions – a hierarchy of thread blocks, shared memory, and barrier synchronization. This model has proven quite successful at programming multithreaded many core GPUs ...
Total amount of shared memory per block: 49152 bytes Total number of registers available per block: 65536 Warp size: 32 Maximum number of threads per multiprocessor: 2048 Maximum number of ...
Shared Memory Application Programming
矩阵乘法,利用shared memory matrixMul_Berkeley 4.7.1 矩阵乘法,利用register reduction 4.7.2 并行归约(缩减)程序 scan 5.1.2 Scan算法,例如计算前缀和 scanLargeArray 5.1.2 ...
矩阵乘法,利用shared memory matrixMul_Berkeley 4.7.1 矩阵乘法,利用register reduction 4.7.2 并行归约(缩减)程序 scan 5.1.2 Scan算法,例如计算前缀和 scanLargeArray 5.1.2 Scan...
3.2.4.1 可分享存储器(portable memory) 34 3.2.4.2 写结合存储器 34 3.2.4.3 被映射存储器 34 3.2.5 异步并发执行 35 3.2.5.1 主机和设备间异步执行 35 3.2.5.2 数据传输和内核执行重叠 36 3.2.5.3 并发内核执行 36...
在不同的并行硬件配置(CPU 线程、GPU 卸载和 CPU 协处理)上运行代码提供了使用常见计算材料科学工作负载的这些工具的基准。 将性能与串行基准进行比较将帮助您做出明智的决定,以决定哪些开发途径适合您的科学...
pyopencl:适用于Python的OpenCL集成,以及出色的功能