PyCUDA 文檔中文翻譯-Tutorial 中文版

說起來很逗樂，我今天大概弄完了杜克大學STA663全部課件的中文翻譯，然後想著翻譯一下 PyCUDA 的文檔，因為最近在弄建模的一些事情，需要用到。

然後我就上網搜一下，看看是不是有現成的翻譯：

當時我就驚呆了，原來我自己去年（2016-10-13）就翻譯過 PyCUDA 的 Tutorial ，而我自己卻忘了。

所以就趁機開一個新坑：PyCUDA-文檔中文翻譯項目。

下面的就是最基礎的第一部分 Tutorial 的中文譯稿，不過是去年的，等我抽空再潤色潤色啥的吧。

唉，自己開的坑多了，居然不知道自己翻譯過這部分。。。

人生啊，真是太逗樂了。

PyCUDA Tutorial 英文原文

CycleUser 翻譯

開始使用

在你使用PyCuda之前，要先用import命令來導入並初始化一下。

import pycuda.driver as cudaimport pycuda.autoinitfrom pycuda.compiler import SourceModule

這裡要注意，你並不是必須使用pycuda.autoinit,初始化、內容的創建和清理也都可以手動實現。

轉移數據

接下來就是要把數據轉移到設備（device）上了。一般情況下，在使用PyCuda的時候，原始數據都是以NumPy數組的形式存儲在宿主系統（host）中的。（不過實際上，只要符合Python緩衝區介面的數據類型就都可以使用的，甚至連字元串類型str都可以。）

（譯者註：宿主系統host，就是處理器-內存-外存組成的常規Python運行環境;設備device，就是你要拿來做CUDA運算的顯卡或者運算卡，可以是單卡也可以是陣列。）

下面這行示例代碼創建了一個隨機數組成的4*4大小的數組a：

import numpya = numpy.random.randn(4,4)

不過要先暫停一下—咱們剛剛創建的這個數組a包含的是雙精度浮點數，但大多數常用的NVIDIA顯卡只支持單精度浮點數，所以需要轉換一下類型：

（譯者註：原作者的這篇簡介主要針對使用常規普通顯卡的用戶，比如GeForce系列的各種大家平時用到的都是這個範圍的，相比專門的計算卡，在雙精度浮點數等方面進行了閹割，所以作者才建議轉換類型到單精度浮點數。如果你使用的是專門的計算卡，就不用這樣了。）

a = a.astype(numpy.float32)

接下來，要把已有的數據轉移過去，還要設定一個目的地，所以我們要在顯卡中分配一段顯存：

（譯者註：原文說的是設備，這裡就直接說成顯卡了，畢竟大家用顯卡的比較多。另外下面這個代碼中的a.nbytes是剛剛生成的數組a的大小，這裡作者是按照數組大小來分配的顯存，新入門的用戶要注意這裡，後續的使用中，顯存的高效利用是很重要的。）

a_gpu = cuda.mem_alloc(a.nbytes)

最後，咱們把剛剛生成的數組a轉移到GPU裡面吧：

cuda.memcpy_htod(a_gpu, a)

運行一個內核函數（kernel）

咱們這篇簡介爭取說的都是最簡單的內容：咱們寫一個代碼來把a_gpu這段顯存中存儲的數組的每一個值都乘以2. 為了實現這個效果，我們就要寫一段CUDA C代碼，然後把這段代碼提交給一個構造函數，這裡用到了pycuda.compiler.SourceModule:

mod = SourceModule(""" __global__ void doublify(float *a) { int idx = threadIdx.x + threadIdx.y*4; a[idx] *= 2; } """)

（譯者註：上面這段代碼需要對C有一定了解。這裡的threadIDx是CUDA C語言的內置變數，這裡藉此確定了a數組所在的位置，然後通過指針對數組中每一個元素進行了自乘2的操作。）

這一步如果沒有出錯，就說明這段代碼已經編譯成功，並且載入到顯卡中了。然後咱們可以使用一個到咱們這個pycuda.driver.Function的引用，然後調用此引用，把顯存中的數組a_gpu作為參數傳過去，同時設定塊大小為4x4：

func = mod.get_function("doublify")func(a_gpu, block=(4,4,1))

最後，咱們就把經過運算處理過的數據從GPU取回，並且將它和原始數組a一同顯示出來對比一下：

a_doubled = numpy.empty_like(a)cuda.memcpy_dtoh(a_doubled, a_gpu)print (a_doubled)print (a)

（譯者註：原作者在原文中使用的是Python2，我這裡用的是帶括弧的print，這樣同時能在python2和python3上運行。）

輸出的效果大概就是如下所示：

（譯者註：上面這個是從顯存中取回的翻倍過的數組，下面的是原始數組。）

