# 3.10。示例
> 原文: [http://numba.pydata.org/numba-doc/latest/cuda/examples.html](http://numba.pydata.org/numba-doc/latest/cuda/examples.html)
## 3.10.1。矩陣乘法
這是使用 CUDA 內核的矩陣乘法的簡單實現:
```py
@cuda.jit
def matmul(A, B, C):
"""Perform square matrix multiplication of C = A * B
"""
i, j = cuda.grid(2)
if i < C.shape[0] and j < C.shape[1]:
tmp = 0.
for k in range(A.shape[1]):
tmp += A[i, k] * B[k, j]
C[i, j] = tmp
```
這種實現很簡單直觀但性能很差,因為相同的矩陣元素將從設備內存中多次加載,這很慢(某些設備可能有透明的數據緩存,但它們可能不夠大,不能一次保存整個輸入)。
如果我們使用阻塞算法來減少對設備內存的訪問,則會更快。 CUDA 為塊中的線程提供快速[共享內存](memory.html#cuda-shared-memory),以便在任務上協同計算。以下實現了使用共享內存的方形矩陣乘法的更快版本:
```py
from numba import cuda, float32
# Controls threads per block and shared memory usage.
# The computation will be done on blocks of TPBxTPB elements.
TPB = 16
@cuda.jit
def fast_matmul(A, B, C):
# Define an array in the shared memory
# The size and type of the arrays must be known at compile time
sA = cuda.shared.array(shape=(TPB, TPB), dtype=float32)
sB = cuda.shared.array(shape=(TPB, TPB), dtype=float32)
x, y = cuda.grid(2)
tx = cuda.threadIdx.x
ty = cuda.threadIdx.y
bpg = cuda.gridDim.x # blocks per grid
if x >= C.shape[0] and y >= C.shape[1]:
# Quit if (x, y) is outside of valid C boundary
return
# Each thread computes one element in the result matrix.
# The dot product is chunked into dot products of TPB-long vectors.
tmp = 0.
for i in range(bpg):
# Preload data into shared memory
sA[tx, ty] = A[x, ty + i * TPB]
sB[tx, ty] = B[tx + i * TPB, y]
# Wait until all threads finish preloading
cuda.syncthreads()
# Computes partial product on the shared memory
for j in range(TPB):
tmp += sA[tx, j] * sB[j, ty]
# Wait until all threads finish computing
cuda.syncthreads()
C[x, y] = tmp
```
由于共享內存是有限的資源,因此代碼一次從輸入數組預加載小塊。然后,它調用 [`syncthreads()`](../cuda-reference/kernel.html#numba.cuda.syncthreads "numba.cuda.syncthreads") 等待所有線程完成預加載并在共享內存上進行計算之前。它在計算后再次同步,以確保所有線程在共享內存中完成數據,然后在下一次循環迭代中覆蓋它。
- 1. 用戶手冊
- 1.1。 Numba 的約 5 分鐘指南
- 1.2。概述
- 1.3。安裝
- 1.4。使用@jit 編譯 Python 代碼
- 1.5。使用@generated_jit 進行靈活的專業化
- 1.6。創建 Numpy 通用函數
- 1.7。用@jitclass 編譯 python 類
- 1.8。使用@cfunc 創建 C 回調
- 1.9。提前編譯代碼
- 1.10。使用@jit 自動并行化
- 1.11。使用@stencil裝飾器
- 1.12。從 JIT 代碼 中回調到 Python 解釋器
- 1.13。性能提示
- 1.14。線程層
- 1.15。故障排除和提示
- 1.16。常見問題
- 1.17。示例
- 1.18。會談和教程
- 2. 參考手冊
- 2.1。類型和簽名
- 2.2。即時編譯
- 2.3。提前編譯
- 2.4。公用事業
- 2.5。環境變量
- 2.6。支持的 Python 功能
- 2.7。支持的 NumPy 功能
- 2.8。與 Python 語義的偏差
- 2.9。浮點陷阱
- 2.10。 Python 2.7 壽命終止計劃
- 3. 用于 CUDA GPU 的 Numba
- 3.1。概述
- 3.2。編寫 CUDA 內核
- 3.3。內存管理
- 3.4。編寫設備功能
- 3.5。 CUDA Python 中支持的 Python 功能
- 3.6。支持的原子操作
- 3.7。隨機數生成
- 3.8。設備管理
- 3.10。示例
- 3.11。使用 CUDA 模擬器 調試 CUDA Python
- 3.12。 GPU 減少
- 3.13。 CUDA Ufuncs 和廣義 Ufuncs
- 3.14。共享 CUDA 內存
- 3.15。 CUDA 陣列接口
- 3.16。 CUDA 常見問題
- 4. CUDA Python 參考
- 4.1。 CUDA 主機 API
- 4.2。 CUDA 內核 API
- 4.3。內存管理
- 5. 用于 AMD ROC GPU 的 Numba
- 5.1。概述
- 5.2。編寫 HSA 內核
- 5.3。內存管理
- 5.4。編寫設備功能
- 5.5。支持的原子操作
- 5.6。代理商
- 5.7。 ROC Ufuncs 和廣義 Ufuncs
- 5.8。示例
- 6. 擴展 Numba
- 6.1。高級擴展 API
- 6.2。低級擴展 API
- 6.3。示例:間隔類型
- 7. 開發者手冊
- 7.1。貢獻給 Numba
- 7.2。 Numba 建筑
- 7.3。多態調度
- 7.4。關于發電機的注意事項
- 7.5。關于 Numba Runtime 的注意事項
- 7.6。使用 Numba Rewrite Pass 獲得樂趣和優化
- 7.7。實時變量分析
- 7.8。上市
- 7.9。模板注釋
- 7.10。關于自定義管道的注意事項
- 7.11。環境對象
- 7.12。哈希 的注意事項
- 7.13。 Numba 項目路線圖
- 8. Numba 增強建議
- 9. 術語表