Caffe 초학 픽업 (6) 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 에서 SM 내부 의 Shared Memory (한 번 방문 시 0.5 시간 주기) 로 한 번 에 읽 고 내부 에서 처리 하면 반복 적 으로 읽 는 시간 을 절약 할 수 있다.

#include "cuda_runtime.h"  
#include "device_launch_parameters.h"  
  
#include   
  
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
\t1^2+2^2+3^2+4^2+5^2 = %d
\t1*2*3*4*5 = %d





", 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!
", 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 내 Thread 의 수량 이 너무 적 으 면 안 됩 니 다. 그렇지 않 으 면 하드웨어 자원 에 낭 비 를 초래 할 것 입 니 다.
또한 블록 당 SM 을 최대 1 개 만 차지 할 수 있 기 때문에 여러 블록 에 임 무 를 할당 해 야 GPU 의 능력 을 발휘 할 수 있다.
하나의 이상 적 인 방안 은 N 개의 Block 을 병행 하 는 것 이다. 각 Block 은 512 개의 Thread 를 포함 하고 문 제 를 분해 처리 하 며 효율 은 단일 한 스 레 드 병행 처리 나 단일 블록 병행 처리 보다 훨씬 높다.