【CUDA优化入门实战】用CUDA做个简单的加法_C

注：本markdown代码部分格式高亮基于c
代码详见(https://github.com/H-Freax/CUDA_optimization)

用CUDA做个简单的加法(cuda_sum1.cu)

该部分在first_cuda.cu的基础上进行

目标是计算一堆数字的平方和

首先把开头部分改成

#include
#include
#include

#define DATA_SIZE 1048576

int data[DATA_SIZE];

编写生成随机数的函数

void GenerateNumbers(int *number, int size){
    for(int i = 0; i < size; i++){
        number[i] = rand() % 10;
    }
}

CUDA执行计算前需要把数据从主记忆体复制到显示记忆体中，才能让显示晶片使用，因此需要获取一块大小适当的显示记忆体来复制数据。

在main函数中加入以下部分

   GenerateNumbers(data,DATASIZE);
    int* gpudata, *result;

    cudaMalloc((void**) &gpudata,sizeof(int) * DATA_SIZE);
    cudaMalloc((void**) &result,sizeof(int));
    //从主记忆体复制到显示记忆体，所以使用 cudaMemcpyHostToDevice。

    //如果是从显示记忆体复制到主记忆体，则使用 cudaMemcpyDeviceToHost
    cudaMemcpy(gpudata, data,sizeof(int) * DATA_SIZE,
        cudaMemcpyHostToDevice);

首先引用GenerateNumbers函数产生数据，使用cudaMalloc取得大小合适的记忆体，gpudata变量用来存储数据，result用来存储计算结果，通过cudaMemcpy将数据复制到显示记忆体中。

cudaMalloc和cudaMemcpy的用法和一般的malloc以及memcpy类似，不过cudaMemcpy多出一个参数，指示复制记忆体的方向，如果是从主记忆体复制到显示记忆体，使用 cudaMemcpyHostToDevice，如果是从显示记忆体复制到主记忆体，则使用 cudaMemcpyDeviceToHost。

接下来编写在显示晶片上的函数，在CUDA中，在函数前面加上__global__表示这个函数是在显示晶片上执行的。

__global__ static void sumOfSquares(int *num, int* result)
{
    int sum = 0;
    int i;
    for(i = 0; i < DATA_SIZE; i++) {
        sum += num[i] * num[i];
    }

    *result = sum;
}

在显示晶片上执行程序有一定的限制条件，例如不能传回值。

接下来是让CUDA执行函数，在CUDA中执行函数需要使用以下的语法

函数名称<<<block 数目, thread 数目, shared memory 大小>>>(参数);

执行完毕后需要将结果从显示晶片复制回主记忆体，在main函数中加入以下部分

    sumOfSquares<<<1, 1, 0>>>(gpudata, result);

    int sum;
    cudaMemcpy(&sum, result, sizeof(int), cudaMemcpyDeviceToHost);
    cudaFree(gpudata);
    cudaFree(result);

    printf("sum: %d\n", sum);

编写的部分设置了只适用一个thread，所以block数目跟thread数目都为1，没有使用到shared memory，设为0。

如果要验证结果是否正确，可以加入CPU的执行代码进行验证

    sum = 0;
    for(int i = 0; i < DATA_SIZE; i++) {
        sum += data[i] * data[i];
    }
    printf("CPU sum: %d\n", sum);

CUDA提供了clock函数可以取得timestamp，适合用来判断程序执行的时间，单位为GPU执行单元的时间，可以在优化程序时间时提供参考，如果需要记录时间的话，需要更改sumOfSquares为以下内容

__global__ static void sumOfSquares(int *num, int* result, clock_t* time){
    int sum = 0;
    int i;
    clock_t start = clock();
    for(i =0; i<DATA_SIZE; i++){
        sum+= num[i]*num[i];
    }

    *result = sum;
    *time = clock()-start;
}

main函数的部分修改如下：

    int* gpudata, *result;
    clock_t* time;
    cudaMalloc((void**) &gpudata, sizeof(int) * DATA_SIZE);
    cudaMalloc((void**) &result, sizeof(int));
    cudaMalloc((void**) &time, sizeof(clock_t));
    cudaMemcpy(gpudata, data, sizeof(int) * DATA_SIZE,
        cudaMemcpyHostToDevice);

    sumOfSquares<<<1, 1, 0>>>(gpudata, result, time);

    int sum;
    clock_t time_used;
    cudaMemcpy(&sum, result, sizeof(int), cudaMemcpyDeviceToHost);
    cudaMemcpy(&time_used, time, sizeof(clock_t),
        cudaMemcpyDeviceToHost);
    cudaFree(gpudata);
    cudaFree(result);

    printf("sum: %d time: %d\n", sum, time_used);