一种基于cuda技术对栅格化数据进行抽阶的方法
【专利说明】一种基于CUDA技术对栅格化数据进行抽阶的方法
[0001]
技术领域
[0002]本发明涉及直写式光刻机数据处理技术，具体来说是一种基于CUDA技术对栅格化数据进行抽阶的方法。
[0003]
【背景技术】
[0004]CUDA是NVIDIA公司2007年提出的支持GPU进行通用计算的编程模型和开发环境，CUDA编程的思想是用海量的线程来开发程序中的并行性，海量线程以层次化的方式组织，单个的线程被映射到标量核SP上执行，一组线程被组织成一个线程块(Block)被映射到一个流处理单位SM上执行，最后由线程块组成的线程栅格(Grid)映射到一个GPGPU(GPU)上执行。由于GPU具有远超CPU的计算核心数以及海量的并行计算资源，适合进行计算密集型、高度并行化的计算任务。同时，由于GPU的价格远远低于同等性能的并行计算系统，由CPU和GPGPU (GPU)组成的异构系统已经越来越广的应用到生物医学、流体力学等诸多工程应用领域。
[0005]直写式光刻机的数据处理过程是将用户提供的矢量数据，转化为图形发生器能接受的图像数据，数据处理过程中涉及到数据的分析、计算和传输。目前实际应用处理得到的栅格化数据中，一个像素用一个字节来表示(8阶灰度)，而下位机只需要其中的1、2、4位即可满足显示的灰度要求，因此如果能够针对栅格化数据进行抽阶处理，去掉其中冗余的数据，提取出有效的灰度值，即可降低了数据规模，降低传输链路带宽。如针对同样一幅图，抽阶后的数据量变少了，传输这些数据要求的时间不变，所以需要的带宽(传输速率)降低了。实际应用中选择成本低，速度慢的传输链路也可满足传输时间要求，则相当于降低了生产成本。
[0006]对栅格化数据进行抽阶处理是根据实际需要来处理，如针对4位的灰度要求，在对数据栅格化时便可以将需要的4位排列在8位字节(一个像素)的前4位，在做抽阶工作时，直接抽取0-3位即可；或针对2位的灰度要求，在数据栅格化时将需要的2位排在8位字节的第2位和第3位，在做抽阶工作时，直接抽取1-2位即可。但是目前栅格化数据的数据量过于庞大，导致抽阶工作较慢，抽阶工作的分析、计算和传输均比较耗时，难以满足产能要求，如何利用CUDA技术的特点，实现栅格化数据抽阶的多线程并行处理已经成为急需解决的技术问题。
[0007]
【发明内容】
[0008]本发明的目的是为了解决现有技术中栅格化数据抽阶效率较低的缺陷，提供一种基于CUDA技术对栅格化数据进行抽阶的方法来解决上述问题。
[0009]为了实现上述目的，本发明的技术方案如下:
一种基于CUDA技术对栅格化数据进行抽阶的方法，包括以下步骤:
CPU分配显存和计算资源，CPU根据进行抽阶数据的规模，结合当前GPU可供使用的硬件资源，计算出最优的GPU线程分配方式；
根据GPU线程分配方式申请显存空间，将输入数据由内存拷贝到显存中；
GPU在每一个线程Thread上进行核函数计算，对每个字节进行抽阶操作；
待所有线程Thread的核函数计算完后，将显存中的结构数据拷贝回内存，此结构数据为抽阶后的栅格化数据，完成抽阶过程。
[0010]所述的CPU分配显存和计算资源包括以下步骤:
输入栅格化处理后的二维位图像素阵列，其宽度定义为width，高度定义为height ； 将每个二维线程块Block的宽度定义为blockDim.X、高度定义为blockDim.y ；
计算出线程栅格Grid的宽度gridDim.X，其计算公式如下: gridDim.x = width/blockDim.x ；
计算出线程栅格Grid的高度gridDim.y，其计算公式如下: gridDim.y = height/b1ckDim.y ；
计算出显存分配总大小length，其计算公式如下:
Iength=Width 氺height 氺(1+N/8)，
其中，N=l、2或4 ;
获得线程分配方式，
线程分配方式中二维线程块Block为Block (blockDim.x, blockDim.y)；出线程栅格Grid 为 Grid(gridDim.x, gridDim.y)。
[0011]所述的GPU在每一个线程Thread上进行核函数计算包括以下步骤:
计算当前线程偏移量，根据当前线程的栅格坐标计算全局线程编号，通过全局线程编号计算出当前线程在缓冲区中的偏移量；
根据线程的偏移量从显存中取出数据存入线程块内共享显存；
针对满足tid.x=0的所有线程来执行其所在二维线程块Block输入数据的抽阶操作；根据栅格化处理规则从每一个字节取出特定的位，将抽阶的结果数据暂存到结果缓存mask中，其中设缓存mask为4个字节；
设int型为4字节，使用tid.x<4的线程Thread将结果数据从结果缓存mask中拷贝至对应的显存中。
[0012]所述的计算当前线程偏移量包括以下步骤:
计算当前线程所在块的编号Md，其计算公式如下: bid=gridDim.x^blockldx.y+blockldx.x ；
其中，blockldx.y为当前线程所在线程块Block中的列号，blockldx.x为当前线程所在线程块Block中的行号；
计算当前块内线程编号cur_tid，其计算公式如下: cur_tid=b1ckDim.x*threadldx.y+threadldx.x ；
其中，blockldx.y为当前线程所在线程块Block中的列号，blockldx.x为当前线程所在线程块Block中的行号； 计算全局线程编号total_tid，其计算公式如下: total_tid=bid*b1ckDim.x*blockDim.y+cur_tid ；
根据全局线程编号确定当前线程输入输出数据在缓冲区中的偏移量offset，其计算公式如下:
offset=total_tid*(blockDim.x*blockDim.y)*(N/8)，
其中8是一个字节的位数，N为抽取的阶数，N=l、2或4。
[0013]
有益效果
本发明的一种基于CUDA技术对栅格化数据进行抽阶的方法，与现有技术相比通过并行化提高了计算效率，增加了直写式光刻机的产能，同时降低了数据规模，减少了对计算能力及传输带宽的依赖，降低了成本。
[0014]本发明在直写式光刻机的数据处理过程中，发掘出了其中的计算密集型、高度并行化过程，并将这个过程并行化，通过CUDA将其部署到GPU上并行执行，极大的提高了处理速度。并且，在实现并行化的过程中，充分地利用了 GPU以及CUDA框架的特性，实现了最大化的加速比；在分配线程资源时，根据当前硬件的最大线程数、线程块中的最优线程数来确定Block和Grid的尺寸；核函数在处理时先将输入数据拷贝到共享显存，充分利用了共享显存的高带宽特性，提高了处理速度；读写全局显存时根据线程编号同步操作，有效的屏蔽了访存延时，进一步提高了处理效率。
[0015]
【附图说明】
图1为本发明的方法顺序图；
图2为本发明中CUDA线程栅格示意图。
[0016]
【具体实施方式】
[0017]为使对本发明的结构特征及所达成的功效有更进一步的了解与认识，用以较佳的实施例及附图配合详细的说明，说明如下:
如图1所示，本发明所述的一种基于CUDA技术对栅格化数据进行抽阶的方法，抽阶的过程在GPU中并行进行，利用CUDA框架实现。其包括以下步骤:
第一步，CPU分配显存和计算资源，CPU根据进行抽阶数据的规模，结合当前GPU可供使用的硬件资源，计算出最优的GPU线程分配方式，为后面的步骤分配显存和计算资源。其具体包括以下步骤:
(I)输入栅格化处理后的二维位图像素阵列，二维位图像素阵列即为栅格化数据，其宽度定义为width，高度定义为height。
[0018](2)根据CUDA的技术要求，每个二维线程块Block的宽度为blockDim.X、高度为blockDim.yD
[0019](3)计算出线程栅格Grid的宽度gridDim.X，其计算公式如下: gridDim.x = width/blockDim.χ0
[0020]计算出线程栅格Grid的高度gridDim.y，其计算公式如下: gridDim.y = height/b1ckDim.y。
[0021]计算出显存分配总大小length，其计算公式如下:
Iength=Width 氺height 氺(1+N/8)，
其中，N=l、2或4。
[0022]如图2所示，在此步骤后获得最优的GPU线程分配方式，在GPU (device端)共有gridDim.X * gridDim.y个二维线程块Block并行执行，每个二维线程块Block中有BlockDim.x * BlockDim.y个线程Thread并行执行，通过这种高度并行化极大地提高了执行效率。
[0023](4)获得线程分配方式，
线程分配方式中二维线程块Block为Block (blockDim.x, blockDim.y)；出线程栅格Grid 为 Grid(gridDim.x, gridDi