CUDA1.1-函数类型限定符与变量类型限定符

这部分来自于《CUDA_C_Programming_Guide.pdf》,看完《GPU高性能变成CUDA实战》的第四章,觉得这本书还是很好的,是一种循序渐进式的书,值得看,而不是工具书那种,适合入门,看完这章,觉得应该先简单的列下函数类型限定符,顺带列下变量类型限定符。知识是“积少成多”的。

ps;极力推荐使用编辑器之神-vim来写代码,正打算没事一点一点的使用这个神器,抛却其他编辑器,每天不需要学新东西,如果能够使用超过半年,我想有了熟悉感,学习其他的就不难了,这也是我极力倡导的:软件、算法、代码、数学等知识,不是一蹴而就的,可以每天稍微接触点,时间长了除却了陌生感,再次拾起的时候也会上手很快(那些有时间这么干的一群人,比如学生。切记贪多嚼不烂,主打几个方面,以1-3个方面为主,其他为辅,可是为辅不代表等用的时候再去接触,那时候肯定手忙脚乱一团糊,为辅可以不精通,可是最好需要了解,入门)。

一、函数类型限定符

函数类型限定符是用来指定这个函数的运行极其调用是在host(主机)还是device(设备)上运行的。主机:cpu及内存条;设备:gpu及显存。

1、__device__

这个限定符声明的函数:a、在设备上运行;b、只能被设备调用;

2、__global__

这个限定符声明的函数是可以作为核函数看待,该函数:a、在设备上运行;b、可以从主机上被调用;c、可以被计算能力为3.x的设备调用。

声明的函数的返回值必须是void类型。

任何调用这种函数的函数必须指定它的运行配置(在执行配置部分介绍,这里略)。

这种函数是异步的,也就是在设备完全完成它的运行之前就返回了。

3、__host__

这个限定符声明的函数:a、在主机上运行;b、只能被主机调用;

当函数前面没有任何函数类型限定符的时候,默认的就是这个限定符。

__global__和__host__不能一起声明同一个函数。

__device__和__host__可以一起声明同一个函数,这是为了让编译器编译出两个不同的版本,在Application Compatibility中的红__CUDA_ARCH__可以用来区分代码的路径是来自于主机还是设备:

