CUDAExample-0-cppOverload
2015-12-20 23:22
393 查看
官方提供的例程中有GPU和CPU两部分程序,这两部分程序完成相同的工作,其中,GPU部分是用共享内存完成,且提供了三种不同的方法,CPU部分同样提供了三种不同的计算方法,完成与GPU程序相同的运算,用于做对比。overload指的是一个块内sharememory超出了max shared memory per block的size。但是在cuda7.0,gtx980显卡下max shared memory per block为48k,当把一块内的共享内存提升到48k以上时,就会报错,提示share memory溢出。但是当没有超过share memory时,程序运行正确且对比结果为真。换句话说,overload问题已经通过自身软件进行了限制,不允许程序中出现overload问题,否则直接报错。
以上程序在每一部分都相同,且在kernel函数运行之前就进行显示,而且simple—kernel函数是复用的,区分复用哪一个kernel函数用
分辨。
其中,cudaFuncSetCacheConfig和cudaFuncGetAttributes的代码未给出,但是通过传入传出参数和运行代码可以知道,此段代码是用于设置和返回当前正在使用的kernel函数的状态,如shared size 、constant size、max threads per block、number of register等等,其中shared size是根据程序中设置的值来确定的,是个变量,在文档中有提及
Note that some function attributes such as
ref ::cudaFuncAttributes::maxThreadsPerBlock “maxThreadsPerBlock”
may vary based on the device that is currently being used
可以看出此部分代码虽然放到kernel函数之前,但是却能返回kernel中设置的信息。
测试程序
GPU部分
/* *share memory为THREAD—N*sizeof(int)的大小。 */ __global__ void simple_kernel(const int *pIn, int *pOut, int a) { __shared__ int sData[THREAD_N]; int tid = threadIdx.x + blockDim.x*blockIdx.x; sData[threadIdx.x] = pIn[tid]; __syncthreads(); pOut[tid] = sData[threadIdx.x]*a + tid;; } /* *share memory为THREAD—N*sizeof(int2)的大小。 *int2结构体类型,包含x,y两维度 */ __global__ void simple_kernel(const int2 *pIn, int *pOut, int a) { __shared__ int sData6[3]; int tid = threadIdx.x + blockDim.x*blockIdx.x; sData[threadIdx.x] = pIn[tid]; __syncthreads(); pOut[tid] = (sData[threadIdx.x].x + sData[threadIdx.x].y)*a + tid;; } /* *share memory为THREAD—N*sizeof(int)* 2的大小。 */ __global__ void simple_kernel(const int *pIn1, const int *pIn2, int *pOut, int a) { __shared__ int sData1[THREAD_N]; __shared__ int sData2[THREAD_N]; int tid = threadIdx.x + blockDim.x*blockIdx.x; sData1[threadIdx.x] = pIn1[tid]; sData2[threadIdx.x] = pIn2[tid]; __syncthreads(); pOut[tid] = (sData1[threadIdx.x] + sData2[threadIdx.x])*a + tid; }
CPU部分
bool check_func1(int *hInput, int *hOutput, int a) { for (int i = 0; i < N; ++i) { int cpuRes = hInput[i]*a + i; if (hOutput[i] != cpuRes) { return false; } } return true; } bool check_func2(int2 *hInput, int *hOutput, int a) { for (int i = 0; i < N; i++) { int cpuRes = (hInput[i].x + hInput[i].y)*a + i; //printf("hInput[i].x %d - hInput[i].y %d = %d\n", hInput[i].x , hInput[i].y, hInput[i].x - hInput[i].y); if (hOutput[i] != cpuRes) { return false; } } return true; } bool check_func3(int *hInput1, int *hInput2, int *hOutput, int a) { for (int i = 0; i < N; i++) { if (hOutput[i] != (hInput1[i] + hInput2[i])*a + i) { return false; } } return true; }
GPU状态提取部分
struct cudaFuncAttributes attr; // overload function 1 func1 = simple_kernel; memset(&attr, 0, sizeof(attr)); checkCudaErrors(cudaFuncSetCacheConfig(*func1, cudaFuncCachePreferShared)); checkCudaErrors(cudaFuncGetAttributes(&attr, *func1)); OUTPUT_ATTR(attr);
以上程序在每一部分都相同,且在kernel函数运行之前就进行显示,而且simple—kernel函数是复用的,区分复用哪一个kernel函数用
void (*func1)(const int *, int *, int); void (*func2)(const int2 *, int *, int); void (*func3)(const int *, const int *, int *, int);
分辨。
其中,cudaFuncSetCacheConfig和cudaFuncGetAttributes的代码未给出,但是通过传入传出参数和运行代码可以知道,此段代码是用于设置和返回当前正在使用的kernel函数的状态,如shared size 、constant size、max threads per block、number of register等等,其中shared size是根据程序中设置的值来确定的,是个变量,在文档中有提及
Note that some function attributes such as
ref ::cudaFuncAttributes::maxThreadsPerBlock “maxThreadsPerBlock”
may vary based on the device that is currently being used
可以看出此部分代码虽然放到kernel函数之前,但是却能返回kernel中设置的信息。
测试程序
由于GTX980每一块中的共享内存较多,线程只有1024,所以通过在每一块中额外定义和使用Shared memory来验证程序是否会返回false,但是当块中使用的shared memory超出限制时,就会在编译时报错。总结
通过以上程序以及测试程序可以分析得知overload问题已经在在此版本中通过编译前检查是否溢出来解决。相关文章推荐
- C#委托与C语言函数指针及函数指针数组
- more effinctive c++(operator+=)
- 基于C++高性能、跨平台日志模块的分析与实现
- 魔兽世界之二:装备(C++程序设计第5周)
- C++ 实现单向链表
- C++实现双向链表的创建,插入,修改,删除
- 继承与派生编程题1(C++程序设计第5周)
- C++临时变量什么时候销毁
- 一起talk C栗子吧(第七十八回:C语言实例--创建进程)
- C++学习笔记-容器
- 初步C++: g++: error: CreateProcess: No such file or directory
- windows下sublime编译调试c++的配置代码
- libc++和libstdc++
- C++模板类之理解编译器的编译模板过程
- 有关C语言中输入语句的问题
- C++中的typedef
- C++类型转换
- C++类型转换
- Effective C++ 笔记目录
- c++学习笔记