CUDA中文教程03之心得体会

    技术2022-05-12  11

     

    定义在GPU上的变量:

    1、使用关键字__device__ __local__  int X,则意味着该变量是定义在thread中的,它的生存周期跟它所在的thread一致。实际上定义为__local__ 的变量会存在global memory 中,所以速度也会很慢,一般不采用__local__关键字定义变量。关键字缺省情况下的变量是存在register中的,速度比存在global memory中快,只有当register存满了变量之后,系统才会自动把变量定义为__local__的,所以不要随意的采用__local__关键字,这其实是下下策。

    2、使用关键字__device__ __shared__  int X,则意味着该变量是定义在block中的,它的生存周期跟它所在的block一致,并且为该block里的512个thread共享,都可以访问到这个变量。

    3、如果缺省了第二个关键字,即__device__              int X,则是定义在grid中的,不仅在GPU中可见,CPU也可见。

    4、使用关键字__device__ __constant__  int X,则意味着该变量是定义在grid中的,是一常量,在run的过程中不会改变其值,且GPU及CPU均可见,在CPU中可见,是因为CPU要把该值传入GPU中。

    *使用这些关键字时,如果是用了__local__,__shared__,__constant__,则不必要在前面加__device__,系统就会知道定义的是GPU上的变量,如果是定义了global memory 的变量,则需写__device__关键字即可。

    *自动变量,即没有任何限定词的,会自动的放到register中,除了数组,数组会存在local memory中,所以当声明数组时,必须存到thread中去run。

    怎样选择关键字呢?

    第一步,考虑该关键字是否被CPU可见:“是”,进入第二步;“不是”,进入第三步。

    第二步,如果要被CPU可见,则选择关键字__global__或者__constant__。而且在声明变量时,必须写在所有函数体外,保证全局性。

    第三步,如果不被CPU可见,则选择关键字__shared__或__local__或者缺省不写关键字(存在register中),这时变量声明必须在kernel 函数中。

    shared memory 是一个很重要的概念,因为如果我们每次都去access global memory的话,就要很久的时间,所以我们要提高速度,就得想怎样换到shared memory中去。这里我们采用的是tile data ,即把数据切片的方法,变成一个个subset,使其刚好满足shared memory的大小,处理完数据之后再从shared memory传到global memory去。

    *指针只能指向定义在global memory中的函数或变量。


    最新回复(0)