Caffe 初学拾遗(六) CUDA 线程通信

Original Source:

http://blog.csdn.net/augusdi/article/details/12833235

一些CUDA编程的简单示例程序，笔者在此进行了整理说明：

1.线程通信：

理论上，各个线程所进行的处理是互不相关的，即两个线程不会产生交集（向量点积，向量加运算）。

在实际中，各个线程不再是相互独立的，而是具有一定关联，线程2可能会用到线程1的结果，这时需要利用线程通信。

线程通信在CUDA中有三种实现方式：
1. 共享存储器；
2. 线程同步；
3. 原子操作；
最常用的是前两种方式。

Shared Memory，是位于SM（流多处理器）中的特殊存储器。

一个SM中不仅包含若干个SP（流处理器），还包括一部分高速Cache，寄存器组，共享内存等。

一个SM内有M个SP，Shared Memory由这M个SP共同占有。

指令单元也被这M个SP共享，即SIMT架构（单指令多线程架构），一个SM中所有SP在同一时间执行同一代码。

需要同步机制才能使线程之间实现有序处理。

当线程A需要线程B计算的结果作为输入时，需要确保线程B已经将结果写入共享内存中，然后线程A再从共享内存中读出。

同步机制可以用CUDA内置函数：

__syncthreads()：

当某个线程执行到该函数时，进入等待状态，直到同一线程块（Block）中所有线程都执行到这个函数为止。

即一个__syncthreads()相当于一个线程同步点，确保一个Block中所有线程都达到同步，然后线程进入运行状态。

NOTE：

由于SIMT特性，所有线程都执行同样的代码，所以在线程中需要判断自己的ThreadID，以免误操作。
位于同一个Block中的线程才能利用Shared Memory实现通信。

不同Block中的线程不能通过共享内存、同步进行通信，而应采用原子操作或主机介入方式通信。

E.g.

Aim：

分别求出1~5这5个数字的和，平方和，连乘积。

Introduction：

输入数据原本位于Host Memory，通过cudaMemcpy API拷贝到 Device Memory（也被称为Global Memory，泛指显存），

每个线程运行时需要从Global Memory读取输入数据，然后完成计算，最后将结果写回Global Memory。

当我们计算需要多次相同输入数据时，反复多次读数据会耗费大量时间（一次访问需几十甚至上百个时钟周期）。

可以将数据从Global Memory一次性读到SM内部的Shared Memory(一次访问需0.5个时钟周期)，然后在内部进行处理，这样可以节省反复读取的时间。

#include "cuda_runtime.h"  
#include "device_launch_parameters.h"  
  
#include <stdio.h>  
  
cudaError_t addWithCuda(int *c, const int *a, size_t size);  
  
__global__ void addKernel(int *c, const int *a)  
{  
    int i = threadIdx.x;  
	extern __shared__ int smem[];
	smem[i] = a[i];  
    __syncthreads();  
    if(i == 0)  // 0号线程做平方和  
    {  
        c[0] = 0;  
        for(int d = 0; d < 5; d++)  
        {  
            c[0] += smem[d] * smem[d];  
        }  
    }  
    if(i == 1)  // 1号线程做累加  
    {  
        c[1] = 0;  
        for(int d = 0; d < 5; d++)  
        {  
            c[1] += smem[d];  
        }  
    }  
    if(i == 2)  // 2号线程做累乘  
    {  
        c[2] = 1;  
        for(int d = 0; d < 5; d++)  
        {  
            c[2] *= smem[d];  
        }  
    }  
	// 从示例来看 申请5个线程 只有3个在执行
	// Thread ID = 4/5 的在同步函数处等待
}  
  
int main()  
{  
    const int arraySize = 5;  
    const int a[arraySize] = { 1, 2, 3, 4, 5 };  
    int c[arraySize] = { 0 };  
    // Add vectors in parallel.  
    cudaError_t cudaStatus = addWithCuda(c, a, arraySize);  
    if (cudaStatus != cudaSuccess)   
    {  
        fprintf(stderr, "addWithCuda failed!");  
        return 1;  
    }  
    printf("\t1+2+3+4+5 = %d\n\t1^2+2^2+3^2+4^2+5^2 = %d\n\t1*2*3*4*5 = %d\n\n\n\n\n\n", c[1], c[0], c[2]);  
    // cudaThreadExit must be called before exiting in order for profiling and  
    // tracing tools such as Nsight and Visual Profiler to show complete traces.  
    cudaStatus = cudaThreadExit();  
    if (cudaStatus != cudaSuccess)   
    {  
        fprintf(stderr, "cudaThreadExit failed!");  
        return 1;  
    }  
    return 0;  
}  
  
// Helper function for using CUDA to add vectors in parallel.  
cudaError_t addWithCuda(int *c, const int *a,  size_t size)  
{  
    int *dev_a = 0;  
    int *dev_c = 0;  
    cudaError_t cudaStatus;  
  
    // Choose which GPU to run on, change this on a multi-GPU system.  
    cudaStatus = cudaSetDevice(0);  
    if (cudaStatus != cudaSuccess)   
    {  
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");  
        goto Error;  
    }  
  
    // Allocate GPU buffers for three vectors (two input, one output)    .  
    cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));  
    if (cudaStatus != cudaSuccess)   
    {  
        fprintf(stderr, "cudaMalloc failed!");  
        goto Error;  
    }  
  
    cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));  
    if (cudaStatus != cudaSuccess)   
    {  
        fprintf(stderr, "cudaMalloc failed!");  
        goto Error;  
    }  
    // Copy input vectors from host memory to GPU buffers.  
    cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);  
    if (cudaStatus != cudaSuccess)   
    {  
        fprintf(stderr, "cudaMemcpy failed!");  
        goto Error;  
    }  
    // Launch a kernel on the GPU with one thread for each element.  
	addKernel<<<1, size, size * sizeof(int), 0>>>(dev_c, dev_a);  
	// size * seizof(int) : 共享内存大小(字节)

    // cudaThreadSynchronize waits for the kernel to finish, and returns  
    // any errors encountered during the launch.  
    cudaStatus = cudaThreadSynchronize();  
    if (cudaStatus != cudaSuccess)   
    {  
        fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus);  
        goto Error;  
    }  
  
    // Copy output vector from GPU buffer to host memory.  
    cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);  
    if (cudaStatus != cudaSuccess)   
    {  
        fprintf(stderr, "cudaMemcpy failed!");  
        goto Error;  
    }  
  
    cudaFree(dev_c);  
    cudaFree(dev_a);      
    return cudaStatus;  
}  

｛笔者注：此示例程序效率并不高，一个Block拆分出多Thread以及多task并不推荐，只作示例｝

2.性能提升

问题规模不是很大，那么采用Thread并行是比较合适的。

大问题分多个Block处理时，每个Block内Thread数目不宜太少，否则将对硬件资源的造成浪费。

同时，每个Block最多只能占用一个SM，因此，需要分配任务到多个Block上才能发挥出GPU的能力。

一个理想的方案是，N个Block并行，每个Block包含512个Thread，将问题分解处理，效率往往比单一的线程并行处理或单一块并行处理高很多。