cuda入门程序

• 虽然是作业，但是我也觉得了解一下cuda还是很有必要的，之前本科想做硬件加速时也不知道应该先学学cuda是怎么工作的。
• Linux上安装cuda后, cuda其实提供了很多例程, 如下图. 但是还是感觉太复杂了,所以这里总结几个简单的cuda程序.
  
• 需要注意的是使用了cuda的c程序后缀名应为.cu , 同时编译器为nvcc ,如nvcc check_device.cu -o check_deviec

1 查看设备

在cuda安装目录下的./include/cuda_runtime_api.h中可以看到对结构体cudaDeviceProp的说明.

// check_device.cu
#include <stdio.h>
//#include <cuda_runtime.h>
int main(int argc, char **argv)
{
   
        int deviceCount;
        cudaGetDeviceCount(&deviceCount);
        printf("Device count: %d\n", deviceCount);

        int device;
        for(device = 0; device < deviceCount; ++device)
        {
   
                cudaDeviceProp deviceProp;
                cudaGetDeviceProperties(&deviceProp,device);
                printf("Device id: %d, \n\tDevice name: %s\n", device, deviceProp.name);
                printf("\tCompute capability %d.%d.\n",deviceProp.major,deviceProp.minor);
                printf("\tNum of Stream Multiprocess(SM): %d\n", deviceProp.multiProcessorCount);
                printf("\t每个线程块的最大线程数: %d\n",  deviceProp.maxThreadsPerBlock);
                printf("\t每个SM的最大线程数: %d, 最大线程束数: %d\n", 
                        deviceProp.maxThreadsPerMultiProcessor, deviceProp.maxThreadsPerMultiProcessor/32);

        }
        return 0;
}

2 向量加法

• cuda编写简单的程序,不着重考虑性能的话只需要了解一下grid和block的概念,以及__global__和__device__关键字,参考1中写的很简明了.编写核函数时思维和在cpu上编程非常不同, 要考虑其如何并行, **即每个thread不能依赖于其他thread的结果。**但实际代码其实只需稍作改动就可以。

// vector_adder.cu
// 程序有点问题,运行后显存不释放
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>

   
// 每个thread只做一个元素的加法
__global__ void add_1( int *a, int *b, int *c, int num) 
{
   

    int idx = threadIdx.x + blockIdx.x * blockDim.x;

    c[idx] = a[idx] + b[idx];
    //  Fermi以后的架构,可以在核函数中调用printf
    // printf("%d\n",c[idx]);
}

// thread数量有限,当数据量大的时候, 就需要每个thread做多个元素的加法.
// 这里的代码是类比参考2 CUDA编程入门极简教程，这个代码逻辑没问题，也能并行计算，但我认为不符合cuda的编程思想。应该按照参考1 线程块调度的思想来写程序。也就是核函数中只关注一次计算就可以了。
__global__ void add_2( int *a, int *b, int *c, int num) 
{
   
    // 获取thread在整个grid中的索引.
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    // 计算thread总数
    int stride = blockDim.x * gridDim.x;

    // 每个thread处理 num/stride 个元素
    for (int i = idx; i < num; i += stride)
    {
   
        c[i] = a[i] + b[i];
    }
}

/* 我们最好是参考这种形式来写,然后分配足够多的block和thread，显然之后大于实际可用的thread数量，但是gpu会按照block进行调度。当gpu中的计算单元完成核函数的计算后就会被释放，此时等候的block就得以分配到SM上。
*/
/*
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N])
{
	int i = blockIdx.x * blockDim.x + threadIdx.x;
	int j = blockIdx.y * blockDim.y + threadIdx.y;
	if (i < N && j < N)
		C[i][j] = A[i][j] + B[i][j];
}
*/

int main(int argc, char **argv)
{
   
    // 向量长度和字节数
    int N = 20;
    int bytes = N*sizeof(int);
    int pass = 1;

    int *dev_c1;
    int *dev_c2;
    int *dev_b;
    int *dev_a;
    // 申请host内存,即CPU接的内存
    int *vec_a = (int *)malloc(bytes); 
    int *vec_b = (int *)malloc(bytes);     
    int *vec_c1 = (int *)malloc(bytes);
    int *vec_c2 = (int *)<