# 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 內核對象。調用時，內核對象將專門為給定的參數（如果沒有合適的專用版本已經存在）＆amp;計算功能，并在與當前上下文關聯的設備上啟動。 內核對象不是由用戶構造的，而是使用 [`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`指令類似。