《GPU高性能程式設計CUDA實戰》學習筆記 (1)

《GPU高性能程式設計CUDA實戰》中代碼整理

https://blog.csdn.net/fengbingchun/article/details/45954355

《GPU高性能程式設計CUDA實戰》學習筆記（三）

https://blog.csdn.net/w09103419/article/details/52484969

CUDA程式設計——GPU架構，由sp，sm，thread，block，grid，warp說起

https://blog.csdn.net/junparadox/article/details/50540602

《CUDA并行程式設計-GPU程式設計指南》讀書筆記–(1)線程網格、線程塊以及線程

https://blog.csdn.net/u011475134/article/details/71577770

SP（streaming Process），SM（streaming multiprocessor）是硬體（GPU hardware）概念。而thread，block，grid，warp是軟體上的（CUDA）概念。

• SP：最基本的處理單元，streaming processor，也稱為CUDA core。最後具體的指令和任務都是在SP上處理的。GPU進行并行計算，也就是很多個SP同時做處理。
• SM：多個SP加上其他的一些資源組成一個streaming multiprocessor。也叫GPU大核，其他資源如：warp scheduler，register，shared memory等。SM可以看做GPU的心髒（對比CPU核心），register和shared memory是SM的稀缺資源。CUDA将這些資源配置設定給所有駐留在SM中的threads。是以，這些有限的資源就使每個SM中active warps有非常嚴格的限制，也就限制了并行能力。

CUDA架構專門為GPU計算設計了一種全新的子產品，目的是減輕早期GPU計算中存在的一些限制，而正是這些限制使得之前的GPU在通用計算中沒有得到廣泛的應用。

使用CUDA C來編寫代碼的前提條件包括：

(1)、支援CUDA的圖形處理器，即由NVIDIA推出的GPU顯示卡，要求顯存超過256MB；

(2)、NVIDIA裝置驅動程式，用于實作應用程式與支援CUDA的硬體之間的通信，確定安裝最新的驅動程式，注意選擇與開發環境相符的圖形卡和作業系統；

(3)、CUDA開發工具箱即CUDA Toolkit，此工具箱中包括一個編譯GPU代碼的編譯器；

(4)、标準C編譯器，即CPU編譯器。CUDA C應用程式将在兩個不同的處理器上執行計算，是以需要兩個編譯器。其中一個編譯器為GPU編譯代碼，而另一個為CPU編譯代碼。

第三章 第一段CUDA C代碼+Host/Device

3.1 第一個程式

3.1.1 Hello World！

一般，将CPU以及系統的記憶體稱為主機(Host)，而将GPU及其記憶體稱為裝置(Device)。在GPU裝置上執行的函數通常稱為核函數(Kernel)。

#include<stdio.h>
#include <iostream>

__global__ void kernel(void) {
}

int main() {
  kernel <<<1, 1>>>();
  printf("hello world!");
  getchar();
  return 0;
}
           

這個程式與最初的“ Hello， World！”相比，多了兩個值得注意的地方：

• 一個空的函數kernel()，并且帶有修飾符__global__ 。這個修飾符将告訴編譯器，函數應該編譯為在裝置而不是主機上運作。在這個簡單的示例中，函數kernel()将被交給編譯裝置代碼的編譯器，而main()函數将被交給主機編譯器。CUDA編譯器和運作時将負責實作從主機代碼中調用裝置代碼。
• 對這個空函數的調用，并且帶有修飾字元 <<<1,1>>> 。

3.1.2 傳遞參數給核函數

#include <stdio.h>
#include <iostream>
#include <device_launch_parameters.h>

__global__ void add(int a, int b, int *c) {
  *c = a + b;
}

