技术开发 频道

CUDA初探:GPU的并行计算

    【IT168 技术】请先看一段教程:

  到目前为止,我们的程序并没有做什么有用的工作。所以,现在我们加入一个简单的动作,就是把一大堆数字,计算出它的平方和。

  首先,把程序最前面的 include 部份改成:

#include <stdio.h>
#include
<stdlib.h>
#include
<cuda_runtime.h>

#define DATA_SIZE 1048576

int data[DATA_SIZE];

  并加入一个新函式 GenerateNumbers:

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

  这个函式会产生一大堆 0 ~ 9 之间的随机数。

  要利用 CUDA 进行计算之前,要先把数据复制到显卡内存中,才能让显示芯片使用。因此,需要取得一块适当大小的显卡内存,再把产生好的数据复制进去。在 main 函式中加入:

GenerateNumbers(data, DATA_SIZE);

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

  上面这段程序会先呼叫 GenerateNumbers 产生随机数,并呼叫 cudaMalloc 取得一块显卡内存(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(
"sum (CPU): %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);

  编译后执行,就可以看到执行所花费的时间了。

0
相关文章