10287人阅读
编程语言（17）
& & &最近初试cuda编程，作为一个新手，遇到了各种各样的问题，然后花费了大量时间解决这些匪夷所思的问题。为了避免后来人重蹈覆辙，现把自己遇到的问题总结如下。
（一）、cudaMalloc
& & &初次使用该函数，感觉没有什么困难，和c语言的malloc类似。但是在具体应用中却出了一个很难找的错误，花费了很多时间。该函数使用是需要注意的就是，它分配的内存空间单位是字节，所以需要我们在使用时用sizeof指定具体分配的变量类型，这样才能正确分配空间。例：
& & & & & & & & & & & & & &cudaMalloc((void**)&gpu_data,sizeof(float)*1024);
（二）、函数的执行位置
& & cuda程序的一大特色是程序的核心部分在GPU上执行，所以cuda函数就分为不同的类别：host、global、device三类。所以我们在编写函数时一定要分清楚当前正在编写的是哪类函数，可以调用什么库函数。
host函数：在CPU上调用，在CPU上执行，可以调用global函数，不能调用device函数；global函数：只能在host函数中调用，但是执行是在GPU上执行，例如cudaMalloc之类的内存操作库函数，可以调用device函数；device函数：只能在GPU上调用和执行，只能被global函数引用。
& & 关于函数类别容易出现的错误就是内存分配时CPU和GPU的混淆。我们只需要记住，在host函数中可以直接使用的内存都是CPU上的内存，GPU上的内存需要通过cudaMemcpy函数调用拷贝到CPU内存空间；在global和device函数中使用的内存都是在GPU内存空间，使用之前需要分配。
（三）、共享内存
& & &共享内存是提升程序性能很重要的一部分，能不能用好共享内存是是否掌握cuda编程的一个重要依据。在此只想强调一点：共享内存没有初始化！下面是自己写的一个数组求和程序，用到了共享内存：
__device__ int count=0;
__global__ static void sum(int* data_gpu,int* block_gpu,int *sum_gpu,int length)
extern __shared__ int blocksum[];
__shared__
const int tid=threadIdx.x;
const int bid=blockIdx.x;
blocksum[tid]=0;
for(int i=bid*THREAD_NUM+i&i+=BLOCK_NUM*THREAD_NUM)
blocksum[tid]+=data_gpu[i];
__syncthreads();
offset=THREAD_NUM/2;
while(offset&0)
if(tid&offset)
blocksum[tid]+=blocksum[tid+offset];
offset&&=1;
__syncthreads();
if(tid==0)
block_gpu[bid]=blocksum[0];
__threadfence();
int value=atomicAdd(&count,1);
islast=(value==gridDim.x-1);
__syncthreads();
if(islast)
if(tid==0)
for(int i=0;i&BLOCK_NUM;i++)
s+=block_gpu[i];
*sum_gpu=s;
& &特别注意第11八行代码，不对要访问的共享内存进行初始化将得不到正确的结果。
（四）、原子函数调用
& &在调用原子函数时，需要指定当前显卡的计算能力，否则会报错“atomic*** is undefined.”。 linux下解决方案是在编译源代码时为nvcc编译器指定一个计算能力的选项。例如计算能力时1.3，则可以添加参数：-arch sm_13，这样就可以顺利编译。
（五）、CUDA语法
& &很多参考书都介绍说CUDA采用的是C扩展语法，所以一开始我们很容易认为采用C语法就够了。但是这样也容易让我们陷入一个误区：只能是C语法，而不能是其他。其实CUDA是C和C++的混合体，有时候采用C++的语法会更便利：
for循环内可以定义变量，标准C语言不支持，所以我们可以直接用(for int i=0;i&i++)，这样的好处是可以节省一个寄存器；变量定义位置无限制，可以在任意位置定义变量；CUDA支持多态，所以我们可以定义多个名称相同，参数不同的函数，这个没有问题；有时多态可以用模版（template）来合并代码，达到简化编程的目的；
（六）、block和thread号的正确使用
& & 为了调度不同的线程，我们通常需要利用内置变量threadIdx和blockIdx作为循环中的增量。但是切记在循环内部要正确使用内置变量，两天debug的教训！下面是一个示例代码：
__global__ static void saliencefunc(float *peaks_gpu,int *index_gpu,float *saliencebins_gpu,int framenumber)
__shared__ float peaks[HALF_PEAK_NUM];
__shared__ int index[HALF_PEAK_NUM];
int tid=threadIdx.x;
int bid=blockIdx.x;
for(int i=i&i+=BLOCK_NUM)
if(tid&HALF_PEAK_NUM)
peaks[tid]=peaks_gpu[HALF_PEAK_NUM*i+tid];
index[tid]=index_gpu[HALF_PEAK_NUM*i+tid];
__syncthreads();
注意代码第十三和十四行的赋值操作HALF_PEAK_NUM*i+tid，笔者之前的写法是HALF_PEAK_NUM*bid+tid，结果花了两天的时间找问题，所以要正确使用，在可以替换的情况下就用i或者j这样的变量，尽量少用内置变量。
（七）、空间释放
& & 在GPU上分配的空间，在使用完成之后要及时释放。对于运行一次的程序，不释放空间没有什么大碍，毕竟程序结束空间自动会被释放掉。但是当程序不间断运行多次的时候，不释放空间会导致非常严重的GPU内存泄露。第一个问题是随着程序的运行，GPU内存耗尽，导致后续内存分配失败；第二个问题是，程序运行会越来越慢。所以我们一定要养成用完及时释放空间的习惯。
参考知识库
* 以上用户言论只代表其个人观点，不代表CSDN网站的观点或立场
访问：643919次
积分：5403
积分：5403
排名：第3923名
原创：67篇
转载：19篇
评论：227条
阅读：103466
阅读：53061C中 malloc()分配堆内存实际的大小 - 博客频道 - CSDN.NET
《西决》 —— 克里斯·保罗
分类：-[小西南]-
以前相关笔记：、。
C动态分配的实际大小
Figure1：内存中的堆内存空间
假设从《The &C &Programming &Language》中推测正确，从未经动态分配的堆内存呈现上图形式。不连续的堆内存以“链”的形式联系：Heap1 -& Heap2 -&Heap3 -&Heap4-&Heap1。笔记将构成“堆链”的每个堆内存（如Heap1）称为“堆块”。malloc()/free()将每个堆块看作由两部分构成：“Header”和“可用堆内存”。在Header中包含了“指向下一个堆内存块的指针”、“本堆块的大小”。这样malloc()/free()就能更好地管理堆。
2 堆内存分配
[1] mallco()分配机制
根据C中malloc(n)函数动态分配堆的机制：分配堆内存的时候就依序由低到高的地址搜索“堆链”中的堆块，搜索到“可用堆内存”满足n的堆块（如Heap1）为止。若Heap1的“可用堆内存”刚好满足n，则将Heap1从“堆链”中删除，同时重新组织各堆块形成新的“堆链”；若Heap1的“可用堆内存”大小大于n，则将malloc(n)申请到的“Header” + &可用堆内存&部分从Heap1中分裂，将剩余的Heap1堆内存块重新加入“堆链”中。经分裂后的堆内存也包含“Header”和“可用堆内存”两部分（如图Figure
2），然后将由malloc()分配得到的“可用堆内存”返回给用户。若某块堆内存空间比较大（如Heap1），能够满足较小内存的多次申请，那么由malloc(n)多次申请的堆内存块都是连续被分配给用户的（因为有Header，所以用户使用的堆地址不连续）。
为方便free()释放堆空间，经malloc(n)分配给用户的堆空间也隐含一个Header。如下图所示：
Figure2：malloc()分配的堆内存结构
由于Header的构成的内存对齐，C中malloc(n)函数分配的堆内存会大于等于Header + n。
[2] 程序验证
为了测试C语言malloc()函数所分配堆的实际大小(Header 大小和内存对齐)，编写了如下宏：
#define malloc_no_free(type, type_str)
printf(&%s\tsizeof(%s)=%d\t&, type_str, type_str, sizeof(type));\
type *pthis = NULL, *plast = NULL;
for(i = 0; i & 6; ++i){
if( NULL != (pthis = (type *)malloc(sizeof(type))) ){
printf(&%d\t&, (int)pthis - (int)plast);\
plast =/*a litte logic to p*/
/*free(pthis);*/
pthis = NULL;
printf(&\n\n&);
用宏写这一段函数式为了在C中能够传递任何数据类型的变量。如果传递给malloc()函数的参数较小，那么malloc()函数是极有可能在图中的Heap1对区域进行多次分配的，如此这段代码就能够测试出malloc()函数分配堆内存的实际大小。
将此宏在主函数中调用：
struct _sc{ };
struct _sic{ };
struct _sip{ struct _sip *p; };
struct _sdc{};
struct _s17{ char a[17]; };
int main()
malloc_no_free(char, &char&);
malloc_no_free(int, &int&);
malloc_no_free(double, &double&);
malloc_no_free(struct _sc, &_sc&);
malloc_no_free(struct _sic, &_sic&);
malloc_no_free(struct _sip, &_sip&);
malloc_no_free(struct _sdc, &_sdc&);
malloc_no_free(struct _s17, &_17&);
在codeblocks + GNU GCC compiler下编译通过并运行得到一下结果：
Figure3：程序运行结果
分析运行结果：
(1) 如果将宏中” free(pthis);”的注释去掉，则所有的地址差都为0。这恰好验证了free()堆地址的机制：假设将要被释放的堆内存的地址为p，free(p)将以p为首地址的堆内存块释放到堆链中的某个地址上，且此地址跟p值最邻近，这样就可以保证尽可能的使堆内存以大块的形式存在而不至于让堆内存成为碎片。
(2) 将” free(pthis);”注释
各结构体所占内存大小再一次体现了。在一大块堆之上用malloc()函数分配内存，真实得到的内存要比实际大。得到的真实内存块的大小为8（实为内存对齐大小值）的整数倍。当数据结构所占内存为16的倍数时，malloc()分配的内存会多出8个字节，这8个字节就是header。（将struct _s17内数组维数改为16倍数时可验证）。
从这里可以看出：Header所占堆内存大小为8字节，堆内存对齐也为8（会看内存对齐就不难观察出来）。
3 动态分配内存对齐
造就实际分配内存大于用户所需原因：
每次malloc(n)分配的堆内存由Header + n构成。Header造就的内存对齐。
[1] Header
在的第八章”The& UNIX System Interface”的第七节”Example – A Strorage Allocator”中用以下结构体描述了Header：
typedef long A
union header {
union header *
union &header用其成员s.ptr作为指向堆链表中下一个堆块，s.size保存此堆的大小。x本意是用于内存对齐。malloc()/free()就是通过以堆内存为存储空间隐藏于用户的Header结构体来合理地管理堆内存的。
[2] 内存对齐
《C圣经》中想用Align(long)作为内存对齐的值，但由于win32上sizeof(long) = 4，由于sizeof(union header) = sizeof(s) = 8，故而在malloc()返回的堆内存中存储数据的内存对齐数值为8。
所以，以上为每种数据类型所分配堆空间实际大小就得到了合理的解释（排除第一个为32 bytes）：
Figure4:C中实际的动态分配
所以在适当的明白malloc()的实际分配之后，在以后程序中的动态分配过程中就要利用这个过程避免堆内存的浪费：许多单独的堆内存块可以合并在一起分配，这样程序在运行过程中使用的堆内存会减少，有利于程序压缩空间。
4 malloc()分配总结
对于C中的malloc(n)分配，有以下进一步的结论：
实际分配的堆内存是Header + n结构。返回给用户的是n部分的首地址。由于内存对齐值8，实际分配的堆内存大于等于sizeof(Header) + n。
所以在编程过程中，需要用内存对齐的知识合理的让malloc()分配的内存变得小一些。
C Note Over。
排名：第559名
（3）（425）查看: 1734|回复: 0
专家解答：分配一个二维数组，cudaMallocPitch（）的参数该如何设置，pitch设置多少
论坛徽章:3
网友提问：我想分配一个5000*21的二维数组，浮点型的数据。请问&&cudaMallocPitch（）的参数该如何设置，尤其是pitch该设置多少？谢谢
专家解答：
A:cudaMallocPitch()中的pitch不是输入参数而是输出参数。pitch参数的地方，需要传入一个size_t*类型的变量指针。调用cudaMallocPitch()成功后，该变量中保存的是该块内存对齐后的pitch值。以后在对该数组的访问中，需要用这个pitch值修正访问地址。具体的例子详见CUDA 2.1 Programming Guide 4.5.2.3
以该问题中的数据为例，5000×21的浮点数组（假设是5000行x21列），那么正确的调用方法应该是：
& && && & float* d_
& && && & size_
& && && & cudaMallocPitch((void**)&d_matrix, &pitch, 21 * sizeof(float), 5000);
在G80系列GPU上，返回的pitch值是128，这意味着每行的数据元素是128byte，而不是21 × 4 ＝ 84 byte。
itpub.net All Right Reserved. 北京皓辰网域网络信息技术有限公司版权所有　　　　
　北京市公安局海淀分局网监中心备案编号：　广播电视节目制作经营许可证：编号（京）字第1149号