int main()
{
  int a = 2, b = 3, c = 0;
  int* dev_c;
  cudaMalloc((void**)&dev_c, sizeof(int));

  //尖括号表示要将一些參數傳遞給CUDA編譯器和運作時系統
  //尖括号中這些參數并不是傳遞給裝置代碼的參數，而是告訴運作時如何啟動裝置代碼,
  //傳遞給裝置代碼本身的參數是放在圓括号中傳遞的，就像标準的函數調用一樣
  add << <1, 1 >> > (a, b, dev_c);
  cudaMemcpy(&c, dev_c, sizeof(int), cudaMemcpyDeviceToHost);

  printf("%d + %d = %d\n", a, b, c);
  cudaFree(dev_c);
  getchar();
  return 0;
}
           

注意這裡增加了多行代碼，在這些代碼中包含兩個概念：

• 可以像調用C函數那樣将參數傳遞給核函數。

• 當裝置執行任何有用的操作時，都需要配置設定記憶體，例如将計算值傳回給主機。

通過cudaMalloc() 來配置設定記憶體。

這個函數調用的行為非常類似于标準的C函數malloc()，但該函數的作用是告訴CUDA運作時在裝置上配置設定記憶體。

• 第一個參數是一個指針，指向用于儲存新配置設定記憶體位址的變量，
• 第二個參數是配置設定記憶體的大小。除了配置設定記憶體的指針不是作為函數的傳回值外，這個函數的行為與malloc()是相同的，并且傳回類型為void*。

cudaMalloc函數使用限制總結：

1.可以将cudaMalloc()配置設定的指針傳遞給在裝置上執行的函數。

2.可以在裝置代碼中使用cudaMalloc()配置設定的指針進行記憶體讀/寫操作。

3.可以将cudaMalloc()配置設定的指針傳遞給在主機上執行的函數。

4.不能在主機代碼中使用cudaMalloc()配置設定的指針進行記憶體讀/寫操作。

要釋放cudaMalloc()配置設定的記憶體，需要調用cudaFree()

裝置指針的使用方式與标準C中指針的使用方式完全一樣。 主機指針隻能通路主機代碼中的記憶體，而裝置指針也隻能通路裝置代碼中的記憶體。

3.1.3 查詢裝置

有可能在單塊卡上包含了兩個或多個GPU。

在主機代碼中可以通過調用cudaMemcpy() 來通路裝置上的記憶體，在裝置和主機之間複制資料。

獲得CUDA裝置的數量，可以調用cudaGetDeviceCount()

int count;
HANDLE_ERROR( cudaGetDeviceCount( &count ) );
           

在內建的GPU上運作代碼，可以與CPU共享記憶體。

計算功能集的版本為1.3或者更高的顯示卡才能支援雙精度浮點數的計算。

3.1.4 裝置屬性的使用

cudaGetDeviceProperties()

查找主版本号大于1，或者主版本号為1且次版本号大于等于3的裝置：

首先，找出我們希望裝置擁有的屬性并将這些屬性填充到一個cudaDeviceProp結構。

cudaDeviceProp prop;
memset( &prop, 0, sizeof( cudaDeviceProp ) );
prop.major = 1;
prop.minor = 3;
           

在填充完 cudaDeviceProp 結構後，将其傳遞給 cudaChooseDevice()，這樣CUDA運作時将查找是否存在某個裝置滿足這些條件。 cudaChooseDevice()函數将傳回一個裝置ID，然後我們可以将這個ID傳遞給 cudaSetDevice()。随後，所有的裝置操作都将在這個裝置上執行。

#include "../common/book.h"
int main( void )
{
	cudaDeviceProp prop;
	int dev;
	HANDLE_ERROR( cudaGetDevice( &dev ) );
	printf( "ID of current CUDA device: %d\n", dev );
	memset( &prop, 0, sizeof( cudaDeviceProp ) );
	prop.major = 1;
	prop.minor = 3;
	HANDLE_ERROR( cudaChooseDevice( &dev, &prop ) );
	printf( "ID of CUDA device closest to revision 1.3: %d\n", dev );
	HANDLE_ERROR( cudaSetDevice( dev ) );
}