[[ 0.51360393 1.40589952 2.25009012 3.02563429] [-0.75841576 -1.18757617 2.72269917 3.12156057] [ 0.28826082 -2.92448163 1.21624792 2.86353827] [ 1.57651746 0.63500965 2.21570683 -0.44537592]][[ 0.25680196 0.70294976 1.12504506 1.51281714] [-0.37920788 -0.59378809 1.36134958 1.56078029] [ 0.14413041 -1.46224082 0.60812396 1.43176913] [ 0.78825873 0.31750482 1.10785341 -0.22268796]]

出現上面這樣輸出就說明成功了！整個攻略就完成了。另外很值得慶幸的是，運行輸出之後PyCuda就會把所有清理和內存回收工作做好了，咱們的簡介也就完畢了。不過你可以再看一下接下來的內容，裡面有一些有意思的東西。

(本文的代碼在PyCuda源代碼目錄下的examples/demo.py文件中。)

簡化內存拷貝

PyCuda提供了pycuda.driver.In, pycuda.driver.Out, 以及pycuda.driver.InOut 這三個參數處理器（argument handlers），能用來簡化內存和顯存之間的數據拷貝。例如，咱們可以不去創建一個a_gpu，而是直接把a移動過去，下面的代碼就可以實現：

func(cuda.InOut(a), block=(4, 4, 1))

有準備地調用函數

使用內置的 pycuda.driver.Function.call() 方法來進行的函數調用，會增加類型識別的資源開銷（參考顯卡介面）。 要實現跟上面代碼同樣的效果，又不造成這種開銷，這個函數就需要設定好參數類型（如Python的標準庫中的結構體模塊struct所示），然後再去調用該函數。這樣也就不用需要再使用numpy.number類去制定參數的規模了：

grid = (1, 1)block = (4, 4, 1)func.prepare("P")func.prepared_call(grid, block, a_gpu)

抽象以降低複雜度

使用 pycuda.gpuarray.GPUArray，同樣效果的代碼實現起來就更加精簡了：

import pycuda.gpuarray as gpuarrayimport pycuda.driver as cudaimport pycuda.autoinitimport numpya_gpu = gpuarray.to_gpu(numpy.random.randn(4,4).astype(numpy.float32))a_doubled = (2*a_gpu).get()print a_doubledprint a_gpu

進階內容

結構體

（由Nicholas Tung提供，代碼在examples/demo_struct.py文件中)

假如我們用如下的構造函數，對長度可變的數組的每一個元素的值進行翻倍：

mod = SourceModule(""" struct DoubleOperation { int datalen, __padding; // so 64-bit ptrs can be aligned float *ptr; }; __global__ void double_array(DoubleOperation *a) { a = &a[blockIdx.x]; for (int idx = threadIdx.x; idx datalen; idx += blockDim.x) { a->ptr[idx] *= 2; } } """)

網格grid中的每一個塊block（這些概念參考CUDA的官方文檔）都將對各個數組進行加倍。for循環允許比當前線程更多的數據成員被翻倍，當然，如果能夠保證有足夠多的線程的話，這樣做的效率就低了。接下來，基於這個結構體進行封裝出來的一個類就產生了，並且有兩個數組被創建出來：

class DoubleOpStruct: mem_size = 8 + numpy.intp(0).nbytes def __init__(self, array, struct_arr_ptr): self.data = cuda.to_device(array) self.shape, self.dtype = array.shape, array.dtype cuda.memcpy_htod(int(struct_arr_ptr), numpy.getbuffer(numpy.int32(array.size))) cuda.memcpy_htod(int(struct_arr_ptr) + 8, numpy.getbuffer(numpy.intp(int(self.data)))) def __str__(self): return str(cuda.from_device(self.data, self.shape, self.dtype))struct_arr = cuda.mem_alloc(2 * DoubleOpStruct.mem_size)do2_ptr = int(struct_arr) + DoubleOpStruct.mem_sizearray1 = DoubleOpStruct(numpy.array([1, 2, 3], dtype=numpy.float32), struct_arr)array2 = DoubleOpStruct(numpy.array([0, 4], dtype=numpy.float32), do2_ptr)print("original arrays", array1, array2)

上面這段代碼使用了pycuda.driver.to_device() 和 pycuda.driver.from_device() 這兩個函數來分配內存和複製數值，並且演示了在顯存中如何利用從已分配塊位置進行的偏移。最後咱們執行一下這段代碼；下面的代碼中演示了兩種情況：對兩個數組都進行加倍，以及只加倍第二個數組：

func = mod.get_function("double_array")func(struct_arr, block = (32, 1, 1), grid=(2, 1))print("doubled arrays", array1, array2)func(numpy.intp(do2_ptr), block = (32, 1, 1), grid=(1, 1))print("doubled second only", array1, array2, "
")

接下來的征程

當你對這些基礎內容感到足夠熟悉了，就可以去深入探索一下顯卡介面。更多的例子可以再PyCuda的源碼目錄下的examples子目錄。這個文件夾裡面也包含了一些測試程序，可以用來比對GPU和CPU計算的差別。另外PyCuda源代碼目錄下的test子目錄裡面由一些關於功能如何實現的參考。