cuda学习笔记

cuda学习笔记

标签（空格分隔）： 学习笔记

---

一、学习平台

1.1学习平台搭建与学习平台基本信息

1.1.1学习平台搭建

此篇文档的整理基于nvidia公司出品的GeForce GTX 950 GPU，在电脑主机当中安装好独立显卡之后，安装cuda7.0至软件盘（不用再单独安装显卡驱动程序）。在vs下新建cuda工程，就可以编写cuda程序了。

1.1.2学习平台基本信息

在编写cuda程序时，程序的头文件应该包括 “cuda_runtime.h”和”device_launch_parameters.h”；以下一段代码用来查看显卡gpu的计算性能和架构

int main()
{
    cudaDeviceProp prop;

    int count;
    cudaGetDeviceCount(&count);

    for (int i = 0; i < count; ++i){
        cudaGetDeviceProperties(&prop, i);
        printf("     --- Genaral Information for Device %d ---\n", i);
        printf("Name : %s\n",prop.name);
        printf("Compute capability : %d.%d\n", prop.major, prop.minor);
        printf("Clock rate:  %d\n",prop.clockRate);
        printf("Device copy overlap: ");

        if (prop.deviceOverlap){
            printf("Enabled\n");
        }
        else{
            printf("Disabled\n");
        }

        printf("Kernel execition timeout : ");
        if (prop.kernelExecTimeoutEnabled)
            printf("Enabled\n");
        else
            printf("Disabled\n");

        printf("  ---Memory Information for Device %d ---\n",i);
        printf("Total global mem: %ld\n",prop.totalGlobalMem);
        printf("Total const mem : %ld\n", prop.totalConstMem);
        printf("Max mem pitch : %ld\n", prop.memPitch);
        printf("Texture Alignment : %ld\n",prop.textureAlignment);


        printf("   ---MP Information for device %d ---\n", i);
        printf("Multiprocessor count : %d\n", prop.multiProcessorCount);
        printf("shared mem per mp: %d\n", prop.sharedMemPerBlock);
        printf("Register per mp: %d\n",prop.regsPerBlock);
        printf("Threads in warp: %d\n", prop.warpSize);
        printf(" Max threads per block :%d\n", prop.maxThreadsPerBlock);
        printf("Max thread dimentions : (%d, %d, %d)\n", prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);
        printf("Max grid dimensions:(%d, %d, %d)\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2] );

        printf("\n\n\n");

    }
}

950显卡运行结果如下图所示：

二、cuda线程模型

2.1线程索引相关问题

我们把在GPU上启动的线程块集合称为一个线程格。从名字的含义可以看出，线程格既可以使一维的线程块集合，也可以是二维的线程块集合。核函数的每个副本都可以通过内置变量blockIdx来判断哪个线程块正在执行它。同样，它还可以通过内置变量gridDim来获得线程块的大小。通过这两个变量来计算每个线程块需要的数据索引。
则当前线程块的索引如下：
线程块的索引=行索引*线程格的数目+列索引
blockIdx.y * gridDim.x+blockIdx.x;
同样的：
线程索引 = 行索引*线程块的数目+列索引
threadIdx.y * blockIdx.x + threadIdx.x;
int offset = x + y * Dim;在这里Dim表示线程块的大小（也就是线程的数量），y为线程块索引，并且x为线程块中的线程索引，所以计算得到如下索引:

int tid = blockDim.x*blockIdx.x + threadIdx.x;
tid += blockDim.x*gridDim.x;
//每个线程块中的数量乘以线程格中线程块的总数量，即为当前线程格中运行的线程总数量。

对于二维线程的索引，有如下代码：

int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
int offset = x + y * blockDim.x * gridDim.x;

2.2线程开辟

一种解决方案是将线程块的大小设置为某个固定数值BLOCKSIZE，然后启动N/BLOCKSIZE个线程块，这样就相当启动了N个线程同时运行。通常我们设置的线程块的个数为（N+BLOCKSIZE-1）/BLOCKSIZE来防止0线程的开辟问题。

2.3多维数据线程调用

线性存储器也可以通过cudaMallocPitch（）和cudaMalloc3D分配。在分配二维和三维数组的时候，推荐使用，因为上述调用保证了GPU的最佳性能。返回的（pitch,stride）必须用于访问数组元素。下面的代码分配了一个尺寸为weight*height的二维浮点数组，同时演示了怎么在设备代码中遍历数组元素

//host code
int width =64,height = 64;
float *dexPtr;
int pitch;
cudaMallocPitch((void **)&devPtr,&pitch,width*sizeof(float),height);
kernel<<<100,512>>>(devPtr,pitch,widtf,height);

//device code
__global__void kernel(float* devPtr,int pitch,int width,int height)
{
for(int i =0;i<helght;++i){
float* row =(float*) ((char*)devPtr+i*pitch);
for(int j = 0;j<width;++j){
float element = row[i];
        }
    }
}

下面的代码演示分配一个尺寸为width*height*depth的三维浮点数组，同时演示了怎么在设备代码中遍历数组元素。

//host code
cudaPitchedPtr devPitchedPtr;
cudaExtent extent = make_cudaExtent(64,64,64);
cudaMalloc3D(&devPitchedPtr,extent);
kernel<<<100,512>>>(devPitchedPtr,extent);

//device code
__global__ void kernel(cudaPitchedPtr devPitchedPtr,cudaExtent extent){
char *devPtr=devPitchedPtr.ptr;
size_t pitch =devPitchedPtr.pitch;
size_t slicePitch = pitch*extent.height;

for(int i=0;i<extent.depth;++i){
char *slice = devPtr + i*slicePitch;
for(int j=0;j<extent.height;++j){
float *row=(float*)(slice + y*pitch);
for(int x = 0;x<extent.width;++x){
float element = row[x];
        }
    }
}

三、矩阵乘法

3.1使用全局内存

#include<iostream>
#include<fstream>
using namespace std;
#define AROWS 16
#define ACOLS 4
#define BROWS 4
#define BCOLS 16
void randomInit(float* _data, int _size)
{
    for (int i = 0; i < _size; ++i)
    {
        _data[i] = rand() %5;
        //_data[i] = (float)(rand() / (float)RAND_MAX);
    }
}
__global__ void MatMul(float *dev_a, float *dev_b, float *dev_c)
{

    int col = threadIdx.x;
    int row = threadIdx.y;
    float cvalue = 0;
    for (int i = 0; i<ACOLS; i++)
    {
        cvalue += dev_a[row*ACOLS+i] * dev_b[BCOLS*i + col];
    }
    dev_c[row*BCOLS + col] = cvalue;
}

int main()
{
    float *a, *b, *c, *dev_a, *dev_b, *dev_c;
    a = (float*)malloc(sizeof(float)*AROWS*ACOLS);
    b = (float*)malloc(sizeof(float)*BROWS*BCOLS);
    c = (float*)malloc(sizeof(float)*AROWS*BCOLS);
    memset(c, 0, sizeof(float)*AROWS*BCOLS);
    cudaMalloc((void **)&dev_a, sizeof(float)*AROWS*ACOLS);
    cudaMalloc((void **)&dev_b, sizeof(float)*BROWS*BCOLS);
    cudaMalloc((void **)&dev_c, sizeof(float)*AROWS*BCOLS);
    cudaMemset(dev_c, 0, sizeof(float)*AROWS*BCOLS);
    //给矩阵AB赋初值
    randomInit(a, AROWS*ACOLS);
    randomInit(b, BROWS*BCOLS);
    for (int i = 0; i < AROWS*ACOLS; i++)
    {
        cout << "a[" << i << "]=" << a[i] << endl;
    }
    for (int i = 0; i < BROWS*BCOLS; i++)
    {
        cout << "b[" << i << "]=" << b[i] << endl;
    }

    cudaMemcpy(dev_a, a, sizeof(float)*AROWS*ACOLS, cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, sizeof(float)*BROWS*BCOLS, cudaMemcpyHostToDevice);
    dim3 blocks(16, 16);
    dim3 threads(AROWS+15, BCOLS+15);
    MatMul << <blocks, threads >> >(dev_a, dev_b, dev_c);
    cudaMemcpy(c, dev_c, sizeof(float)*AROWS*BCOLS, cudaMemcpyDeviceToHost);

    //打印C矩阵到文件中
    ofstream outFile;
    outFile.open("result.txt");
    outFile << fixed;
    outFile.precision(2);
    outFile.setf(ios_base::showpoint);
    for (int i = 0; i < AROWS*BCOLS; i++)
    {
        outFile << "C[" << i << "]=" << c[i] << endl;
    }

    system("pause");
    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);
    free(a);
    free(b);
    free(c);
    return 0;
}

使用全局内存实现矩阵的一维向量乘法如上程序所示，但这种实现方式并没有充分利用gpu的优势。下面的代码是使用共享内存实现矩阵乘法。

3.2使用共享内存

四、关于cuda的基本架构

在CUDA架构下，线程的最小单元是thread，多个thread组成一个block，多个block再组成一个grid。每一个block中开辟的所有线程共享同一个shared memory。block里面的thread之间的通信和同步所带来的开销是比较大的。SM以 32 个 Thread 为一组的 Warp 来执行 Thread。Warp内的线程是静态的，即在属于同一个warp内的thread之间进行通信，不需要进行栅栏同步(barrier)。

4.1共享内存

每个block中开辟的所有线程共享一个shared memory。使用共享内存变量的时候，需要在声明的时候加上shared关键词修饰，使用共享内存的时候应注意同一个线程块中的线程都执行结束才能进行下一步操作，所以需要使用__syncthread关键词使得block中的线程同步。

4.2常量内存

常量内存用于保存在核函数执行期间不会发生变化的数据，定义常量内存的时候应该使用关键词constant进行修饰。当从主机内存复制到GPU上的常量内存时，需要使用cudaMemcpyToSymbol()复制数据。常量内存读取数据可以节约内存带宽。注：只要当一个warp的半线程束中的所有16个线程有相同的读取请求时，才值得使用常量内存。

4.3纹理内存

纹理内存是专门为那些在内存访问模式中存在大量空间局部性的图形应用程序而设计的。
以下内容参考博文http://www.cnblogs.com/traceorigin/archive/2013/04/11/3015755.html

4.3.1、概述

纹理存储器中的数据以一维、二维或者三维数组的形式存储在显存中，可以通过缓存加速访问，并且可以声明大小比常数存储器要大的多。
在kernel中访问纹理存储器的操作称为纹理拾取(texture fetching)。将显存中的数据与纹理参照系关联的操作，称为将数据与纹理绑定(texture binding).
显存中可以绑定到纹理的数据有两种，分别是普通的线性存储器和cuda数组。
注：线性存储器只能与一维或二维纹理绑定,采用整型纹理拾取坐标，坐标值与数据在存储器中的位置相同；
CUDA数组可以与一维、二维、三维纹理绑定，纹理拾取坐标为归一化或者非归一化的浮点型，并且支持许多特殊功能。

4.3.2、纹理缓存：

（1）、纹理缓存中的数据可以被重复利用
（2）、纹理缓存一次预取拾取坐标对应位置附近的几个象元，可以实现滤波模式。

4.3.3纹理存储器的使用：

使用纹理存储器时，首先要在主机端声明要绑定到纹理的线性存储器或CUDA数组
(1)声明纹理参考系

texture<Type, Dim, ReadMode> texRef;
//Type指定数据类型，特别注意：不支持3元组
//Dim指定纹理参考系的维度，默认为1
//ReadMode可以是cudaReadModelNormalizedFloat或cudaReadModelElementType(默认)

注：纹理参照系必须定义在所有函数体外
(2) 声明CUDA数组，分配空间
CUDA数组可以通过cudaMalloc3DArray()或者cudaMallocArray()函数分配。前者可以分配1D、2D、3D的数组，后者一般用于分配2D的CUDA数组。使用完毕，要用cudaFreeArray()函数释放显存。

 //1数组
 cudaMalloc((void**)&dev_A, data_size);
 cudaMemcpy(dev_A, host_A, data_size, cudaMemcpyHostToDevice);
 cudaFree(dev_A);

 //2维数组
 cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>()
 cudaArray *cuArray;
 cudaMallocArray(&cuArray, &channelDesc, 64, 32); //64x32
 cudaMemcpyToArray(cuArray, 0, 0, h_data, sizeof(float)*width*height, cudaMemcpyHostToDevice);
 cudaFreeArray(cuArray);

 //3维数组 64x32x16
 cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<uchar>();
 cudaArray *d_volumeArray;
 cudaMalloc3DArray(&d_volumeArray, &channelDesc, volumSize);

 cudaMemcpy3DParms copyParams = {0};
 copyParams.srcPtr   = make_cudaPitchedPtr((void*)h_volume, volumeSize.width*sizeof(uchar), volumeSize.width, volumeSize.height);
 copyParams.dstArray = d_volumeArray;
 copyParams.extent   = volumeSize;
 copyParams.kind     = cudaMemcpyHostToDevice;
 cudaMemcpy3D(&copyParams);

 tex.normalized = true;
 tex.filterMode = cudaFilterModeLinear;
 tex.addressMode[0] = cudaAddressModeWrap;
 tex.addressMode[1] = cudaAddressModeWrap;
 tex.addressMode[2] = cudaAddressModeWrap;

（3）设置运行时纹理参照系属性

struct textureReference
{
    int normalized;
    enum cudaTextureFilterMode filterMode;
    enum cudaTextureAddressMode addressMode[3];
    struct cudaChannelFormatDesc channelDesc;
}

normalized设置是否对纹理坐标归一化
filterMode用于设置纹理的滤波模式
addressMode说明了寻址方式

(4)纹理绑定
通过cudaBindTexture() 或 cudaBindTextureToArray()将数据与纹理绑定。
通过cudaUnbindTexture()用于解除纹理参照系的绑定
注：与纹理绑定的数据的类型必须与声明纹理参照系时的参数匹配
(I).cudaBindTexture() //将1维线性内存绑定到1维纹理

cudaError_t cudaBindTexture
(    
    size_t *     offset,
    const struct textureReference *     texref,
    const void *     devPtr,
    const struct cudaChannelFormatDesc *     desc,
    size_t     size = UINT_MAX     
)

(II).cudaBindTexture2D //将1维线性内存绑定到2维纹理

cudaError_t cudaBindTexture2D(    
    size_t *     offset,
    const struct textureReference *     texref,
    const void *     devPtr,
    const struct cudaChannelFormatDesc *     desc,
    size_t     width,
    size_t     height,
    size_t     pitch     
)

(III). cudaBindTextureToArray() //将cuda数组绑定到纹理

 cudaError_t cudaBindTextureToArray   
 (    
 const struct textureReference *     texref,
 const struct cudaArray *     array,
 const struct cudaChannelFormatDesc *     desc     
 )

(5)纹理拾取
　　对于线性存储器绑定的纹理，使用tex1Dfetch()访问，采用的纹理坐标是整型。由cudaMallocPitch() 或者 cudaMalloc3D()分配的线性空间实际上仍然是经过填充、对齐的一维线性空　　　间，因此也用tex1Dfetch()
　　对与一维、二维、三维cuda数组绑定的纹理，分别使用tex1D(), tex2D() 和 tex3D()函数访问，并且使用浮点型纹理坐标。

4.4cuda流

五、使用事件对cuda程序进行计时

创建事件

cudaEvent_t  start,stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start,0);

