# 4.2。 CUDA 內核 API
> 原文: [http://numba.pydata.org/numba-doc/latest/cuda-reference/kernel.html](http://numba.pydata.org/numba-doc/latest/cuda-reference/kernel.html)
## 4.2.1。內核聲明
`@cuda.jit`裝飾器用于創建 CUDA 內核:
```py
numba.cuda.jit(func_or_sig=None, argtypes=None, device=False, inline=False, bind=True, link=[], debug=None, **kws)
```
JIT 編譯符合 CUDA Python 規范的 python 函數。如果提供了簽名,則返回一個函數進行編譯的函數。如果
| 參數: |
* **func_or_sig** (_function_ _or_ _numba.typing.Signature_) –
JIT 編譯的函數,或要編譯的函數的簽名。如果提供了函數,則返回`AutoJitCUDAKernel`。如果提供了簽名,則返回一個函數,該函數接受函數編譯并返回`AutoJitCUDAKernel`。
注意
內核不能有任何返回值。
* **設備**( [_bool_](https://docs.python.org/3/library/functions.html#bool "(in Python v3.7)")) - 表示這是否是設備功能。
* **結合**( [_bool_](https://docs.python.org/3/library/functions.html#bool "(in Python v3.7)")) - 立即強制結合 CUDA 環境
* **鏈接**( [_ 列表 _](https://docs.python.org/3/library/stdtypes.html#list "(in Python v3.7)") ) - 包含 PTX 源的文件列表,用于鏈接功能
* **debug** - 如果為 True,檢查執行內核時拋出的異常。由于這會降低性能,因此應僅用于調試目的。默認為 False。 (可以通過設置環境變量`NUMBA_CUDA_DEBUGINFO=1`來覆蓋默認值。)
* **fastmath** - 如果為 true,則啟用 flush-to-zero 和 fusion-multiply-add,禁用精確除法和平方根。此參數對設備功能沒有影響,其 fastmath 設置取決于調用它們的內核函數。
* **max_registers** - 限制內核每個線程最多使用這個數量的寄存器。有助于增加入住率。
|
| --- | --- |
```py
class numba.cuda.compiler.AutoJitCUDAKernel(func, bind, targetoptions)
```
CUDA 內核對象。調用時,內核對象將專門為給定的參數(如果沒有合適的專用版本已經存在)&計算功能,并在與當前上下文關聯的設備上啟動。
內核對象不是由用戶構造的,而是使用 [`numba.cuda.jit()`](#numba.cuda.jit "numba.cuda.jit") 裝飾器創建的。
```py
extensions
```
必須具有 <cite>prepare_args</cite> 函數的對象列表。當調用專用內核時,每個參數將傳遞給 <cite>prepare_args</cite> (從此列表中的最后一個對象到第一個對象)。 <cite>prepare_args</cite> 的參數是:
* <cite>ty</cite> numba 類型的參數
* <cite>val</cite> 參數值本身
* <cite>stream</cite> 用于當前調用內核的 CUDA 流
* <cite>retr</cite> 一個零 arg 函數列表,你可能想要將調用后的清理工作附加到。
<cite>prepare_args</cite> 函數必須返回一個元組<cite>(ty,val)</cite>,它將依次傳遞給下一個最右側<cite>擴展名</cite>。在調用所有擴展之后,生成的<cite>(ty,val)</cite>將被傳遞到 Numba 的默認參數編組邏輯中。
```py
inspect_asm(signature=None, compute_capability=None)
```
返回到目前為止遇到的所有簽名的生成的匯編代碼,或者返回 LLVM IR 以獲取特定簽名和 compute_capability(如果給定)。
```py
inspect_llvm(signature=None, compute_capability=None)
```
返回到目前為止遇到的所有簽名的 LLVM IR,或者給出特定簽名和 compute_capability 的 LLVM IR。
```py
inspect_types(file=None)
```
生成此函數的 Python 源代碼的轉儲,并使用相應的 Numba IR 和類型信息進行注釋。如果 _ 文件 _ 為 _ 無 _,轉儲將寫入 _ 文件 _ 或 _sys.stdout_ 。
```py
specialize(*args)
```
編譯并綁定當前上下文專用于給定 _args_ 的此內核版本。
各個專用內核是 [`numba.cuda.compiler.CUDAKernel`](#numba.cuda.compiler.CUDAKernel "numba.cuda.compiler.CUDAKernel") 的實例:
```py
class numba.cuda.compiler.CUDAKernel(llvm_module, name, pretty_name, argtypes, call_helper, link=(), debug=False, fastmath=False, type_annotation=None, extensions=[], max_registers=None)
```
CUDA 內核專門用于一組給定的參數類型。調用時,此對象將驗證參數類型是否與其專用的參數類型匹配,然后在設備上啟動內核。
```py
bind()
```
強制綁定到當前的 CUDA 上下文
```py
device
```
獲取當前活動上下文
```py
inspect_asm()
```
返回此內核的 PTX 代碼。
```py
inspect_llvm()
```
返回此內核的 LLVM IR。
```py
inspect_types(file=None)
```
生成此函數的 Python 源代碼的轉儲,并使用相應的 Numba IR 和類型信息進行注釋。如果 _ 文件 _ 為 _ 無 _,轉儲將寫入 _ 文件 _ 或 _sys.stdout_ 。
```py
ptx
```
該內核的 PTX 代碼。
## 4.2.2。內在屬性和函數
本節中的其余屬性和函數只能在 CUDA 內核中調用。
### 4.2.2.1。線程索引
```py
numba.cuda.threadIdx
```
當前線程塊中的線程索引,通過屬性`x`,`y`和`z`訪問。每個索引是一個整數,范圍從 0 到 0 到 [`numba.cuda.blockDim`](#numba.cuda.blockDim "numba.cuda.blockDim") 不包含的屬性的相應值。
```py
numba.cuda.blockIdx
```
線程塊網格中的塊索引,通過屬性`x`,`y`和`z`訪問。每個索引是一個整數,范圍從 0 到 0 到 [`numba.cuda.gridDim`](#numba.cuda.gridDim "numba.cuda.gridDim") 不包含的屬性的相應值。
```py
numba.cuda.blockDim
```
線程塊的形狀,在實例化內核時聲明。對于給定內核中的所有線程,該值是相同的,即使它們屬于不同的塊(即每個塊都是“滿”)。
```py
numba.cuda.gridDim
```
塊網格的形狀,通過屬性`x`,`y`和`z`訪問。
```py
numba.cuda.laneid
```
當前 warp 中的線程索引,作為一個整數,范圍從 0 到包含 [`numba.cuda.warpsize`](#numba.cuda.warpsize "numba.cuda.warpsize") 不包括。
```py
numba.cuda.warpsize
```
GPU 上的 warp 線程的大小。目前這總是 32。
```py
numba.cuda.grid(ndim)
```
返回整個塊網格中當前線程的絕對位置。 _ndim_ 應該對應于實例化內核時聲明的維數。如果 _ndim_ 為 1,則返回單個整數。如果 _ndim_ 為 2 或 3,則返回給定數量的整數的元組。
第一個整數的計算如下:
```py
cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x
```
并且與其他兩個索引類似,但使用`y`和`z`屬性。
```py
numba.cuda.gridsize(ndim)
```
返回整個塊網格的線程中的絕對大小(或形狀)。 _ndim_ 應該對應于實例化內核時聲明的維數。
第一個整數的計算如下:
```py
cuda.blockDim.x * cuda.gridDim.x
```
并且與其他兩個索引類似,但使用`y`和`z`屬性。
### 4.2.2.2。內存管理
```py
numba.cuda.shared.array(shape, dtype)
```
使用給定的`shape`和`dtype`在 CUDA 內核的本地內存空間中創建一個數組。
返回其內容未初始化的數組。
注意
同一線程塊中的所有線程都看到相同的數組。
```py
numba.cuda.local.array(shape, dtype)
```
使用給定的`shape`和`dtype`在 CUDA 內核的本地內存空間中創建一個數組。
返回其內容未初始化的數組。
注意
每個線程都看到一個唯一的數組
```py
numba.cuda.const.array_like(ary)
```
在編譯時將`ary`復制到 CUDA 內核上的常量內存空間。
返回類似`ary`參數的數組。
注意
所有線程和塊都看到相同的數組。
### 4.2.2.3。同步和原子操作
```py
numba.cuda.atomic.add(array, idx, value)
```
執行`array[idx] += value`。僅支持 int32,int64,float32 和 float64。 `idx`參數可以是整數或整數索引的元組,用于索引到多維數組。 `idx`中的元素數必須與`array`的維數相匹配。
在存儲新值之前返回`array[idx]`的值。表現得像原子載荷。
```py
numba.cuda.atomic.max(array, idx, value)
```
執行`array[idx] = max(array[idx], value)`。僅支持 int32,int64,float32 和 float64。 `idx`參數可以是整數或整數索引的元組,用于索引到多維數組。 `idx`中的元素數必須與`array`的維數相匹配。
在存儲新值之前返回`array[idx]`的值。表現得像原子載荷。
```py
numba.cuda.syncthreads()
```
同步同一線程塊中的所有線程。此函數實現與傳統多線程編程中的障礙相同的模式:此函數等待,直到塊中的所有線程調用它,此時它將控制權返回給所有調用者。
```py
numba.cuda.syncthreads_count(predicate)
```
[`numba.cuda.syncthreads`](#numba.cuda.syncthreads "numba.cuda.syncthreads") 的擴展,其中返回值是`predicate`為真的線程數。
```py
numba.cuda.syncthreads_and(predicate)
```
[`numba.cuda.syncthreads`](#numba.cuda.syncthreads "numba.cuda.syncthreads") 的擴展,如果`predicate`對所有線程都為真,則返回 1,否則返回 0。
```py
numba.cuda.syncthreads_or(predicate)
```
[`numba.cuda.syncthreads`](#numba.cuda.syncthreads "numba.cuda.syncthreads") 的擴展,如果任何線程的`predicate`為真,則返回 1,否則返回 0。
警告
所有 syncthreads 函數必須由線程塊中的每個線程調用。如果這樣做可能會導致未定義的行為。
### 4.2.2.4。記憶柵欄
內存屏障用于保證內存操作的效果可由同一線程塊內的其他線程,相同的 GPU 設備和相同的系統(跨全局內存的 GPU)看到。內存加載和存儲保證不會通過優化傳遞在內存柵欄中移動。
警告
內存柵欄被認為是高級 API,大多數用戶使用線程屏障(例如`syncthreads()`)。
```py
numba.cuda.threadfence()
```
設備級別的存儲器圍欄(在 GPU 內)。
```py
numba.cuda.threadfence_block()
```
線程塊級別的內存柵欄。
```py
numba.cuda.threadfence_system()
```
系統級別的內存柵欄(跨 GPU)。
### 4.2.2.5。 Warp Intrinsics
所有 warp 級操作至少需要 CUDA 9.參數`membermask`是一個 32 位整數掩碼,每個位對應于 warp 中的一個線程,1 表示該線程位于函數調用中的線程子集中。如果 GPU 計算能力低于 7.x,則`membermask`必須全為 1。
```py
numba.cuda.syncwarp(membermask)
```
在 warp 中同步屏蔽的線程子集。
```py
numba.cuda.all_sync(membermask, predicate)
```
如果`predicate`對于屏蔽 warp 中的所有線程都為 true,則返回非零值,否則返回 0。
```py
numba.cuda.any_sync(membermask, predicate)
```
如果`predicate`對于屏蔽 warp 中的任何線程為 true,則返回非零值,否則返回 0。
```py
numba.cuda.eq_sync(membermask, predicate)
```
如果 boolean `predicate`對于屏蔽 warp 中的所有線程都相同,則返回非零值,否則返回 0。
```py
numba.cuda.ballot_sync(membermask, predicate)
```
返回 warp 中`predicate`為 true 并且在給定掩碼內的所有線程的掩碼。
```py
numba.cuda.shfl_sync(membermask, value, src_lane)
```
在屏蔽的扭曲中隨機播放`value`并從`src_lane`返回`value`。如果這在 warp 之外,則返回給定的`value`。
```py
numba.cuda.shfl_up_sync(membermask, value, delta)
```
在屏蔽的扭曲中隨機播放`value`并從`laneid - delta`返回`value`。如果這在 warp 之外,則返回給定的`value`。
```py
numba.cuda.shfl_down_sync(membermask, value, delta)
```
在屏蔽的扭曲中隨機播放`value`并從`laneid + delta`返回`value`。如果這在 warp 之外,則返回給定的`value`。
```py
numba.cuda.shfl_xor_sync(membermask, value, lane_mask)
```
在屏蔽的扭曲中隨機播放`value`并從`laneid ^ lane_mask`返回`value`。
```py
numba.cuda.match_any_sync(membermask, value, lane_mask)
```
返回與掩碼 warp 中給定`value`具有相同`value`的線程掩碼。
```py
numba.cuda.match_all_sync(membermask, value, lane_mask)
```
返回(mask,pred)的元組,其中 mask 是掩碼 warp 中與給定`value`具有相同`value`的線程的掩碼,如果它們都具有相同的值,則為 0。是掩碼 warp 中的所有線程是否具有相同 warp 的布爾值。
### 4.2.2.6。整數內在函數
可以使用 CUDA Math API 的整數內在函數的子集。有關進一步的文檔,包括語義,請參閱 [CUDA 工具包文檔](https://docs.nvidia.com/cuda/cuda-math-api/group__CUDA__MATH__INTRINSIC__INT.html)。
```py
numba.cuda.popc()
```
返回給定值中的設置位數。
```py
numba.cuda.brev()
```
反轉整數值的位模式,例如 0b10110110 變為 0b01101101。
```py
numba.cuda.clz()
```
計算值中前導零的數量。
```py
numba.cuda.ffs()
```
在整數中查找設置為 1 的最低有效位的位置。
### 4.2.2.7。浮點內在函數
可以使用 CUDA Math API 的浮點內部函數的子集。有關進一步的文檔,包括語義,請參閱 CUDA Toolkit 文檔的[單](https://docs.nvidia.com/cuda/cuda-math-api/group__CUDA__MATH__SINGLE.html)和[雙](https://docs.nvidia.com/cuda/cuda-math-api/group__CUDA__MATH__DOUBLE.html)精度部分。
```py
numba.cuda.fma()
```
執行融合乘法 - 加法運算。以 C api 中的`fma`和`fmaf`命名,但映射到`fma.rn.f32`和`fma.rn.f64`(舍入到最近 - 偶數)PTX 指令。
### 4.2.2.8。控制流程說明
CUDA 控制流指令的子集可直接作為內在函數使用。避免分支是提高 CUDA 性能的關鍵方法,使用這些內在函數意味著您不必依賴`nvcc`優化器來識別和刪除分支。有關進一步的文檔,包括語義,請參閱[相關的 CUDA 工具包文檔](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#comparison-and-selection-instructions)。
```py
numba.cuda.selp()
```
根據第一個參數的值,在兩個表達式之間進行選擇。與 LLVM 的`select`指令類似。
- 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. 術語表