python中一个简单的cuda/opencl内核调谐器
kernel_tuner的Python项目详细描述
内核调谐器通过启用基于python的gpu代码单元测试并使开发用于自动调整gpu内核的脚本变得容易,从而简化了优化和自动调整gpu程序的软件开发。这也意味着在内核代码中不需要大量的更改和新的依赖项。内核仍然可以从任何宿主编程语言编译和正常使用。
内核调谐器为gpu程序的自动调整提供了一个全面的解决方案,支持在主机和设备代码中自动调整用户定义的参数,支持在调整期间对所有基准内核的输出验证,以及许多加快调整过程的优化策略。
文档
有完整的文档 here。
安装
安装内核调谐器的最简单方法是使用pip:
调整CUDA内核:
- 首先,确保安装了CUDA Toolkit
- 然后键入:pip install kernel_tuner[cuda]
要优化opencl内核:
- 首先,确保您的opencl平台有一个opencl编译器
- 然后键入:pip install kernel_tuner[opencl]
或两者皆有:
- pip install kernel_tuner[cuda,opencl]
有关如何安装内核调谐器及其 依赖项可以在installation guide
示例用法
下面是一个调整CUDA内核的简单示例:
kernel_string=""" __global__ void vector_add(float *c, float *a, float *b, int n) { int i = blockIdx.x * block_size_x + threadIdx.x; if (i<n) { c[i] = a[i] + b[i]; } } """size=10000000a=numpy.random.randn(size).astype(numpy.float32)b=numpy.random.randn(size).astype(numpy.float32)c=numpy.zeros_like(b)n=numpy.int32(size)args=[c,a,b,n]tune_params=dict()tune_params["block_size_x"]=[32,64,128,256,512]tune_kernel("vector_add",kernel_string,size,args,tune_params)
完全相同的python代码可用于优化opencl内核:
kernel_string=""" __kernel void vector_add(__global float *c, __global float *a, __global float *b, int n) { int i = get_global_id(0); if (i<n) { c[i] = a[i] + b[i]; } } """
内核调谐器将检测内核语言并选择正确的编译器和 运行时。对于参数空间中的每个内核,内核调谐器将插入c 预处理器为可调参数定义、编译和基准测试内核。这个 计时结果将打印到控制台,但也由tune_kernel返回到 允许进一步分析。注意,这只是默认行为,什么和如何 tune-kernel所做的一切都是通过它的多个optional arguments来控制的。
搜索调整策略
内核调谐器支持许多优化算法来加速自动调整过程。目前 实现的搜索算法有:暴力(默认)、nelder mead、powell、cg、bfgs、l-bfgs-b、tnc, cobyla,slsqp,随机搜索,basinhopping,差分进化,遗传算法,粒子群 优化、萤火虫算法和模拟退火。
使用搜索策略很简单,只需指定tune_kernel哪种策略和方法 您想要使用,例如strategy="genetic_algorithm"或strategy="basinhopping",method="Powell"。有关支持的搜索策略和方法的完整概述,请参见user api documentation。
调整主机和内核代码
可以在 主机和内核代码。这就允许了很多强大的东西, 例如,为使用CUDA流的内核调整流的数量 或opencl命令队列以重叠主机和设备之间的传输 内核执行。这可以与调整 内核代码中的参数。见convolution_streams example code 以及 documentation 有关内核调谐器python脚本的详细说明。
正确性验证
或者,您可以让内核调谐器验证 内核它通过传递一个answer列表来编译和基准测试。这个 list与内核的参数列表匹配,但包含 内核的预期输出。输入参数将替换为“无”。
answer=[a+b,None,None]# the order matches the arguments (in args) to the kerneltune_kernel("vector_add",kernel_string,size,args,tune_params,answer=answer)
贡献
引文
一篇关于内核调谐器的科学论文已被接受出版,请引用 内核调谐器如下:
@article{kerneltuner, author = {Ben van Werkhoven}, title = {Kernel Tuner: A search-optimizing GPU code auto-tuner}, journal = {Future Generation Computer Systems}, year = {2019}, volume = {90}, pages = {347-358}, doi = {https://doi.org/10.1016/j.future.2018.08.004}, }