GPU/CUDA的使用日志-可能适合初学者

本人在10年初的时候，研究过一段时间。当时记录的一些东西，也不知道有没有用。

10.2.1 一些初学的问题

Q1：dim3 block即为分块的时候，会不会有问题？

ThreadIdx Idx是index的简写。

Q2：在Block中分配Thread的时候，如果不是16一组，例如分成9一组会不会有问题

不会，但是整个线程是按照Warp来进行的，一个warp有32个线程，但是，在实际的运算中，实际上只有16个半warp在运行，所以，即使分成9个thread，还是会16组在一起运行。同理，我们也可以看到GPU不是绝对意义上的同步，也就是说，看到的所有的thread的运行其实是分成16一组一组的线程，然后这些16一组的线程才是绝对意义上的同步。所以说，数据量大执行的时间也会加长，不是我们想象的那种，只用了一个线程工作所用的时间。

 

 

Q3：block中最多分配512个Thread，一个Grid中block的数量有没有最大的限制

最多有512×512×512×64threads

一个grid里边，可以有512×512×64个blocks，而一个block中最多有512个线程。

 

Q4：纹理

纹理是把Global的内存变为Thread的Local内存;其实内存就在Global内存中放置,只不过用Texture来格式化它，这种机制，使存取Global中的数据速度更快。这里猜想，在Texture中，寻址不是在整个显存编址中进行，而是在Texture的相对地址中进行的，所以，每次都是在寻找相对地址，所以，速度会快。

    但是按照纪文的说法，还是没有__shared__格式快。

 

Q5：存储体的类型和应用

