CUDA C 开发入门

"Hello, world" 通常是我们编写的第一个程序。我们可以对 CUDA 做同样的事情。

1、最小CUDA C程序

在文件 hello.cu 中输入如下代码：

#include "stdio.h"

int main()
{
    printf("Hello, world\n");
    return 0;
}

在我们的 安装了CUDA SDK的机器上，你可以使用以下命令进行编译：

$ nvcc hello.cu
$ ./a.out

可以使用 -o 标志更改输出文件名： nvcc -o hello hello.cu

如果不想一直输入前面的 . （它指的是当前工作目录），可以编辑 .bashrc 文件，将当前目录添加到路径中：

export PATH=$PATH:.

不过有些人建议出于安全目的不要这样做。

2、使用核函数

你可能认为这个程序是作弊，因为它实际上不使用任何 CUDA 功能，一切都在主机上运行。但是，重点是 CUDA C 程序可以完成常规 C程序可以做的所有事情。

下面是一个稍微有趣一点（但效率低下且仅用作示例）的程序，它使用核函数 add() 将两个数字相加：

#include "stdio.h"

__global__ void add(int a, int b, int *c)
{
    *c = a + b;
}

int main()
{
    int a,b,c;
    int *dev_c;
    
    a=3;
    b=4;
    cudaMalloc((void**)&dev_c, sizeof(int));
    
    add<<<1,1>>>(a,b,dev_c);
    cudaMemcpy(&c, dev_c, sizeof(int), cudaMemcpyDeviceToHost);
    printf("%d + %d is %d\n", a, b, c);
    
    cudaFree(dev_c);
    return 0;
}

请浏览上面的代码，它运行得更快！考虑一下为什么？

如果成功， cudaMalloc 将返回 cudaSuccess；可以检查以确保程序能够正确运行。

3、示例：向量求和

这是一个简单的问题。给定两个向量（即数组），我们希望将它们相加到第三个数组中。例如：

A = {0, 2, 4, 6, 8}
B = {1, 1, 2, 2, 1}

那么

  C = A + B = {1, 3, 6, 8, 9}

在此示例中，数组长度为 5 个元素，因此我们创建 5 个不同的线程（thread）。

第一个线程负责计算 C[0] = A[0] + B[0]。第二个线程负责计算 C[1] = A[1] + B[1]，依此类推。

以下是我们如何使用传统 C 代码执行此操作：

#include "stdio.h"

#define N 10

void add(int *a, int *b, int *c)
{
    int tID = 0;
    while (tID < N)
    {
        c[tID] = a[tID] + b[tID];
        tID += 1;
    }
}

int main()
{
    int a[N], b[N], c[N];
    
    // Fill Arrays
    for (int i = 0; i < N; i++)
    {
        a[i] = i,
        b[i] = 1;
    }
    add (a, b, c);
    for (int i = 0; i < N; i++)
    {
        printf("%d + %d = %d\n", a[i], b[i], c[i]);
    }
    return 0;
}

这是两个数组求和的一种相当迂回的方法——我们这样做的原因是，这样可以更好地转换为 CUDA 版本。要编译和运行它，我们必须使用 g++（因为它使用了一些在 C 中不起作用的 C++ 样式符号）。

以下是 CUDA 版本：

#include "stdio.h"

#define N 10

__global__ void add(int *a, int *b, int *c)
{
     int tID = blockIdx.x;
     if (tID < N)
     {
	     c[tID] = a[tID] + b[tID];
     }
}
int main()
{
    int a[N], b[N], c[N];
    int *dev_a, *dev_b, *dev_c;
    cudaMalloc((void **) &dev_a, N*sizeof(int));
    cudaMalloc((void **) &dev_b, N*sizeof(int));
    cudaMalloc((void **) &dev_c, N*sizeof(int));

	// Fill Arrays
    for (int i = 0; i < N; i++)
    {
        a[i] = i,
        b[i] = 1;
    }
    cudaMemcpy(dev_a, a, N*sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, N*sizeof(int), cudaMemcpyHostToDevice);
    
    add<<<N,1>>>(dev_a, dev_b, dev_c);
    
    cudaMemcpy(c, dev_c, N*sizeof(int), cudaMemcpyDeviceToHost);
    
    for (int i = 0; i < N; i++)
    {
    	printf("%d + %d = %d\n", a[i], b[i], c[i]);
    }
    return 0;
}

