GPU(CUDA)学习日记（三）------ CUDA基本架构介绍以及编程入门！~~

首先，介绍CUDA的架构：在CUDA架构下，线程的最小单元是thread，多个thread组成一个block，多个block再组成一个grid，不同block之间的thread不能读写同一shared memory共享内存，因此，block里面的thread之间的通信和同步所带来的开销是比较大的。SM以 32 个 Thread 为一组的 Warp 来执行 Thread。Warp内的线程是静态的，即在属于同一个warp内的thread之间进行通信，不需要进行栅栏同步(barrier)。Fermi的设计根据G80和GT200的架构作出的很多缺陷来改变。在Fermi中，每个SM中的数量不再是GT200的8个SP，而是变成了32个SP，NVIDIA现在又称之为CUDA Core，总共具有16个SM，所以总共有512个SP。而在GT200中，是30个SM，240个SP。

    其次，介绍CUDA的初始化：首先，先建立一个档案 first_cuda.cu。如果是使用 Visual Studio 的话，则请先按照这里的设定方式设定 project。要使用 runtime API 的时候，需要 include cuda_runtime.h。所以，在程序的最前面，加上

#include <stdio.h>

 #include <cuda_runtime.h>

接下来是一个 InitCUDA 函式，会呼叫 runtime API 中，有关初始化 CUDA 的功能：
bool InitCUDA()

 {

      int count;

      cudaGetDeviceCount(&count); 

      if(count == 0) 

      { 

         fprintf(stderr, "There is no device.\n"); 

         return false; 

       }

 int i; 

 for(i = 0; i < count; i++) 

 { 

    cudaDeviceProp prop;

     if(cudaGetDeviceProperties(&prop, i) == cudaSuccess) 

       { 

            if(prop.major >= 1)

                 {  

                    break; 

                 }

        }

  } 

  if(i == count) 

   { 

       fprintf(stderr, "There is no device supporting CUDA 1.x.\n");

       return false;

    } 

   cudaSetDevice(i); 

   return true;

 } 

    这个函式会先呼叫 cudaGetDeviceCount 函式，取得支持 CUDA 的装置的数目。如果系统上没有支持 CUDA 的装置，则它会传回 1，而 device 0 会是一个仿真的装置，但不支持 CUDA 1.0 以上的功能。所以，要确定系统上是否有支持 CUDA 的装置，需要对每个 device 呼叫cudaGetDeviceProperties 函式，取得装置的各项数据，并判断装置支持的 CUDA 版本（prop.major 和 prop.minor 分别代表装置支持的版本号码，例如 1.0 则 prop.major 为 1 而prop.minor 为 0）。

    nvcc 是 CUDA 的 compile 工具，它会将 .cu 檔拆解出在 GPU 上执行的部份，及在 host 上执行的部份，并呼叫适当的程序进行 compile 动作。在 GPU 执行的部份会透过 NVIDIA 提供的 compiler 编译成中介码，而 host 执行的部份则会透过系统上的 C++ compiler 编译（在Windows 上使用 Visual C++ 而在 Linux 上使用 gcc）。

    最后，来看一个例子，这个例子对于我来说很有帮助！

在主函数中，

int* gpudata, *result; clock_t* time; 

cudaMalloc((void**) &gpudata, sizeof(int) * DATA_SIZE);      // 此句子用来开辟GPU的空间；

cudaMalloc((void**) &result, sizeof(int) * THREAD_NUM * BLOCK_NUM); 

cudaMalloc((void**) &time, sizeof(clock_t) * BLOCK_NUM * 2);    //此句为定义了一个CUDA提供的计算时间的方法，新手忽略~

cudaMemcpy(gpudata, data, sizeof(int) * DATA_SIZE, cudaMemcpyHostToDevice);   

//此句用来将数据data以gpudata的名字送入GPU中，其中cudaMemcpyHostToDevice表示从CPU传入GPU；

sumOfSquares<<<BLOCK_NUM, THREAD_NUM, 0>>>(gpudata, result, time); 

int sum[THREAD_NUM * BLOCK_NUM]; clock_t time_used[BLOCK_NUM * 2]; 

cudaMemcpy(&sum, result, sizeof(int) * THREAD_NUM * BLOCK_NUM, cudaMemcpyDeviceToHost); 

cudaMemcpy(&time_used, time, sizeof(clock_t) * BLOCK_NUM * 2, cudaMemcpyDeviceToHost); 

cudaFree(gpudata); 

cudaFree(result);

cudaFree(time); 

int final_sum = 0;

 for(int i = 0; i < THREAD_NUM * BLOCK_NUM; i++) 

{ 

    final_sum += sum[i];

 }

clock_t min_start, max_end; 

min_start = time_used[0];

 max_end = time_used[BLOCK_NUM]; 

for(int i = 1; i < BLOCK_NUM; i++) 

{ 

    if(min_start > time_used[i]) min_start = time_used[i]; 

    if(max_end < time_used[i + BLOCK_NUM]) max_end = time_used[i + BLOCK_NUM];}

printf("sum: %d time: %d\n", final_sum, max_end - min_start);

这个程序例子中，主要的东西是主函数中的调用方法，具体要做的是：

1、定义一个内存指针，并为自己将要传到GPU函数中的变量开辟一个空间（cudaMalloc()）；

2、利用cudaMemcpy()将数据从CPU中传到GPU中；

3、定义内核函数，注意内核函数的调用方式以及获得线程序号的方法~