你好,游客 登录 注册 搜索
背景:
阅读新闻

CUDA入门教程

[日期:2014-07-17] 来源:Linux社区  作者:hipercomer [字体: ]

1. 准备makefile

为了避免每次都要键入nvcc的命令,要准备一个makefile。makefile如下:

CUFLAG = -g  -Xcompiler -v \
        -gencode=arch=compute_20,code=sm_20\
        -gencode=arch=compute_20,code=compute_20\
        -O2
IFLAGS = -I$(CUDA_DIR)/include -I$(CUDA_SDK_DIR)/C/common/inc -I../include
LFLAGS = -L$(CUDA_DIR)/lib64 -L$(CUDA_SDK_DIR)/C/lib
PRG = cuda_test
$(PRG) : main.cu
 nvcc main.cu -o $(PRG) $(CUFLAG) $(IFLAGS) $(LFLAGS)

Ubuntu 12.04 下 CUDA 编程 http://www.linuxidc.com/linux/2014-06/103056.htm

Ubuntu 12.04 安装 CUDA-5.5  http://www.linuxidc.com/Linux/2013-10/91101.htm

Ubuntu 11.10 上安装CUDA开发环境 http://www.linuxidc.com/Linux/2012-04/58913.htm

Fedora 15系统下配置CUDA环境 http://www.linuxidc.com/Linux/2011-12/49874.htm

Ubuntu 11.04 安装 NVIDIA CUDA 4.0 RC2 http://www.linuxidc.com/Linux/2011-10/46304.htm

Linux Mint 13/Ubuntu 12.04 配置CUDA 4.2 & OpenCV 2.4.2 方法 http://www.linuxidc.com/Linux/2013-10/91102.htm

2 异构计算(Heterogeneous Computing)

以下为几个技术名词的简单介绍:

 

  • 主机(host):CPU及其内存(host memory)。
  • 设备(device):GPU及其内存(device memory)。
  • 主机代码(host code):运行在CPU上的(一般来说「串行执行」的)代码。
  • 设备代码(device code):运行在GPU上的并行执行的代码。
  • 异构计算:由主机代码(host code)和设备代码(device code)协同执行完成的计算。

宏观上看,GPU执行代码的流程如下:

 

 

  1. 将输入数据通过PCI总线从CPU内存拷贝到GPU的DRAM中。
  2. 从内存中加载需要执行的代码到GPU后。
  3. 数据和指令都就绪后,就可以执行了。注意,在执行的过程中,GPU会在片上缓存数据以提升性能。
  4. 计算完毕后,将结果从GPU的DRAM中拷回CPU的Memory中。

例1: Hello World

#include<stdio.h>
#include<stdlib.h>
#include<cuda.h>
#include<cutil.h>

__global__ void mykernel(void) {
}

int main(void) {
 mykernel<<<1,1>>>();
 printf("Hello World!\n");
 return 0;
}

上述代码编译后运行生成可执行文件cuda_test,运行cuda_test后将输出:

 

Hello World!

注意:

  1. 调用kernel时需要三个尖括号
  2. 包含必要的头文件

CUDA C/C++中引入的新关键字__global__所修饰的函数有以下两方面含义:

  • 此函数代码由设备执行
  • 此函数由主机代码调用

nvcc将源代码分为设备函数和主机函数两大类:

  • 设备函数由NVIDA编译器编译
  • 主机函数由主机上配置的编译器编译
三个尖括号标志着一个从主机代码调用设备代码的函数,称为“启动内核”(kernel launch)

 

例2: 整数相加

#include<stdio.h>
#include<stdlib.h>
#include<cuda.h>
#include<cutil.h>

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

int main(void) {
 int a,b,c;
 int * d_a, * d_b, * d_c;
 int size = sizeof(int);
 cudaMalloc((void**)&d_a,size);
 cudaMalloc((void**)&d_b,size);
 cudaMalloc((void**)&d_c,size);
 printf("Enter two integers with a space to separate them:\n");
 scanf("%d %d",&a,&b);
 cudaMemcpy(d_a,&a,size,cudaMemcpyHostToDevice);
 cudaMemcpy(d_b,&b,size,cudaMemcpyHostToDevice);
 integer_add<<<1,1>>>(d_a,d_b,d_c);
 cudaMemcpy(&c,d_c,size,cudaMemcpyDeviceToHost);
 cudaFree(d_a);
 cudaFree(d_b);
 cudaFree(d_c);
 printf("Sum is %d\n",c);
 return 0;
}