blockIDx.x 为我们提供 Block ID，范围从 0 到 N-1。如果我们改用 add<<<1,N>>> 会怎么样？然后我们可以通过 ThreadID 访问，即变量 threadIDx.x。

再举一个例子，让我们求两个 2D 数组的和。我们可以定义一个 2D 整数数组，如下所示：

int c[2][3];

以下代码说明了 2D 数组在内存中的布局方式：

for (int i=0; i < 2; i++)
    for (int j=0; j< 3; j++)
    printf("[%d][%d] at %ld\n",i,j,&c[i][j]);

输出如下：

[0][0] at 140733933298160
[0][1] at 140733933298164
[0][2] at 140733933298168
[1][0] at 140733933298172
[1][1] at 140733933298176
[1][2] at 140733933298180

我们可以看到，我们有一个布局，其中 j 维中的下一个单元格占据内存中的下一个连续整数，其中一个 int 为 4 个字节：

一般来说，单元格的地址可以通过以下方式计算：

&c + [(sizeof(int) * sizeof-j-dimension * i] + (sizeof(int)) * j 

在我们的示例中，j 维的大小为 3。例如， c[1][1] 处的单元格将合并为：
基地址 + (431) + (41) = &c+16。

如果我们使用数组表示法，C 将为我们进行寻址，因此如果 INDEX=iWIDTH + J
那么我们可以通过以下方式访问元素： c[INDEX]

CUDA 要求我们将内存分配为一维数组，因此我们可以使用上面的映射到2D 数组。

为了使内核函数中的映射更容易一些，我们可以将块声明为与 2D 数组尺寸相同的网格。这将创建与数组宽度和高度相对应的变量 blockIdx.x 和 blockIdx.y 。

#include "stdio.h"

#define COLUMNS 3
#define ROWS 2

__global__ void add(int *a, int *b, int *c)
{
    int x = blockIdx.x;
    int y = blockIdx.y;
    int i = (COLUMNS*y) + x;
    c[i] = a[i] + b[i];
}
int main()
{
    int a[ROWS][COLUMNS], b[ROWS][COLUMNS], c[ROWS][COLUMNS];
    int *dev_a, *dev_b, *dev_c;
    
    cudaMalloc((void **) &dev_a, ROWS*COLUMNS*sizeof(int));
    cudaMalloc((void **) &dev_b, ROWS*COLUMNS*sizeof(int));
    cudaMalloc((void **) &dev_c, ROWS*COLUMNS*sizeof(int));
    
    for (int y = 0; y < ROWS; y++) // Fill Arrays
        for (int x = 0; x < COLUMNS; x++)
        {
            a[y][x] = x;
            b[y][x] = y;
        }

	cudaMemcpy(dev_a, a, ROWS*COLUMNS*sizeof(int),
    cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, ROWS*COLUMNS*sizeof(int),
    cudaMemcpyHostToDevice);
    
    dim3 grid(COLUMNS,ROWS);
    add<<<grid,1>>>(dev_a, dev_b, dev_c);
    cudaMemcpy(c, dev_c, ROWS*COLUMNS*sizeof(int),
    cudaMemcpyDeviceToHost);
    
    for (int y = 0; y < ROWS; y++) // Output Arrays
    {
        for (int x = 0; x < COLUMNS; x++)
        {
	        printf("[%d][%d]=%d ",y,x,c[y][x]);
        }
        printf("\n");
    }
    return 0;
}

原文链接：CUDA – First Programs

BimAnt翻译整理，转载请标明出处