__host__ __device__ func()
{
#if __CUDA_ARCH__ >= 300
// Device code path for compute capability 3.x
#elif __CUDA_ARCH__ >= 200
// Device code path for compute capability 2.x
#elif __CUDA_ARCH__ >= 100
// Device code path for compute capability 1.x
#elif !defined(__CUDA_ARCH__)
// Host code path
#endif
4、__noinline__ 和__forceinline__

当为计算能力为1.x的设备编译代码的时候,__device__函数默认是内联的。当为计算能力为2.x或者更高的设备编译代码的时候,__device__函数只有当编译器认为适当的时候才会是内联的。

__noinline__函数限定符可以用来作为编译器不将函数作为内联函数处理的标识。函数体必须在与被调用的文件中,(即调用的函数和函数的定义要在同一个文件中)。对于计算能力为1.x的设备来说,即使使用__noinline__函数限定符,编译器也不会执行带有指针参数和带有很长的参数列表的函数的。

__forceinline__函数限定符可以用来强制编译器将函数作为内联函数看待。

二、变量类型限定符

变量类型限定符用来指定变量在设备上的存储位置(因为主机上不需要这个限定符)。

在设备代码中没有任何__device__、__shared__、 __constant__定符的自动变量声明通常是存储在一个寄存器上的。但是在某些情况下,编译器会选择将它存放在局部存储器(会有不良的执行结果)上。

1、__device__

用来声明变量,使其存储在设备上。

大部分情况来说,下面的两个变量限定符中任何的一个通常都会与__device__一起使用,用来指定存储变量的空间,如果没有那两个变量限定符一起使用的话,这个变量:

    存储在全局存储空间中;

    有着与应用程序一样长的生命周期;

    可以被grid中所有的线程和从运行时库中的主机函数所访问。其中的函数例如:cudaGetSymbolAddress() / cudaGetSymbolSize() /cudaMemcpyToSymbol() / cudaMemcpyFromSymbol().

    可以被额外的__managed__限定符所限定,这样的一个变量可以直接被主机代码所引用,例如:它的地址可以直接在一个主机函数中被执行、被读、被写。为了方便 __managed__就表示 __managed__   __device__,当使用__managed__限定符的时候,__device__限定符可以省略。

2、__constant__

这个限定符是作为与__device__限定符可选的额外限定符,它表示的变量:

   存储在常量存储空间中;

   有着与应用程序一样长的生命周期;

   可以被grid中所有的线程和来自运行时库中的主机函数所访问。其中的函数例如:cudaGetSymbolAddress() / cudaGetSymbolSize() /cudaMemcpyToSymbol() / cudaMemcpyFromSymbol().

3、__shared__

  这个限定符也是作为与__device__限定符可选的额外限定符,它声明的变量:

    存储在一个线程块的共享存储器中

    有着与这个块一样的生命周期;

    只能被这个块中的所有线程所访问。

当在共享存储器中声明一个变量作为外部数组,例如:

extern __shared__ float shared[];

这个数组的大小是在运行的时候决定的。所有以这种形式定义的变量,开始与存储器中相同的地址,所以数组中变量的布局必须是通过offsets所显式管理的。例如,如果你想在动态分配共享存储器中得到如下的等价情况:

short array0[128];
float array1[64];
int array2[256];

,你应该按照下面的方式来声明和初始化这些数组:

extern __shared__ float array[];
__device__ void func() // __device__ or __global__ function
{
short* array0 = (short*)array;
float* array1 = (float*)&array0[128];
int* array2 = (int*)&array1[64];
}
注意指针需要预先分配为所指定类型的大小,所以下面的代码不会起效果,因为array1没有被指派成4个字节:

extern __shared__ float array[];
__device__ void func() // __device__ or __global__ function
{
short* array0 = (short*)array;
float* array1 = (float*)&array0[127];
}
4、__restrict__

nvcc支持通过__restrict__关键字来指定受限指针。

在C99标准中引入的受限指针可以缓解存在于C类型语言中的别名问题,这个问题就是:禁止所有来自代码重排序到公共子表达消除中的所有优化。
下面就是一个别名问题,如果使用受限指针,那么就有助于编译器减少指令的数量:

void foo(const float* a,
const float* b,
float* c)
{
c[0] = a[0] * b[0];
c[1] = a[0] * b[0];
c[2] = a[0] * b[0] * a[1];
c[3] = a[0] * a[1];
c[4] = a[0] * b[0];
c[5] = b[0];
...
}
在C类型语言中,指针a,b,c也许会出现别名,所以任何通过c 的写操作会修改a 或 b  中的元素。也就是说为了保证功能的正确性,编译器不能将a[0]和b[0]装载到寄存器中,相乘它们和将结果存储到c[0]和c[1]中,因为当a[0]与c[0]的存储在同一个地方的时候,结果不同于这个抽象执行模型。所以编译器就没有公共子表达的优势了。同样的,编译器不能进行重排序c[4]的计算结果到c[0]和c[1]计算的接近位置,因为对c[3]的写操作可能会改变c[4]计算的输入。

通过设定a,b,c为受限指针,程序员可以对编译器说指针事实上不需要别名,也就是说通过对c的写操作不会重写a和b中的元素。改变后的函数原型为:

void foo(const float* __restrict__ a,const float* __restrict__ b,float* __restrict__ c);
注意到所有的指针参数需要被设置成受限的,增加这个关键字,编译器现在可以重排序并且使用公共子表达消除了,同时保证在抽象执行模型的功能统一性:void foo(const float* __restrict__ a,
const float* __restrict__ b,
float* __restrict__ c)
{
float t0 = a[0];
float t1 = b[0];
float t2 = t0 * t2;
float t3 = a[1];
c[0] = t2;
c[1] = t2;
c[4] = t2;
c[2] = t2 * t3;
c[3] = t0 * t3;
c[5] = t1;
...
}

这里的效果是减少了存储器的访问次数和减少了计算的次数。这里是通过“cached”的装载和公共子表达来增加了寄存器的压力从而平衡得到的。

因为在许多CUDA代码中寄存器压力是一个关键性问题,使用受限指针,减少了占用,所以会在CUDA代码中有着负面性能效果。

notes:在《大规模并行处理器编程实战》第44页说每个块总有有512个线程(不同显卡不同,780ti的是1024);而在《gpu高性能编程cuda实战》中的33页说,启动线程块数组,数组的每一维最大不能超过65535(也是不确定的),而不论是块还是线程数量,都是具有dim3类型的,也就是有x,y,z三个维度。

kernel<<<nblock,nthread>>>(args);这是核函数,告诉运行时创建的函数会有多少个块,每个块有多少个线程。一个核函数就代表一个网格。

上面的意思就是在一个网格中可以有65535×65535×65535个块,每个块最多有8×8×8个线程。(这个说法是错的,在使用matlab的gpuDevice函数返回的结果,其实cuda自己也有查看函数,不过可以从matlab的gpuDevice结果上看,780ti显卡每个块最大线程数1024,但是三个方向上分别最大值不能超过【1024,1024,64】,也就是说我们给定一个块的并发线程总的不能超过1024,意思就是达不到1024×1024×64的结果,而每个网格的最大块数为【2.1475e+09,65535,65535】,所以这些数据不同显卡不同,得及时查证才是)


posted @ 2015-01-05 19:21  仙守  阅读(2425)  评论(0编辑  收藏  举报