……..此处省略cuda程序

事件结束

cudaEventRecord(stop,0);
cudaEventSynchronize(stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime,start,stop);
printf("Time to generate :  %3.1f ms\n",elapsedTime);
cudaEventDestroy(start);
cudaEventDestroy(stop);

六、调试cudaC

Parallel Nsight, visual profiler

4.1线程最优配置tips
最优的cuda线程配置
1 每个SM上面失少要有192个激活线程，寄存器写后读的数据依赖才能被掩盖

2 将 寄存器 的bank冲突降到最低，应尽量使每个block含有的线程数是64的倍数

3 block的数量应设置得令可用的计算资源被充分的利用。由于每个block映射到一个sm上面，所以至少应该让block的数目跟sm的数目一样多。

4 当Block中的线程被同步时或者等待读取设备存储器时，相应的SM会闲置。通常让block的数目是sm的2倍以上，使其在时间轴上重叠

5 如果block的数目足够多，则每个Block里的线程数应设置成warp尺寸的整数倍，以免过小的warp浪费计算资源。

6 给每个block分配越多的线程，能更高效的让他们在时间片上工作。但是相应的每个线程的寄存器也就越少。当寄存器过少，有可能因为访问溢出的寄存器，而导致数据的存储变慢。

7 当每个线程占用的寄存器较多时，不宜在Block内分配过多的线程，否则也会减少block的数目。从而使SM的工作效率降低

8 每个block内的线程数应遵循 相应的 计算能力等级中的规定数目。

9 当线程块的数量为GPU中处理器数量的2倍时，计算性能达到最优。