Numba CUDA Python 参考指南：从 Host API 到 Kernel 内建函数与 Libdevice 完整 API 索引

Numba CUDA Python 参考指南从 Host API 到 Kernel 内建函数与 Libdevice 完整 API 索引【免费下载链接】numbaNumPy aware dynamic Python compiler using LLVM项目地址: https://gitcode.com/gh_mirrors/nu/numbaNumba 的docs/source/cuda-reference/目录下存放着一套完整的CUDA Python ReferenceCUDA Python 参考由index.rst索引页组织为五个相互独立又彼此衔接的参考文档Host API主机端 API、Kernel API内核端 API、CUDA 专用类型、内存管理与 Libdevice 函数。本文以这份参考文档为骨架逐条展开其中的 API 语义、launch 配置语法与使用前提并结合numba/cuda/下的源码实现与测试进行佐证帮助读者把这套参考转化为可以直接编写、配置和调优 CUDA Python 程序的实战知识。参考文档的组织结构docs/source/cuda-reference/index.rst是一个典型的 Sphinxtoctree索引页其核心作用是把整套 CUDA Python 参考组织成五个章节子文档主题覆盖范围host.rstCUDA Host API设备检测、上下文/设备管理、编译入口、Profile 与事件、流管理、运行时查询kernel.rstCUDA Kernel APIcuda.jit声明、Dispatcher 配置与检查、线程索引、内存、原子操作、同步、Warp 内建、FP16 内建等types.rstCUDA 专用类型向量类型float32x3等及其构造规则memory.rst内存管理设备数组分配、数据传输、固定/映射/托管内存libdevice.rstLibdevice 函数被包装的 CUDA libdevice 数学函数全集注意docs/source/cuda-reference/下的所有文档均带有.. cuda-deprecated::标记意为该模块CUDA 目标在 Numba 中被标记为已弃用但这不影响本文对 API 语义本身的讲解。此外types.rst开篇强调该页只介绍CUDA 目标特有的类型其他大量可用类型如标量、数组、元组等内建类型请参考 CUDA 目标的通用类型说明。CUDA Host API主机端编程接口Host API 是运行在 CPU 侧的接口负责设备的发现、上下文管理、内核编译与 launch 之外的调度操作。对应参考文档为 host.rst。设备检测与查询以下函数用于查询当前可用的硬件numba.cuda.is_available()返回布尔值指示 CUDA 驱动与设备是否可用。这是编写无 GPU 也能优雅降级代码的首选入口。numba.cuda.detect()枚举系统内可用的 CUDA 设备。其实现位于 api.py返回一个包含设备信息的列表常用于诊断脚本中输出当前机器上的 GPU 概况。from numba import cuda if cuda.is_available(): print(cuda.detect()) else: print(CUDA 不可用将回退到 CPU 路径)上下文管理CUDA Python 的函数都在CUDA context上下文内执行系统中的每个 CUDA 设备都关联一个 CUDA context而 Numba 目前每个线程只允许一个 context。CUDA Context 的底层语义对应 CUDA Driver API 的 Context Management 与 CUDA C Programming Guide 中关于 Context 的说明在 Numba 中context 是numba.cuda.cudadrv.driver.Context类的实例该类暴露了以下成员reset销毁 context 及其关联资源get_memory_info查询 context 的内存使用情况push/pop将 context 压入或弹出当前线程的 context 栈。获取或选择 context 的函数numba.cuda.current_context()返回当前线程的 context必要时自动创建。numba.cuda.require_context()要求调用必须发生在已有 context 内例如在 callback 中确保 context 已被激活。影响当前 context 的函数numba.cuda.synchronize()阻塞主机线程直到所有排队的工作内核、传输等完成。源码位于 api.py。numba.cuda.close()关闭当前 context 并释放资源。源码位于 api.py。设备管理Numba 维护一个受支持的 CUDA 设备列表numba.cuda.gpus一个可索引的设备列表按整数设备 ID 索引。通过cuda.gpus取到的设备始终是numba.cuda.cudadrv.devices._DeviceContextManager实例——它是一个上下文管理器配合with语句可以把该设备设为当前设备from numba import cuda with cuda.gpus[1]: # 此代码块内使用 GPU 1 ...numba.cuda.gpus.current当前被选中的设备。另外还有三个用于选择/获取设备的函数numba.cuda.select_device(dev)选择指定设备并激活其 context。numba.cuda.get_current_device()返回当前设备。numba.cuda.list_devices()列出所有可用设备。numba.cuda.cudadrv.driver.Device类用于查询所选设备的功能特性其公开属性包括compute_capability(major, minor)元组表示设备支持的计算能力如(7, 5)id设备的整数 IDname设备名称如GeForce GTX 970uuid设备 UUID形如GPU-e6489c45-5b68-3b03-bab7-0e7c8e809643reset删除该设备的 context会销毁该 context 内创建的所有内存分配、事件与流supports_float16返回布尔值指示设备是否支持 float16 运算配合下文 Kernel API 中的cuda.is_float16_supported使用。编译入口Numba 提供了一个不经过 Driver API 的纯编译入口适用于以下场景生成用于嵌入其他 PTX 代码的 PTX例如供 Numba/Python 生态之外使用生成 PTX 或 LTO-IR 与非 Python 翻译单元链接在没有设备的环境下生成代码在 fork 之前、未初始化 CUDA 的情况下生成代码。对应的函数为numba.cuda.compile。需要注意使用编译到 PTX / LTO-IR 的方式时ABI 问题由用户自行负责传入abic关键字参数可以解决大多数此类问题参见 CUDA 目标关于 C ABI 的说明。环境变量NUMBA_CUDA_DEFAULT_PTX_CC用于控制compile默认面向的计算能力见 CUDA 环境变量章节。如果需要生成当前设备计算能力对应的代码则使用numba.cuda.compile_for_current_device()对于只编译到 PTX的旧式代码Numba 还提供两个兼容函数numba.cuda.compile_ptx()numba.cuda.compile_ptx_for_current_device()性能测量Profiling 与事件Profiling。NVIDIA Visual Profiler 可以直接作用于正在执行的 CUDA Python 代码用户并不需要在自己的代码中插入调用但以下函数允许选择性地对代码的特定部分进行 profilingnumba.cuda.profile_start()开始收集 profiling 数据numba.cuda.profile_stop()停止收集numba.cuda.profiling()作为上下文管理器使用with cuda.profiling(): ...包裹的代码段会被纳入 profiling。事件Events。事件用于监控执行进度并记录特定时间点的时间戳。事件创建后立即返回之后可以查询该事件是否已被执行引擎到达。相关函数numba.cuda.event(timingTrue)创建事件对象numba.cuda.event_elapsed_time(start, end)计算两个事件之间的毫秒耗时。事件是numba.cuda.cudadrv.driver.Event类的实例其成员包括query查询事件是否已到达record在流中记录事件时间戳synchronize等待事件完成wait让后续工作等待该事件。start cuda.event() start.record() # ... 执行内核 ... end cuda.event() end.record() end.synchronize() print(耗时(ms):, cuda.event_elapsed_time(start, end))流管理流Stream允许单个设备在同一 context 内并发执行同一流中的工作按序执行不同流中的工作可以并发。大多数涉及 CUDA 设备的操作都可以异步执行包括数据传输与内核启动。Numba 默认使用legacy default stream传统默认流作为默认流。若希望把per-thread default stream每线程默认流设为默认可以设置环境变量NUMBA_CUDA_PER_THREAD_DEFAULT_STREAM1详见 CUDA 环境变量章节。无论该设置如何都可以用下面的函数显式构造这两种默认流对象。流是numba.cuda.cudadrv.driver.Stream类的实例其成员包括synchronize等待流中所有工作完成auto_synchronize作为上下文管理器时自动同步add_callback在流中注册回调async_done异步完成通知。创建与获取流的函数numba.cuda.stream()创建新流源码位于 api.pynumba.cuda.default_stream()获取默认流numba.cuda.legacy_default_stream()显式获取 legacy 默认流numba.cuda.per_thread_default_stream()显式获取 per-thread 默认流numba.cuda.external_stream(ptr)用一个外部已分配的流构造 NumbaStream对象。注意外部流的生命周期由用户负责——Numba 不会释放它且在 NumbaStream对象使用期间该流必须保持有效。运行时版本查询Numba 主要使用 Driver API同时也提供了对 Runtime API 的简单封装用于查询运行时版本。通过cuda.runtime访问它是numba.cuda.cudadrv.runtime.Runtime类的实例成员包括get_version获取当前 CUDA Runtime 版本is_supported_version查询当前版本是否被当前 Numba 官方支持supported_versions列出 Numba 支持的版本集合。此外顶层函数numba.cuda.is_supported_version()也可以直接查询当前运行时是否与当前 Numba 版本匹配。CUDA Kernel API内核端编程接口Kernel API 涵盖内核声明、Dispatcher 配置与检查、以及只能在内核内部调用的内建属性与函数。对应参考文档为 kernel.rst。内核声明与 Dispatcher 配置用cuda.jit装饰器可以把一个 Python 函数编译为 CUDA kernel得到一个可以配置并 launch 的 CUDA dispatcher 对象即numba.cuda.dispatcher.CUDADispatcher。launch 配置使用下标语法# func 是某个被 cuda.jit 装饰的函数 func[griddim, blockdim, stream, sharedmem]griddim、blockdim网格与线程块的尺寸可以是整数也可以是长度至多 3 的元组stream可选参数指定内核在其上 launch 的流sharedmem动态共享内存的字节数。下标返回一个配置对象可以再传入内核参数调用configured func[griddim, blockdim, stream, sharedmem] configured(x, y, z)更地道的写法是配置与调用合并在一条语句中funcgriddim, blockdim, stream, sharedmem这与 CUDA C/C 的 launch 配置类似funcgriddim, blockdim, sharedmem, stream(x, y, z);注意Numba 中stream与sharedmem的顺序与 CUDA C/C 相反——Numba 是先 stream 后 sharedmem而 CUDA C/C 是先 sharedmem 后 stream。Dispatcher 对象还提供一批检查与特化方法见CUDADispatcher的成员inspect_asm/inspect_llvm/inspect_sass查看编译产物汇编、LLVM IR、SASSinspect_types查看类型推断结果get_regs_per_thread查询每线程寄存器数specialize/specialized特化与特化状态检查extensions扩展注册表forall为并行 launch 生成特殊配置get_shared_mem_per_block/get_max_threads_per_block/get_const_mem_size/get_local_mem_per_thread查询内核的资源占用。线程索引内建这些属性与函数只能从 CUDA kernel 内部调用numba.cuda.threadIdx当前线程在其线程块内的索引通过x、y、z属性访问范围为[0, blockDim)numba.cuda.blockIdx当前线程块在网格中的索引通过x、y、z访问范围为[0, gridDim)numba.cuda.blockDim线程块的形状launch 时声明对同一内核的所有线程都相同即使它们属于不同块即每个块都是满的numba.cuda.gridDim网格的形状通过x、y、z访问numba.cuda.laneid当前线程在其 warp 内的索引范围为[0, warpsize)numba.cuda.warpsizewarp 的线程数目前恒为 32。两个核心辅助函数numba.cuda.grid(ndim)返回当前线程在整个网格中的绝对位置。ndim应与 launch 时声明的维度一致ndim1时返回单个整数ndim2/3时返回对应长度的元组。第一个整数的计算方式为cuda.threadIdx.x cuda.blockIdx.x * cuda.blockDim.x其余两个维度同理使用y、z属性。其底层实现位于 intrinsics.py。numba.cuda.gridsize(ndim)返回整个网格的绝对尺寸形状。第一个整数的计算方式为cuda.blockDim.x * cuda.gridDim.x实现位于 intrinsics.py。from numba import cuda cuda.jit def add_kernel(a, b, out): i cuda.grid(1) # 全局线程 ID if i out.shape[0]: out[i] a[i] b[i]内核内存管理内建numba.cuda.shared.array(shape, dtype)在内核的共享内存空间创建数组。返回内容未初始化的数组。同一线程块内的所有线程看到的都是同一个数组。numba.cuda.local.array(shape, dtype)在内核的局部内存空间创建数组。返回内容未初始化。每个线程看到的是自己独有的数组。numba.cuda.const.array_like(ary)在编译时把ary复制到常量内存空间返回一个与原参数类似的数组。所有线程与所有块看到的都是同一个数组。cuda.jit def shared_example(x): tid cuda.threadIdx.x s cuda.shared.array(32, dtypenumba.float32) s[tid] x[tid] cuda.syncthreads() ...同步与原子操作原子操作numba.cuda.atomic提供了一系列原子操作统一的行为模式是返回写入前array[idx]的旧值且都表现为一次原子读atomic load。idx可以是整数也可以是用于多维数组索引的整数元组其元素个数必须与数组维度匹配。各操作的支持类型如下函数语义支持类型atomic.add(array, idx, value)array[idx] valueint32, int64, float32, float64atomic.sub(array, idx, value)array[idx] - valueint32, int64, float32, float64atomic.and_(array, idx, value)array[idx] valueint32, uint32, int64, uint64atomic.or_(array, idx, value)array[idx] \| valueint32, uint32, int64, uint64atomic.xor(array, idx, value)array[idx] ^ valueint32, uint32, int64, uint64atomic.exch(array, idx, value)array[idx] valueint32, uint32, int64, uint64atomic.inc(array, idx, value)array[idx] (0 if array[idx] value else array[idx] 1)uint32, uint64atomic.dec(array, idx, value)array[idx] (value if (array[idx] 0) or (array[idx] value) else array[idx] - 1)uint32, uint64atomic.max(array, idx, value)array[idx] max(array[idx], value)int32, int64, float32, float64atomic.cas(array, idx, old, value)if array[idx] old: array[idx] valueint32, int64, uint32, uint64atomic.cas是唯一的比较并交换操作行为类似一次原子 compare-and-swap。在 numba/cuda/tests/cudapy/test_atomics.py 中可以看到针对这些原子操作的类型覆盖与正确性测试。cuda.jit def histogram(x, bins): i cuda.grid(1) if i x.shape[0]: b int(x[i] * 10) # 简化分箱 cuda.atomic.add(bins, b, 1) # 线程安全地累加线程同步numba.cuda.syncthreads()同步同一线程块内的所有线程等价于传统多线程编程中的屏障barrier——等待块内所有线程都到达该调用点后才向所有调用者返回。底层实现在 intrinsics.py。numba.cuda.syncthreads_count(predicate)syncthreads的扩展返回谓词为真的线程数量numba.cuda.syncthreads_and(predicate)若谓词对所有线程都为真则返回 1否则 0numba.cuda.syncthreads_or(predicate)若谓词对任一线程为真则返回 1否则 0。警告所有syncthreads变体必须被线程块内的每一个线程调用否则可能导致未定义行为。协作组Cooperative Groupsnumba.cuda.cg.this_grid()获取当前网格组GridGroup。numba.cuda.cg.GridGroup网格组类。用户不应直接构造GridGroup而应通过cg.this_grid()获取。其成员方法sync()用于同步当前网格组跨 block 的全网格同步。内存栅栏Memory Fences内存栅栏用于保证内存操作的效果对其他线程可见作用范围可以是同一线程块、同一 GPU 设备、或同一系统跨 GPU 的全局内存。优化阶段保证加载与存储不会越过内存栅栏移动。numba.cuda.threadfence()设备级内存栅栏GPU 内numba.cuda.threadfence_block()线程块级内存栅栏numba.cuda.threadfence_system()系统级内存栅栏跨 GPU。警告内存栅栏属于高级 API大多数使用场景应该使用线程屏障如syncthreads()。Warp 内建函数Warp 内建的membermask是一个 32 位整数掩码每一位对应 warp 中的一个线程1 表示该线程属于本次调用所针对的线程子集。如果 GPU 计算能力低于 7.xmembermask必须全为 1。numba.cuda.syncwarp(membermask)同步 warp 中被掩码选中的线程子集numba.cuda.all_sync(membermask, predicate)掩码内所有线程谓词为真时返回非零值否则返回 0numba.cuda.any_sync(membermask, predicate)掩码内任一线程谓词为真时返回非零值否则返回 0numba.cuda.eq_sync(membermask, predicate)掩码内所有线程的布尔谓词相同时返回非零值否则 0numba.cuda.ballot_sync(membermask, predicate)返回掩码内所有谓词为真线程的位掩码numba.cuda.shfl_sync(membermask, value, src_lane)跨 warp 洗牌返回src_lane的value若src_lane在 warp 外则返回给定的valuenumba.cuda.shfl_up_sync(membermask, value, delta)返回来自laneid - delta的value超出 warp 则返回自身值numba.cuda.shfl_down_sync(membermask, value, delta)返回来自laneid delta的value超出 warp 则返回自身值numba.cuda.shfl_xor_sync(membermask, value, lane_mask)返回来自laneid ^ lane_mask的valuenumba.cuda.match_any_sync(membermask, value, lane_mask)返回掩码内与给定value相同值的线程掩码numba.cuda.match_all_sync(membermask, value, lane_mask)返回(mask, pred)元组——若掩码内所有线程值相同mask为这些线程的掩码否则为 0pred表示掩码内所有线程是否同值numba.cuda.activemask()返回当前 warp 中所有活跃线程的 32 位掩码第 N 位在调用时第 N 个 lane 活跃则置 1已退出内核的线程一律标记为不活跃numba.cuda.lanemask_lt()返回所有 ID 小于当前 lane 的 lane包括不活跃的掩码。cuda.jit def warp_reduce(x, out): lane cuda.laneid v x[cuda.grid(1)] # 使用全 1 掩码的 warp 归约 v v cuda.shfl_down_sync(0xffffffff, v, 16) v v cuda.shfl_down_sync(0xffffffff, v, 8) v v cuda.shfl_down_sync(0xffffffff, v, 4) v v cuda.shfl_down_sync(0xffffffff, v, 2) v v cuda.shfl_down_sync(0xffffffff, v, 1) if lane 0: out[cuda.blockIdx.x] v整数内建CUDA Math API 的整数内建的一个子集可直接使用numba.cuda.popc(x)返回x中置 1 的位数numba.cuda.brev(x)返回x位模式的逆序如0b10110110→0b01101101numba.cuda.clz(x)返回x前导零的个数numba.cuda.ffs(x)返回x中最低位置 1 的位的位置最低位位置记为 1ffs(0)返回 0。浮点内建numba.cuda.fma(a, b, c)融合乘加操作命名源自 C API 的fma/fmaf映射到 PTX 指令fma.rn.f32与fma.rn.f64round-to-nearest-even 舍入模式numba.cuda.cbrt(x)立方根x ** (1/3)命名源自cbrt/cbrtf仅支持 float32 与 float64。16 位浮点内建cuda.fp16cuda.fp16模块中的函数用于操作 16 位浮点操作数返回 16 位浮点结果。使用前需要确认支持情况numba.cuda.is_float16_supported()返回当前配置下 Numba 是否支持编译使用float16的代码设备级支持通过Device.supports_float16属性查询。cuda.fp16提供的运算多数标注了对应的 PTX 指令类别函数语义 / 映射的 PTX 指令算术hfma(a, b, c)(a * b) cfma.rn.f16算术hadd(a, b)a badd.f16算术hsub(a, b)a - bsub.f16算术hmul(a, b)a * bmul.f16算术hdiv(a, b)a / b一元hneg(a)-aneg.f16一元habs(a)|a|三角函数hsin(a)、hcos(a)正弦、余弦对数hlog(a)、hlog10(a)、hlog2(a)自然对数、以 10 为底、以 2 为底的对数指数hexp(a)、hexp10(a)、hexp2(a)自然指数、以 10 为底、以 2 为底的指数舍入hfloor(a)、hceil(a)向下/向上取整开方hsqrt(a)、hrsqrt(a)、hrcp(a)平方根、平方根倒数、倒数舍入hrint(a)、htrunc(a)四舍五入到最近整数、向零截断比较heq、hne、hgt、hge、hlt、hle、!、、、、返回布尔值极值hmax(a, b)a if a b else b极值hmin(a, b)a if a b else b控制流指令numba.cuda.selp(a, b, c)根据第一个参数的值在两个表达式之间选择类似 LLVM 的select指令。避免分支是提升 CUDA 性能的关键手段使用这类内建意味着不必依赖nvcc优化器去识别和消除分支。定时器内建numba.cuda.nanosleep(ns)让线程挂起大约ns纳秒。常用于需要精细控制时序如 spin-wait 调优的场景。CUDA 专用类型向量类型对应参考文档为 types.rst介绍 CUDA 目标特有的向量类型。命名规则与 CUDA C/C 的向量类型相比有两个重要区别推荐命名为base_typexN其中base_type是向量基类型N是元素个数例如int64x3、uint16x4、float32x4。新写的 Numba CUDA kernel 推荐使用这种命名。为方便从 CUDA C/C 迁移也提供了与 C/C 命名一致的别名例如float3别名float32x3long3别名int32x3或int64x3取决于平台。构造方式与 CUDA C/C 使用工厂函数不同Numba 的向量类型直接用构造函数构造from numba.cuda import float32x3 # 在内核中 f3 float32x3(0.0, -1.0, 1.0)向量类型可以由向量与基本类型混合构造只要总分量数匹配目标向量类型即可。以下构造全部合法zero uint32(0) u2 uint32x2(1, 2) # 用基本类型 2 分量向量构造 3 分量向量 u3 uint32x3(zero, u2) # 用两个 2 分量向量构造 4 分量向量 u4 uint32x4(u2, u2)分量访问向量类型的第 1、2、3、4 个分量分别通过字段x、y、z、w访问。在 Numba 当前版本中构造后分量不可变可变分量支持预计在后续版本中加入v1 float32x2(1.0, 1.0) v2 float32x2(1.0, -1.0) dotprod v1.x * v2.x v1.y * v2.y内存管理 API对应参考文档为 memory.rst负责主机与设备之间的数据搬运及各类特殊内存的分配。分配与传输函数numba.cuda.to_device(obj, stream0, copyTrue, toNone)把主机对象拷贝到设备返回设备数组。源码位于 api.py。numba.cuda.device_array(shape, dtype, stridesNone, orderC, stream0)在设备上分配空数组不初始化内容。源码位于 api.py。numba.cuda.device_array_like(ary)按ary的形状与 dtype 分配设备数组。numba.cuda.pinned_array(...)/numba.cuda.pinned_array_like(ary)分配**页锁定pinned**主机内存数组可显著提升与设备之间的拷贝带宽。numba.cuda.mapped_array(...)/numba.cuda.mapped_array_like(ary)分配映射内存数组主机与设备共享同一物理内存支持零拷贝访问。numba.cuda.managed_array(...)分配**托管managed**内存数组由统一内存Unified Memory系统负责页面迁移。numba.cuda.pinned(ary)把现有数组固定pin为页锁定内存上下文管理器用法。numba.cuda.mapped(ary)把现有数组映射到设备地址空间上下文管理器用法。from numba import cuda import numpy as np h_a np.arange(1024, dtypenp.float32) d_a cuda.to_device(h_a) # 主机 → 设备 h_b d_a.copy_to_host() # 设备 → 主机 d_b cuda.device_array_like(h_a) # 按形状分配设备对象设备数组是numba.cuda.cudadrv.devicearray.DeviceNDArray的实例主要方法包括copy_to_device(ary)从主机或其他设备拷贝到该设备数组copy_to_host(aryNone)拷贝回主机is_c_contiguous()/is_f_contiguous()检查 C/F 连续内存布局ravel()/reshape(shape)/split(...)展平、重塑与拆分。DeviceRecord结构化记录类型数组提供copy_to_device与copy_to_hostMappedNDArray映射内存数组提供copy_to_device、copy_to_host与split。Libdevice 函数对应参考文档为 libdevice.rst。Numba 包装了 CUDA 的libdevice数学库除__nv_nan与__nv_nanf外所有 libdevice 函数都被包装并可通过numba.cuda.libdevice模块访问该模块的__init__.py定义了全部包装函数libdevice.py、libdevicedecl.py、libdeviceimpl.py、libdevicefuncs.py分别负责声明、类型与实现映射。两个例外函数返回 quiet NaN 的表示但它们的参数一个指向表示说明对象的指针没有文档且形式特殊——它不是像其他指针参数那样的输出参数。如果确实需要 NaN可以用其他方式获得例如 Python 的math.nan。from numba import cuda from numba.cuda import libdevice cuda.jit def k(x): i cuda.grid(1) x[i] libdevice.fast_sinf(x[i]) # 示例使用 libdevice 三角函数完整函数清单由 libdevice.rst 中的automodule指令自动展开numba.cuda.libdevice的所有成员覆盖了__nv_*前缀下的浮点、整型、三角函数、指数对数、舍入等大量数学原语。小结docs/source/cuda-reference/这套参考为 CUDA Python 编程提供了自下而上的完整 API 索引主机端通过 host.rst 掌握设备/上下文/流/事件/编译入口内核端通过 kernel.rst 掌握 launch 配置与全部内建线程索引、共享/局部/常量内存、原子操作、同步、Warp 洗牌、FP16 运算再配合 types.rst 的向量类型、memory.rst 的各类内存分配与传输以及 libdevice.rst 的数学函数全集即可覆盖从探测设备到编写并调优 kernel的完整开发闭环。实际编写代码时可对照 api.py主机端函数实现与 intrinsics.py内核端内建实现核实函数签名并在 numba/cuda/tests/ 下找到对应的行为测试作为使用范例。【免费下载链接】numbaNumPy aware dynamic Python compiler using LLVM项目地址: https://gitcode.com/gh_mirrors/nu/numba创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考