CUDA减少JIT开销的两种方式

文章揭示了CUDA应用在启动时由于即时编译(JIT)导致的时间延迟问题,并介绍了两种缓解策略:Fat Binaries和JIT缓存。Fat Binaries在二进制文件中包含多架构PTX码,而JIT缓存则保存编译后的二进制代码以避免重复编译。在Docker环境中,首次运行CUDA应用慢是由于缺乏JIT缓存,解决方案包括映射宿主机缓存或打包缓存进镜像。
摘要由CSDN通过智能技术生成

之前测试的时候发现有时从Jenkins上拉下来的的megawise二进制包解压启动之后,执行第一条sql非常慢,需要1.5~3.5分钟不等。

后来docker化之后发现,docker中稳定必现第一条sql长时间卡住的问题。

当时查看日志发现总会卡在zdb_storage的一个MetaAgg的函数中,后来将这个函数注释掉,发现这个问题还是没有解决,会在engine中卡住。

百思不得其解之际,叶富哥一语惊醒梦中人:“好像都是卡在cuda调用的地方”。

从这个方向入手调查,果然找到了问题。

-------------------------------------揭开谜底的分割线------------------------------------------------

随着Nvidia gpu不断发展以支持新的特性,其指令集架构自然也发生了编号。由于应用程序需要在多代gpu上运行,Nvidia编译器工具链在同一个应用的可执行文件或者库中支持对多个不同架构的编译。

CUDA还依赖PTX(Parallel Thread Execution)虚拟GPU ISA(Instruction Set Architecture)来提供向前兼容性,以便已经部署的应用程序可以运行在未来的GPU架构之上。Nvidia编译器nvcc为了提供前向和

后向的兼容性,使用两段编译模型。第一个编译阶段将源代码编译成PTX虚拟汇编码,第二个阶段将PTX码编译成目标架构的二进制代码。CUDA驱动在运行时执行第二阶段编译。这种即时编译就会导致

应用程序在启动时(准确的说是在CUDA上下文创建的时候)产生一定的时间开销。CUDA使用两种方式来减少JIT(just-in-time)编译的时间开销:fat binaries和JIT缓存。

 

Fat Binaries

第一种完全避免

`@cuda.jit` 是一个装饰器,用于在 CUDA 设备上编写 GPU 加速的函数。下面是一个简单的例子: ```python from numba import cuda @cuda.jit def add_kernel(x, y, out): """ A CUDA kernel to add two arrays element-wise. """ i = cuda.grid(1) out[i] = x[i] + y[i] ``` 在这个例子中,我们定义了一个名为 `add_kernel` 的 CUDA kernel,它接受三个参数:两个输入数组 `x` 和 `y`,以及一个输出数组 `out`。函数的主体部分包含一个 `cuda.grid()` 调用,用于获取当前线程的 ID,然后将输入数组中对应位置的元素相加并将结果存储在输出数组中。 要在CUDA设备上运行这个函数,我们需要将输入数据和输出数组分配到设备的全局内存中,并使用 `cuda.jit()` 装饰器来编译内核。下面是一个完整的示例: ```python import numpy as np from numba import cuda # 生成输入数据 x = np.array([1, 2, 3]) y = np.array([4, 5, 6]) out = np.zeros_like(x) # 将数据移动到设备上 x_device = cuda.to_device(x) y_device = cuda.to_device(y) out_device = cuda.to_device(out) # 定义线程块大小和网格大小 threads_per_block = 32 blocks_per_grid = (x.size + (threads_per_block - 1)) // threads_per_block # 调用 CUDA 内核 add_kernel[blocks_per_grid, threads_per_block](x_device, y_device, out_device) # 将结果移回主机内存 out = out_device.copy_to_host() print(out) # 输出结果:[5 7 9] ``` 在这个示例中,我们首先生成了输入数据 `x` 和 `y`,以及一个全零的输出数组 `out`。然后,我们使用 `cuda.to_device()` 函数将这些数组移动到设备上。接下来,我们定义了线程块大小和网格大小,并调用 `add_kernel` 函数来执行加法操作。最后,我们使用 `copy_to_host()` 函数将结果数组移回到主机内存中并打印结果。 请注意,当使用 `@cuda.jit` 装饰器时,函数的参数类型必须是 Numba 支持的类型。如果您需要使用自定义类型,可以使用 Numba 的 `cuda.register_device_memory()` 函数将它们注册到设备内存中。
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值