前言

  本文旨在為讀者介紹CUDA，並制定相應的學習方案。本文以三個問題作為開篇：CUDA是什麼，CUDA被應用在哪些方面，CUDA能做什麼。簡言之，CUDA是一種實現GPU運算的技術，目前這項技術被廣泛應用於影像處理領域和統計AI領域。利用CUDA技術，配合適當的軟體就可以實現降低運算耗時。
  NVIDIA在營銷的時候，往往將編譯器與架構混合推廣；實際上，CUDA是可以相容OpenCL和NVCC編譯器。無論是CUDA C語言或是OpenCL,指令最終都會被驅動程式轉換成PTX程式碼，就像為CPU而寫的程式碼都將轉成彙編程式碼。

CUDA是如何工作的

  CUDA的工作過程是這樣的，CPU端準備好資料，CPU端將資料和GPU端需要執行的程式傳送給GPU，GPU計算結束後把資料傳給CPU。這樣說可能有點抽象，我們通過一個故事來說明。
  中秋將至，迪麗冷巴收到了一份超大的手工月餅訂單，要求在中秋到來的那天供應十萬個手工月餅；迪麗的設計師古力為她設計了B方案，先購買食材；然後把月餅食譜交給新西方學院的三十二名學徒；再接著把食材和月餅食譜交給學徒，由學徒們製作月餅；月餅成型後再派人去收。
  為什麼在大型運算的應用場景中，CUDA相對CPU多核計算是有優勢的呢？我們再回到迪麗的手工月餅生產方案中，再B方案設計之前，迪麗採用的月餅設計方案為A方案，請三個經驗老道的月餅師傅來製作手工月餅。設A方案中的老師傅每人每天製作的月餅量為$x$，B方案中的學徒每人每天製作的月餅數量為$y$；製作十萬個月餅，A方案需要$10000/(3x)$天，B方案需要$T+10000/(36y)$天，其中T為資料傳輸的時間；顯然，如若面對比較耗時的製作，$x$和$y$的值會偏小，傳輸時間T對整體的運算耗時就會相對較小，這時，B方案就更容易取得優勢。

學習CUDA的必備知識

  CUDA不僅是並行的技術，更是並行的藝術；想要掌握這門藝術，開發者需要付出諸多努力，但這些努力又都是值得的。

一定的C++知識

  因為CUDA的語法規則是類C++的，開發者為了避免使用者學習第二種變成語言，竭力讓CUDA的語法和C++相近；掌握C++是踏上CUDA程式設計之路的墊腳石。使用CUDA程式碼需要實現兩部分程式，一部分在CPU上執行，我們稱之為主機程式碼；另一部分在GPU上執行，我們成為裝置程式碼。

瞭解CUDA的架構

  瞭解CUDA的架構對程式設計而言是有幫助的，能夠幫助讀者設計更加穩定可靠的程式。許多CUDA的初學者會把一個功能都寫到一個函式中，對於CUDA而言，這種風險是十分嚴重的。這是因為CUDA給每個執行緒分配的共享記憶體都十分有限，當程式功能比較複雜導致每個執行緒所需的記憶體不能滿足要求時(通常，CUDA給每個裝置分配的共享記憶體只有1.5kb)，必然引發CUDA錯誤。這裡筆者鼓勵讀者採用資料流圖的方式進行CUDA程式設計。

練習CUDA

環境搭建

  CUDA的安裝可以參照筆者的這篇文章CUDA安裝指南。如果沒有什麼特殊情況，筆者建議使用最新版本的CUDA，因為，舊版本總是要被淘汰的，目前最新版的CUDA是CUDA11。筆者仍然堅守在CUDA10，這是因為筆者的工作是依賴於OpenCV4原始碼的，OpenCV4不支援比C++11更新的特性，這點導致其與CUDA11之間存在衝突。

建立工程

  在visual studio中建立CUDA工程，常見的做法有兩種，先介紹操作簡單的一種。
  檔案->新建->專案->NVIDIA中建立專案

  第二種，則是像寫依賴三方庫的程式碼那樣去包含標頭檔案和靜態庫。

注意事項

  總結了一下，CUDA程式設計會犯的這樣幾個錯誤。CPU端呼叫__global__函式使用了傳引用。使用傳引用時，程式拿到的是這個變數的地址，而不是這個變數的拷貝；當GPU端裝置拿到該變數在CPU端的地址，並把這個變數的地址當成GPU端地址訪問時，勢必是要出問題的；而這樣的問題是難以排查的。許多經驗豐富的C++程式設計師會使用模板進行程式設計，但模板程式設計帶來的壞處也是十分明顯的；會讓排查語法錯誤時顯得更為困難。
  想要規避以上錯誤，讀者需要注意以下幾點：第一，慎用傳引用；第二點慎用模板。

擴充閱讀

PyCUDA

  PyCUDA允許使用者從Python方問CUDA的平行計算API。使用Python瞭解一個庫是十分必要的，因為Python環境搭建簡單，同時能夠用更少的程式碼實現更多的功能。
  PyCUDA的官方安裝教程
  使用例程

import pycuda.driver as drv
import numpy

from pycuda.compiler import SourceModule
mod = SourceModule("""
__global__ void multiply_them(float *dest, float *a, float *b)
{
  const int i = threadIdx.x;
  dest[i] = a[i] * b[i];
}
""")

multiply_them = mod.get_function("multiply_them")

a = numpy.random.randn(400).astype(numpy.float32)
b = numpy.random.randn(400).astype(numpy.float32)

dest = numpy.zeros_like(a)
multiply_them(
        drv.Out(dest), drv.In(a), drv.In(b),
        block=(400,1,1), grid=(1,1))

print(dest-a*b)