{

 __device__ float filter[N];

 __global__ void  convolve (float*image)  {

 __shared__ float region[M];

  ...

 region[threadIdx] = image[i];

 __syncthreads() 

  ...

 image[j] = result;

}

 

问题4: 如何使用这些内存

一个Block内最多可以有512个Thread,一个grid内最多可以是512×512×64个Block

一个block内有8192个寄存器，有16kB个SharedMemory,如果申请过多SharedMemory

会自动放入LocalMemory内，Local处于Global Memory中，存取很慢。

 

 

内存的管理：

1-    变量关键字 寄存器（没有关键字） 一个block内有8192个寄存器。只有线程可以访问。

2-   变量关键字 __shared__ 一个block内几个线程共享，最多可以申请16kByte。Block中的线程可以共享。

3-    变量关键字  一共显存512M，但运算的过程中,可以把其变成纹理Texture来判断，还是没有__shared__内存存取速度快。 申请的形式cudaMalloc(void* pointer, size_t size); cudaFree(void* pointer)。所有线程都可以访问。

4-    变量关键字 __constant__内存，是只读类型，也为16kBytes，读取速度快，例如可以把lut放入其中。所有线程都可以访问。

 

函数的申明：

1-       __host__ 表示在CPU程序中可以调用，__device__表示GPU的函数中可以调用(__global__函数表示GPU的内核函数)

这里要注意：__device__在变量申明中和函数申明中都是可以用的。

 

Q6: 内核函数的使用

 

KernelFunc<<< DimGrid, DimBlock,SharedMemBytes,  (事件流/数据流)>>>(...);

//没有看到程序中用到，而且这个内存只有在block中的线程才共享？

第一个参数Grid的维数的dim3结构；

第二个参数Block的维数的dim3结构；

第三个参数是SharedMemory的数量，如果这里有值，那么在内部申请__shared类型内存，就可以写成 extern__shared__ float a[]的形式，表示动态外部申请；如果没有第三个参数，__shared__就必须__shared__ float a[89]形式出现；

第四个参数现在还不知道如何使用；

 

Q7：运算的过程中都是单精度浮点型，为什么在申请内存的过程中也有unsigend short int的形式的？

最大的计算的单位为float类型，也就是sizeof(float)，因为int也是4，所以也可以申请，但是不能申请double的形式的内存单元。

 

Q8: GPU线程间的通讯？

Dividemonolithic thread array into multiple blocks

–  Threads within a block cooperate via shared memory, atomicoperationsand barrier synchronization；

–  Threads in different blocks cannot cooperate;

Q9:CPU必须在一个线程中初始化，且在同一个线程中调用

Q10：要是把所有的函数都写成类，然后这些类在生成对象的过程中，都先申请地址，会不会有问题？

 

10.2.2 相关硬件结构和理解

10.2.2.1 __global__内存的读写存取 memory bandwidth

 

一、方法一

因为带宽和速度的问题，一个典型的做法是将设备中的数据拷贝到__shared__内存中，然后进行计算。

具体的步骤:

1.把global中的数据拷贝到__shared__中。

2.同步，保证所有的__shared__中的数据都是安全的。

3.计算

4.同步计算的结果。

5.得出结论。

在拷贝到__shared__的时候，也是需要注意几点，

比如，对于变量CUDA会生成一个单独的拷贝的指令，比如

__global__ void Func(float* d_in,float* d_out)

{

Const int index = threadIdx.x + blockInx.x * BLOCK_DIM.x;

Float x = d_in[index * 3];

Float y = d_in[index * 3+ 1];

Float z = d_in[index * 3 + 2];

-- 此时threadIdx.x在读取显存中的数据时，会生成3条load instructions，每条load instruction都要符合contiguous的原则。

-- Ok! 第一条load instruction的顺序

显然，不符合三原则。

@1而且，当我们用__align__(8)也不能解决这个问题，但是可以用struct __align__(16) vec3d{x,y,z};来解决这个问题，那么这个时候可以分成两组128字节来读取。但是要浪费一些显存，因为这会让 compiler 在 vec3d 后面加上一个空的 4 bytes，以补齐 16 bytes。

@2 我们还有一个方法，就是把x,y,z三个数分开三块显存来放置，整个程序可能如下：

__global__ void Fun(float* d_inx,float* d_iny,float* d_intz,float*d_out)

{

Const int index = threadIdx.x + blockInx.x * BLOCK_DIM.x;

Float x = d_inx[x];

Float y = d_iny[y];

Float z = d_inz[z];

 

}

这里要提到一个概念，use SOA(structure ofarray) instand of AOS(array of structure)。

@3 但我们更一般的做法是把这些数值拷贝到__shared__内存中，但是__shared__内存是如此的珍贵，以至于一个block中最大为16k字节，所以，如果有256个线程的话，最多有64个字节的byte，也就是16个float类型。其他线程数，可以自行计算。

所以，可以用如下的方法得到这些数值，

__global__ void Fun(float* d_in,float d_out)

{

__shared__ float temp[256 * 3];

Const int index = threadInx.x * blockDim.x + threadIdx.x;

Temp[index] = d_in[index];

Temp[index + 256] = d_in[index + 256];

Temp[index + 256 * 2] = d_in[index + 256 * 2];

__syncthreads();//这样所有的数据就分成3条loadinstruction读入到__shared__memory中。

从上边可以看出，其实和分成3个数组来读取时道理是一样的，只不过@2是在外部就把结构调整好，使其符合coaleseced的情况，但为什么说，后者更一般一些，因为在处理数据的过程中，数据是什么样子的一般是无法由cuda程序员来控制的，那么cudaprogrammer就只能自己改变结构来适应客户导入的数据结构。

这里可以做一个预言，如果在显存中内存的access是coalescsed的，比读入到__shared__中要快。

}

…. …..

}

二、方法二

当compute capability为1.0或者是1.1的版本时

 

原则一：segment的划分原则

当每个线程读取的内存类型大小为4个字节的时候，那么一个segment为64字节；当线程读取的内存类型大小为8个字节时，那segment为128个字节；当读取的数据类型为16个字节的时候，segment还为128字节，不过读取要2次。注：这里的segment的大小，不是在一次__global__内核运算中式固定不变的，segment大小是对于某个需要读取的值而言的。

原则二：所有的half-wrap中的线程，都必须在这个segment中。

原则三：并且这些编号线程必须和内存的编号都是按照队列排列的。

注意这里的Additional restriction – 开始的地址一定是regment的大小的倍数，这样就可以解释为什么thread0都是对应128，因为128肯定是一个这样的一个值，不管segment是64byte或者是128byte。

 

一、当单个线程读取数据量为32bit的时候，那么hwrap就可以合并成64byte来传输

当compute capability版本为1.2或者以上的时

对齐方式如下步骤

一、找到half-warp中的第一个线程需要的地址所在的内存段(segment)，其中，segment的划分是如果是8bit的数据，为32bytes，如果是16bit的数据为64byte，如果是32、64位的数据是128bytes的数据。

二、找其他的线程是否也有需要的数据在此segment中；

三、如果segment是128bytes，但是只有一半的大小需要用到，那么segment就减少到64bytes，如果segment是64bytes就减少到32bytes。

四、执行从显存中传输到GPU的寄存器中，并且把已经完成的内存拷贝的线程赋值成非active的状态。

五、重复这个过程，直到半个wrap中的所有的线程都变成非active状态。

 

如下图就描述了这一个过程：

首先看第一副图像，因为传输的数据是float类型为32bits的，所以，segment预设为是128bytes。threadId为0的时候，为第128个字节处开始，正好为第二个segment的起始地址位置。于是查看其余的线程需要的地址也同样在此segment中，检测结果都是在的。接下来要传输这些数值了，但是发现只用了segment中的低64bytes，所以，减少到64bytes，最后传输。

 

第二副图和上一副图是一样，不过最后发现，要整个128byte的内存段都要传输。

 

第三幅图，首先找到的segment是在第一个segment中，然后，又发现这些线程中的需要的内存地址在这个地址段中的只有32bytes，所以传输32bytes，接下来，查看active的thread，发现最lowestID的thread在第二个segment中，最终查到剩余的所有的thread所占用此segment的ceiling是64，所以，再传输这64个byte的segment。

综上，知，第一次传输的次数最少，而且每次传输的数据量也最少。

第二次虽然也传输了一次，但是数据量要大，占用的内存带宽也大，不过内存带宽很大，比如5G/s，应该不是问题。

第三次传输，传输了2次。所用的时间最长，但是数据量还没有第二次大。

计算能力在1.2版本以上的GPU的global数据读取

 

这里有几个问题，如果拷贝到__shared__类型的内存也会发生如此的作用么？

10.2.2.2 __shared__内存的读取和存取

__shared__如在non-conflict的情况下，access数据的速度和register几乎是一样的，

10.2.2.3 local Memory操作

The constant memory space is cached so a read from constant memorycosts one memory read from device memory only on a cache miss, otherwise itjust costs one read from the constant cache.

以上这段和texture纹理的写的是一样的。