m.y)。
[0024]第二步，根据GPU线程分配方式申请显存空间，将输入数据由内存拷贝到显存中，即将CPU工作移到GPU，在本发明中，在CPU端，利用多核处理器的多线程计算能力，使整个步骤构建成一个流水线结构。同时，根据GPU线程分配方式的计算结果，通过CUDA将此步骤部署到GPU的每一个grid、block以及thread，保证其高度并行化地执行。
[0025]第三步，GPU在每一个线程Thread上进行核函数计算，核函数计算以多线程方式进行，在每个线程中独立运行着一个核函数，对每个字节进行抽阶操作。其具体步骤如下:
(I)计算当前线程偏移量。根据当前线程的栅格坐标计算全局线程编号，通过全局线程编号计算出当前线程在缓冲区中的偏移量。在流处理单元SM中并发的所有线程的执行顺序是按threadO、threadl...threadn的顺序依次执行的，在这个过程中，通过线程同步使得每个block中的线程并发进入访存语句，并按照线程号从小到大按顺序依次执行，使得threadO在等待访存结果时，threadl能够立即开始访存操作，依次类推，屏蔽了大部分线程的访存延时，节约了执行时间。而在此便需要对每个线程计算出其位于缓冲区中的偏移量，实现按照线程号从小到大按顺序依次执行。其具体步骤如下:
A、计算当前线程所在块的编号bid，其计算公式如下: bid=gridDim.x^blockldx.y+blockldx.x ；
其中，blockldx.y为当前线程所在线程块Block中的列号，blockldx.x为当前线程所在线程块Block中的行号。
[0026]B、计算当前块内线程编号cur_tid，其计算公式如下: cur_tid=b1ckDim.x*threadldx.y+threadldx.x ；
同样，blockldx.y为当前线程所在线程块Block中的列号，blockldx.x为当前线程所在线程块Block中的行号。
[0027]C、计算全局线程编号total_tid，其计算公式如下: total_tid=bid*b1ckDim.x*blockDim.y+cur_tid。
[0028]D、根据全局线程编号total_tid确定当前线程输入输出数据在缓冲区中的偏移量offset，其计算公式如下:
offset=total_tid*(blockDim.x*blockDim.y)*(N/8)，
其中8是一个字节的位数，N为抽取的阶数，N=l、2或4。
[0029](2)根据线程的偏移量从显存中取出数据存入线程块内共享显存，显存和线程块内共享显存在物理上均是位于GPU上的存储设备，在逻辑上两者则不同，显存为从内存中拷贝数据，而线程块内共享显存则为开辟于线程块内的缓存，从显存获取数据，并在共享显存内进行线程的核函数计算。
[0030](3)由于CUDA的技术特性，在流处理单元SM中的执行顺序决定了 threadO会优先执行，因此针对满足tid.X=O的所有线程来执行其所在二维线程块Block输入数据的抽阶操作。具体的抽阶操作根据栅格化数据处理来决定，从每一个字节取出特定的位，栅格化时数据的组织方式决定了该取哪一位或者哪几位(目前应用中高四位均为有效数据)，即根据栅格化处理规则从每一个字节取出特定的位，将抽阶的结果暂存到结果mask中。这样便可让所有线程分组执行既节约了运算资源，又避免了线程间数据冲突，保证了并行性。在此可以设缓存mask为4个字节，则处理完毕后，一个块的数据抽阶后只剩下了 32位，被组织成一个int型(4字节)，所以只需要四个线程来拷贝结果数据。
[0031](4)设int型为4字节，使用tid.x〈4的线程Thread将结果数据从结果缓存mask中拷贝至对应的显存中，核函数执行完毕。
[0032]第四步，待所有线程Thread的核函数计算完后，将显存中的结构数据拷贝回内存，此结构数据为抽阶后的栅格化数据，完成抽阶过程。
[0033]本发明利用GPU高并行计算性能，以及CUDA框架的层次化并行特性，将直写式光刻机数据处理中的栅格化数据抽阶过程进行了并行优化，用大量的线程并行执行来加快执行速度，提高了直写式光刻机的产能。同时，在实现并行化的过程中，充分地利用了 GPU的线程资源，实现了最大化的加速比，利用了共享显存的高带宽特性，提高了处理速度，读写全局显存时根据线程编号同步操作，有效的屏蔽了访存延时，进一步提高了处理效率。
[0034]以上显示和描述了本发明的基本原理、主要特征和本发明的优点。本行业的技术人员应该了解，本发明不受上述实施例的限制，上述实施例和说明书中描述的只是本发明的原理，在不脱离本发明精神和范围的前提下本发明还会有各种变化和改进，这些变化和改进都落入要求保护的本发明的范围内。本发明要求的保护范围由所附的权利要求书及其等同物界定。
【主权项】
1.一种基于CUDA技术对栅格化数据进行抽阶的方法，其特征在于，包括以下步骤: 11) CPU分配显存和计算资源，CPU根据进行抽阶数据的规模，结合当前GPU可供使用的硬件资源，计算出最优的GPU线程分配方式； 12 )根据GPU线程分配方式申请显存空间，将输入数据由内存拷贝到显存中； 13)GPU在每一个线程Thread上进行核函数计算，对每个字节进行抽阶操作； 14)待所有线程Thread的核函数计算完后，将显存中的结构数据拷贝回内存，此结构数据为抽阶后的栅格化数据，完成抽阶过程。2.根据权利要求1所述的一种基于CUDA技术对栅格化数据进行抽阶的方法，其特征在于，所述的CPU分配显存和计算资源包括以下步骤: 21)输入栅格化处理后的二维位图像素阵列，其宽度定义为width，高度定义为height ； 22)将每个二维线程块Block的宽度定义为blockDim.X、高度定义为blockDim.y ; 23)计算出线程栅格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 ; 24)获得线程分配方式， 线程分配方式中二维线程块Block为Block (blockDim.x, blockDim.y)；出线程栅格Grid 为 Grid(gridDim.x, gridDim.y)。3.根据权利要求1所述的一种基于CUDA技术对栅格化数据进行抽阶的方法，其特征在于，所述的GPU在每一个线程Thread上进行核函数计算包括以下步骤: 31)计算当前线程偏移量，根据当前线程的栅格坐标计算全局线程编号，通过全局线程编号计算出当前线程在缓冲区中的偏移量； 32)根据线程的偏移量从显存中取出数据存入线程块内共享显存； 33)针对满足tid.x=0的所有线程来执行其所在二维线程块Block输入数据的抽阶操作；根据栅格化处理规则从每一个字节取出特定的位，将抽阶的结果数据暂存到结果缓存mask中，其中设缓存mask为4个字节； 34)设int型为4字节，使用tid.x<4的线程Thread将结果数据从结果缓存mask中拷贝至对应的显存中。4.根据权利要求3所述的一种基于CUDA技术对栅格化数据进行抽阶的方法，其特征在于，所述的计算当前线程偏移量包括以下步骤: 41)计算当前线程所在块的编号bid，其计算公式如下: bid=gridDim.x^blockldx.y+blockldx.x ； 其中，blockldx.y为当前线程所在线程块Block中的列号，blockldx.x为当前线程所在线程块Block中的行号； 42)计算当前块内线程编号cur_tid，其计算公式如下:cur_tid=blockDim.x*threadldx.y+threadldx.x ； 其中，blockldx.y为当前线程所在线程块Block中的列号，blockldx.x为当前线程所在线程块Block中的行号； 43)计算全局线程编号total_tid，其计算公式如下: total_tid=bid*b1ckDim.x*blockDim.y+cur_tid ； 44)根据全局线程编号确定当前线程输入输出数据在缓冲区中的偏移量offset，其计算公式如下:offset=total_tid*(blockDim.x*blockDim.y)*(N/8)， 其中8是一个字节的位数，N为抽取的阶数，N=l、2或4。
【专利摘要】本发明涉及一种基于CUDA技术对栅格化数据进行抽阶的方法，与现有技术相比解决了栅格化数据抽阶效率较低的缺陷。本发明包括以下步骤：CPU分配显存和计算资源，CPU根据进行抽阶数据的规模，结合当前GPU可供使用的硬件资源，计算出最优的GPU线程分配方式；根据GPU线程分配方式申请显存空间，将输入数据由内存拷贝到显存中；GPU在每一个线程Thread上进行核函数计算，对每个字节进行抽阶操作；完成抽阶过程。本发明通过并行化提高了计算效率，增加了直写式光刻机的产能，同时降低了数据规模，减少了对计算能力及传输带宽的依赖，降低了成本。
【IPC分类】G06T1/20, G06F9/50
【公开号】CN105183562
【申请号】CN201510566712
【发明人】陆敏婷
【申请人】合肥芯碁微电子装备有限公司
【公开日】2015年12月23日
【申请日】2015年9月9日