Transcript NVIDIA_CUDA_Programming

Compute Unified Device Architecture
(CUDA)
Overview
•
•
•
•
•
•
•
•
•
•
简介(Introduction)
并行编程(Parallel Programming)
线程协作(Thread Cooperation)
共享内存(Shared Memory)
常量内存(Constant Memory)
纹理内存(Texture Memory）
原子性(Atomics)
流(Streams)
零拷贝主机内存(Zero-Copy Memory)
建议(Suggestion)
Introdoction
CUDA是一种由NVIDIA推出的通用并行
计算架构，该架构使GPU能够解决复杂的
计算问题。
语言：C/C++
开发环境：支持cuda的GPU、NVIDIA设
备驱动程序、CUDA开发工具箱...
第一个程序
#include <iostream>
__global__ void add( int a, int b, int *c ) {
*c = a + b;
}
int main( void ) {
int c;
int *dev_c;
cudaMalloc( (void**)&dev_c, sizeof(int) ) ;
add<<<1,1>>>( 2, 7, dev_c );
cudaMemcpy( &c, dev_c, sizeof(int),
cudaMemcpyDeviceToHost ) ;
printf( "2 + 7 = %d\n", c );
cudaFree( dev_c ) ;
return 0;
}
第一个程序
使用 __device__ 限定符声明的函数：
在设备上执行；
仅可通过设备调用。
使用 __global__ 限定符可将函数声明为内核。
在设备上执行；
只能通过主机调用。
__global__ 函数的返回类型必须为空。
使用 __host__ 限定符声明的函数：
在主机上执行；
仅可通过主机调用。
第一个程序
cudaError_t cudaMalloc( void** devPtr，size_t count )
cudaError_t cudaMemcpy( void* dst，const void* src，
size_t count，enum cudaMemcpyKind kind )
cudaError_t cudaFree (void* devPtr)
cudaMemcpyKind ：
cudaMemcpyHostToHost
cudaMemcpyHostToDevice
cudaMemcpyDeviceToHost
cudaMemcpyDeviceToDevice
Parallel Programming
CPU:
for (int i=0; i<N; i++) {
c[i] = a[i] + b[i];
}
矢量求和
Parallel Programming
CUDA:
__global__ void add( int *a,int *b,int *c)
{
int tid = threadIdx.x;
if (tid < N)
c[tid] = a[tid] + b[tid];
}
cudaMemcpy( dev_a, a, N * sizeof(int),
cudaMemcpyHostToDevice );
cudaMemcpy( dev_b, b, N * sizeof(int),
cudaMemcpyHostToDevice );
int main( void ) {
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) );
cudaMemcpy( c, dev_c, N * sizeof(int),
cudaMemcpyDeviceToHost );
add<<<1,N>>>( dev_a, dev_b, dev_c );
cudaFree( dev_a );
cudaFree( dev_b );
cudaFree( dev_c );
return 0;
}
Parallel Programming
矩阵相加：将两个长度为N*N的矩阵A和B相加，然后
将结果写入矩阵C
__global__ void MatAdd(float **A, float **B, float **C){
int i = threadIdx.x;
int j = threadIdx.y;
C[i][j] = A[i][j] + B[i][j];
}
int main() {
...
dim3 threadsPerBlock(N, N);
MatAdd<<<1, threadsPerBlock>>>(A, B, C);
}
Parallel Programming
线程块中的所有线程必须存在于同一个处理器核心中且共享该
核心有限的存储器资源，因此，一个块内的线程数目是有限的
。
__global__ void MatAdd(float **A, float **B, float **C){
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() {
...
dim3 threadsPerBlock(32, 32);
dim3 numBlocks((N+threadsPerBlock.x)/threadsPerBlock.x,
(N+threadsPerBlock.y)/threadsPerBlock.y);
MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
...
}
Parallel Programming
类型dim3表示一个三维数组，可
以用于指定线程块或线程的数量
threadIdx、blockIdx可以看作线
程或线程块的索引
<<<...>>>指定执行某一指定内核
调用的线程数
Grid of Thread Blocks
Parallel Programming
回顾矢量加法的例子：当矢量维数N特别大时....
__global__ void add( int *a, int *b, int *c ) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
while (tid < N) {
c[tid] = a[tid] + b[tid];
tid += blockDim.x * gridDim.x;
}
}
int main( void ) {
...
add<<<128,128>>>( dev_a, dev_b, dev_c );
...
}
Thread Cooperation/Shared Memory
线程同步：__syncthreads()
共享内存：使用关键字__shared__进行声明，由同一
线程块所有线程共享，使得同一线程块中的线程能够
通信和协作。并且，共享内存缓冲区驻留在物理GPU
上，访问延迟更小。
块内线程可通过共享存储器和同步执行协作，共享存
储器可以共享数据，同步执行可以协调存储器访问。
Thread Cooperation/Shared Memory
点积运算
：
__global__ void dot( float *a, float *b, float *c ) {
__shared__ float cache[threadsPerBlock];
int tid = threadIdx.x + blockIdx.x * blockDim.x;
int cacheIndex = threadIdx.x;
float temp = 0;
if(tid < N) {
temp = a[tid] * b[tid];
}
cache[cacheIndex] = temp;
__syncthreads();
int i = blockDim.x/2;
while (i != 0) {
if (cacheIndex < i)
cache[cacheIndex] +=
cache[cacheIndex + i];
__syncthreads();
i /= 2;
}
if (cacheIndex == 0)
c[blockIdx.x] = cache[0];
}
Thread Cooperation/Shared Memory
归约计算
Constant Memory
用于保存在核函数执行期间不会发生变化的数
据，某些情况中，能有效减少内存带宽。
使用__constant__声明，且变量访问限制为只
读
节约内存带宽：
1、对常量内存的单次读操作可以广播到其他邻
近的线程
2、常量内存将数据缓存起来，因此对相同地址
的连续操作不会产生额外的内存通信量
Texture Memory
适用于那些内存访问模式中存在大量空间局部性
的图形应用程序。
Atomics
计算直方图：假设数组buffer长度为N，数组元素为大于
0小于256的整数，求0～255分布直方图，分布结果放
在result数组中。
CPU：
for (int i=0; i<256; i++)
result[i] = 0;
for (int i=0; i<N; i++)
result[buffer[i]]++;
CUDA：
__global__ void histo_kernel( unsigned char *buffer,
long N,
unsigned int *result)
{
int i = threadIdx.x + blockIdx.x * blockDim.x;
if(i < N) {
//result[buffer[i]]++;
atomicAdd(
&result[buffer[i]], 1 );
}
}
Atomics
Normal
Atomics
Error
Atomics
atomicAdd()：将加法运算后的值写入原地址
atomicSub()：将减法运算后的值写入原地址
atomicExch()：将给定值写入原地址
atomicMin()：把原来的值以及给定的值中较小的
一个写入原地址
atomicMax()：把原来的值以及给定的值中较大的
一个写入原地址
...
Streams
CUDA流表示一个GPU操作队列，并且该队列中的操
作将以指定的顺序执行，我们可以在流中添加一些操
作，例如内核启动、内存复制等等。
作用：实现多任务并行。通过建立多条工作流可以使
多个不同任务在GPU同时进行。
Zero-Copy Memory
页锁定主机内存：使用cudaHostAlloc()进行分配内存
，使用cudaFreeHost()释放内存。GPU可以直接使用
DMA(直接内存访问)在GPU和主机间复制数据。
cudaError_t cudaHostAlloc (void ** pHost, size_t size, unsigned int flags)
cudaError_t cudaHostGetDevicePointer (void**pDevice, void*pHost, 0)
cudaError_t cudaFreeHost (void * ptr)
cudaHostAlloc flags:
cudaHostAllocMapped:核函数可以直接访问这种类型的主机内存，不需
要将数据复制到GPU，因此称为零拷贝内存。
cudaHostAllocDefault
cudaHostAllocPortable
cudaHostAllocWriteCombined
Suggestion
1、减少Host和Device之间的数据拷贝。做到多次使用
同一数据时只拷贝一次。
例子：有两个矢量a、b，（1）求出矢量和，把结果放在数
组c中；（2）进行点积运算，求出结果放在result中。
由于两个小题中输入数据都是数组a、b，并且在两小题中数
组a、b中的数据都不会更改，因此，当使用cudaMemcpy()
把数据a、b的数据传输到GPU内存后，GPU内存中数组a、
b的值也会一直保持不变。因此执行完第一小题的要求后不
要马上使用cudaFree()释放内存，而是继续使用这块内存上
的数据进行点积运算，当完成第二小题之后才将a、b在
GPU内存上的数据释放。这样可以避免两次向GPU传输同
样的数据，节省时间。
Suggestion
2、最大限度并行化。当然，thread也不是越多越好
。
使用CUDA能够使程序加速的主要原因就在于GPU能够执
行大量的并行运算，因此一般来说并行程度越高，速度越
快。
for(int i=0;i<N;i++)
for(int j=0;j<M;j++)
for(int k=0;k<X;k++)
for(int t=0;t<Y;t++)
{
...
}
for(int n=0;n<N*M;n++)
{
int i=n/M;
int j=n%M;
for(int m=0;m<X*Y;m++)
{
int k=m/Y;
int t=m%Y;
}
}
Suggestion
3、存储优化，尽量使用share memory、
constant memory或者texture memory取代global
memory对数据进行访问。
例子：利用sobel算子进行图像边缘检测
，可以使用纹理内存进行数据存储。
Suggestion
4、注意条件分支语句，尽量少用分支；展开小
循环。
__global__ void test(int*a){
int count=0;
for(int i=0;i<5;i++)
{
if(a[i]<0)
count++;
}
}
__global__ void
test(int*a){
int count=0;
if(a[0]<0)
count++;
if(a[1]<0)
count++;
if(a[2]<0)
count++;
if(a[3]<0)
count++;
if(a[4]<0)
count++;
}
Conclusion
•
•
•
•
•
•
•
•
•
•
简介(Introduction)
并行编程(Parallel Programming)
线程协作(Thread Cooperation)
共享内存(Shared Memory)
常量内存(Constant Memory)
纹理内存(Texture Memory）
原子性(Atomics)
流(Streams)
零拷贝主机内存(Zero-Copy Memory)
建议(Suggestion)
Thank you !