__global__修饰的integer_add函数说明:

  • integer_add函数将在设备上执行
  • integer_add函数将被主机调用
由于integer_add函数在设备上执行,所以指针a,b,c应该指向设备内存。这说明需要在设备内存中为变量开辟内存。
 
设备内存和主机内存在物理上是完全分开的不同电子部件:
  • 设备指针指向GPU内存的某个位置。设备指针可以从主机端传给设备端或者从设备端传给主机端,但是设备指针不能在主机端解引用。
  • 主机指针指向CPU内存的某个位置。主机指针可以从设备端传给主机端或者从主机端传给设备端,但是主机指针不能在设备端解引用。
CUDA API提供的用于处理设备内存的函数有cudaMalloc, cudaFree, cudaMemcpy。语义上分别对应于C语言的malloc, free, memcpy函数。这几个函数的具体使用方法如例2所示。
 

3 块(Blocks)

GPU是用来实现大规模并行的,如何实现呢?将上述例子扩展一下,如果我们要实现两个向量相加:
add<<<1,1>>>() ---> add<<<N,1>>>
N表示同时调用N次add函数,这样就可以实现并行的向量相加了。
每个被并行调用的add函数称之为一个(block)。
  • 块的集合称之为网格(grid).
  • 每个块可以使用索引值blockIdx.x
通过使用blockIdx.x作为索引,每个块可以处理数组元素中的一部分。
有了这些基础后,就可以实现并行版本的向量相加了。
例3:向量相加

#include<stdio.h>
#include<stdlib.h>
#include<cuda.h>
#include<cutil.h>
#include<time.h>


#define N 512


__global__ void vec_block_add(int * a, int * b, int * c) {
 c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];
}


void rand_ints(int * arr, int count) {
 srand(time(NULL));
 for(int i=0;i<count;i++) {
  arr[i] = rand() % 100;
 }
}


int main(void) {
 int * a,* b,* c;
 int * d_a, * d_b, * d_c;
 int size = N * sizeof(int);
 cudaMalloc((void**)&d_a,size);
 cudaMalloc((void**)&d_b,size);
 cudaMalloc((void**)&d_c,size);
 
 a = (int *) malloc(size);
 rand_ints(a,N);
 b = (int *) malloc(size);
 rand_ints(b,N);
 c = (int *) malloc(size);
 
 cudaMemcpy(d_a,a,size,cudaMemcpyHostToDevice);
 cudaMemcpy(d_b,b,size,cudaMemcpyHostToDevice);
 vec_block_add<<<N,1>>>(d_a,d_b,d_c);
 cudaMemcpy(c,d_c,size,cudaMemcpyDeviceToHost);
 
#if 1
 for(int i=0;i<N;i++) {
  printf("%-5d: a:%-5d b:%-5d c:%-5d\n",i,a[i],b[i],c[i]);
 }
#endif
 
 cudaFree(d_a);
 cudaFree(d_b);
 cudaFree(d_c);
 
 free(a);
 free(b);
 free(c);
 return 0;
}

例3中最关键的代码为如下几行:

__global__ void vec_block_add(int * a, int * b, int * c) {
 c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];
}

由于函数是并行执行的,和传统的串行程序在integer_add函数中使用循环来完成加法相比,相当于由GPU这个加速器使用硬件的方式进行了循环展开,展开后便可以并行执行了。所以在编写这段代码时,需要使用blockIdx.x来定位当前执行的是循环的哪个部分。

从硬件的角度看,相当于同时有多个块在并行执行:
块0: c[0]=a[0]+b[0]
块1: c[1]=a[1]+b[1]
块2: c[2]=a[2]+b[2]
块3: c[3]=a[3]+b[3]
....

更多详情见请继续阅读下一页的精彩内容http://www.linuxidc.com/Linux/2014-07/104328p2.htm

linux
本文评论   查看全部评论 (0)
表情: 表情 姓名: 字数

       

评论声明
  • 尊重网上道德,遵守中华人民共和国的各项有关法律法规
  • 承担一切因您的行为而直接或间接导致的民事或刑事法律责任
  • 本站管理人员有权保留或删除其管辖留言中的任意内容
  • 本站有权在网站内转载或引用您的评论
  • 参与本评论即表明您已经阅读并接受上述条款