CUDA:CUDAの理解

1.GPUの理解

演算能力を向上させるために、「より多くの、簡単な計算ユニット」

を使うのが好きです.

CPUが解決した問題はLatencyであり、各タスクは最短でどのくらいの時間で

を完成できるかである.

GPUはThroughPutを解決して、1単位の時間ごとにどれだけの任務を解決することができます

GPUは効率的な同時

を得意とする

並列実行大量スレッド

2.CUDA計算モデル3.典型的なGPUプログラム

CPUはGPUにメモリ空間CUDA MALLOC

を割り当てる.

CPUコピー入力データCPU->GPU CUDA memcpy

CPU launches kernel on GPUデータkernel Launch

を計算する

CPUはコピー要求を送信し、GPU->CPU CUDA memcpy

上記の2,4ステップはいずれもデータの転送であり,いずれも非常に高価である.
4. BIG IDEA
Kernels look serial programs.Write your program as if it will run on one thread.The GPU will run that program on many threads
5.最も簡単なCUDAプログラム
これは1組の数字の平方を計算するCUDAプログラムです.global__修飾子はCUDA関数でありkernel functionとも呼ばれる.
全体のプログラムは簡単で分かりやすくて、注釈ははっきりしていて、あまり説明しません

#include <stdio.h>   
  
__global__ void square(float * d_out, float * d_in)
{  
    int idx = threadIdx.x;  
    float f = d_in[idx];  
    d_out[idx] = f * f;  
}  
  
int main(int argc, char ** argv) 
{  
    const int ARRAY_SIZE = 64;  
    const int ARRAY_BYTES = ARRAY_SIZE * sizeof(float);  
      
    // generate the input array on the host   
    float h_in[ARRAY_SIZE];  
    for (int i = 0; i < ARRAY_SIZE; i++)
    {  
        h_in[i] = float(i);  
    }  
    float h_out[ARRAY_SIZE];  
      
    // declare GPU memory pointers   
    float *d_in;  
    float *d_out;  
      
    // allocate GPU memory   
    cudaMalloc((void**) &d_in, ARRAY_BYTES);  
    cudaMalloc((void**) &d_out, ARRAY_BYTES);  
      
    // transfer the array to the GPU   
    cudaMemcpy(d_in, h_in, ARRAY_BYTES, cudaMemcpyHostToDevice);  
      
    // launch the kernel   
    square<<<1, ARRAY_SIZE>>>(d_out, d_in);  
      
    // copy back the result array to the CPU   
    cudaMemcpy(h_out, d_out, ARRAY_BYTES, cudaMemcpyDeviceToHost);  
      
    // print out the resulting array   
    for (int i =0; i < ARRAY_SIZE; i++) {  
        printf("%f", h_out[i]);  
        printf(((i % 4) != 3) ? "\t" : "
");  
    }  
    cudaFree(d_in);  
    cudaFree(d_out);  
      
    return 0;  
}  

6.BLOCKとGRID

多くのスレッドがBLOCKに整理され、threadIdx変数はスレッドがどのスレッドであるかを示します.一方、BLOCKには、1つのBLOCKが最大1024スレッドを含むことができるという制限があります.どうすればいいですか.私たちは複数のBLOCKを使って、複数のBLOCKをどのように組織することができますか.GRIDを使用!

GRIDとBLOCKの関係はBLOCKとTHREADの関係と類似している.

BLOCKおよびGRIDはintタイプであってもよいしdim 3タイプ

であってもよい

<<<...>>> シンボルは、BLOCKおよびGRID

の使用量を指定する

BLOCKは、1次元、2次元、3次元の

であることができる.

kernelはblockIdx変数によって現在そのBLOCK

を取得する.

BLOCKの次元数はblockDimという変数によって

を得ることができる.

組織構造図

Memory
1)Local memory:スレッド内の変数が占めるメモリ
2)shared memory:1つのBLOCK内部のすべてのスレッドが共有するメモリ【注意:_shared__修飾子はこのBLOCK共有メモリを表す】
3)global memory:全スレッド共有メモリ
速度1)2より大きい)3よりはるかに大きい)
ここに練習問題があります:次のこのコード、1,2,3,4行の中で、各行の演算が最も速くて、どの行が最も遅くて、それぞれ各行に点数をつけて、点数は1～4です

__global__ void foo(float *x,float *y,float *z)
{  
    __shared__ float a,b,c;  
    float s,t,u;  
    s = *x;//1   
    t = s;//2   
    a = b;//3   
    *y = *z;//4   
}  

__global__ void foo(float *x,float *y,float *z){
    __shared__ float a,b,c;
    float s,t,u;
    s = *x;//1
    t = s;//2
    a = b;//3
    *y = *z;//4
}

答えは:これによって1の速度は3で、2行の速度は1で、3は2で、4は4です
同期(Synchronise)
コードに障害を設けるように
次に、array[x]=array[x+1]を前に移動したい配列があります.
簡単な関数フレームワークは次のようになります.

__global__ void shift(){  
    int idx = threadIdx.x;  
    __shared__ int array[128];  
    array[idx] = threadIdx.x;  
    if (idx < 127) {  
        array[idx] = array[idx + 1];  
    }  
}  

__global__ void shift(){
    int idx = threadIdx.x;
    __shared__ int array[128];
    array[idx] = threadIdx.x;
    if (idx < 127) {
        array[idx] = array[idx + 1];
    }
}

並行実行を少し理解すると、上記のコードがうまく実行できないことがわかります.例えば、スレッド20が実行する場合、
array[20] = 20;実行可能
array[20] = array[20+1];スレッド21がスレッド20の後に実行される場合、この文の付与は間違いであることは明らかである.
では、上記のプログラムを正しく実行するには、コードにいくつかの「障害」を設定する必要があります.

__global__ void shift(){
    int idx = threadIdx.x;
    __shared__ int array[128];
    array[idx] = threadIdx.x;
    __syncthreads();// ， 
    if (idx < 127) {
        int temp = array[idx + 1];
        __syncthreads();// barrier, ， ， 
        array[idx] = temp;
        __syncthreads();// array 
    }
    
}

__global__ void shift(){  
    int idx = threadIdx.x;  
    __shared__ int array[128];  
    array[idx] = threadIdx.x;  
    __syncthreads();// ，    
    if (idx < 127) {  
        int temp = array[idx + 1];  
        __syncthreads();// barrier, ， ，    
        array[idx] = temp;  
        __syncthreads();// array    
    }  
      
}