驱动执行单元的机器码产生方法以及装置的制造方法
【技术领域】
[0001]本发明涉及一种图形处理单元技术，特别是一种驱动执行单元的机器码产生方法以及装置。
【背景技术】
[0002]在通用计算图形处理单元(GPGPU,General-Purpose Computing on GraphicsProcessing Unit)模型，例如，开放计算语言(OpenCL, Open Computing Language)、清染脚本(rendering scripts)、DirectX (Direct extens1n)及开放图形语言(OpenGL, OpenGraphics Language)等计算着色器(CS, Compute Shader)，大部分的输出输入都由存储器的加载(load)、储存(store)、缩减(reduct1n)及采样(sample)操作数完成。如何提升加载、储存、缩减及取样作业的效能，是个重要议题。因此，本发明提出一种驱动执行单元的机器码产生方法以及装置，不仅减少存储器地址的运算成本，更可以提升暂存数据的命中率。
【发明内容】
[0003]本发明的实施例提出一种驱动执行单元的机器码产生方法，使用编译器执行，包含下列步骤。搜集核心程序中关联于存储器表面的多个数据存取指令；分析数据存取指令关联的地址样式，用以产生全域识别码地址；以及产生包含全域识别码地址的机器码。
【附图说明】
[0004]图1是依据本发明实施例的驱动执行单元的机器码产生装置的硬件架构图。
[0005]图2是依据本发明实施例的驱动执行单元的机器码产生方法流程图。
[0006]图3A至3C是依据本发明实施例的驱动执行单元的机器码产生方法流程图。
【具体实施方式】
[0007]以下说明为完成发明的较佳实现方式，其目的在于描述本发明的基本精神，但并不用以限定本发明。实际的
【发明内容】
必须参考之后的权利要求范围。
[0008]必须了解的是，使用于本说明书中的“包含”、“包括”等词，用以表示存在特定的技术特征、数值、方法步骤、作业处理、组件以及/或组件，但并不排除可加上更多的技术特征、数值、方法步骤、作业处理、组件、组件，或以上的任意组合。
[0009]在权利要求中使用如“第一”、“第二”、“第三”等词是用来修饰权利要求中的组件，并非用来表示之间具有优先权顺序，先行关系，或者是一个组件先于另一个组件，或者是执行方法步骤时的时间先后顺序，仅用来区别具有相同名字的组件。
[0010]在通用计算图形处理单元(GPGPU)模型，执行于电子装置的主程序(mainfunct1n)称为核心程序(Kernel)。以下为范例核心程序的伪码，用以将二个存储器平面(surfaces)相加并输出:
[0011]_kernel void VectorAdd(_global float4*A, _global float4*B, _globalfloat4*C, int WidthA, int WidthB, int WidthC){
[0012]int x = get_global_id (0)；
[0013]int y = get_global_id (1)；
[0014]int addressA = (y+2) *ffidthA+x+5 ；
[0015]int addressB = (y+3)*ffidthB+x+6 ；
[0016]int addressC = y*ffidthC+x ；
[0017]C[addressC] = A[addressA]+B[addressB] ；}
[0018]核心程序的加载、储存及取样作业的地址以全域识别码(global-1ds)为基础，为二维/三维(2D/3D)空间地址。在一些实施方式中，编译器直接依据二维/三维地址为加载、储存及取样作业计算一维地址，并产生机器码。范例的机器码如下所示:
[0019]IMUL R0, C[4], R0
[0020]IADD R0，R4，R0
[0021]IADD R0，C[7]，R0 //x = get_global_id (0)
[0022]IMUL Rl, C[5], R1
[0023]IADD R1，R5，R1
[0024]IADD Rl, C[8], Rl //y = get_global_id(1)
[0025]IADD R2, Rl, 2
[0026]IMUL R2, CB[16], R2
[0027]IADD R2，R2，R0 //(y+2) *ffidthA+x ；
[0028]IADD R2, R2, 5 //(y+2)*ffidthA+x+5 ；
[0029]IADD R2, C[13], R2?4
[0030]LD R4.xyzw, R2 // 从 A [addressA]加载数据
[0031]IADD R3, Rl, 3
[0032]IMUL R3, C[17], R3
[0033]IADD R3，R3，R0 //(y+3) *ffidth+x ；
[0034]IADD R3，R3，6 //(y+3) *ffidth+x+6 ；
[0035]IADD R3, C[14], R3?4
[0036]LD R8.xyzw, R3 // 从 B [addressA]加载数据
[0037]FADD.rp3 R4, R4, R8 //A[addressA]+B[addressB]
[0038]IMUL Rl, C[18], Rl
[0039]IADD Rl, Rl, R0 //y*ffidth+x ；
[0040]IADD Rl, C[15], Rl?4
[0041 ] ST R4.xyzw, Rl //储存数据至 C [addressC]
[0042]IMUL代表整数乘法操作数，IADD代表整数加法操作数，FADD代表浮点数加法操作数，LD代表载入操作数，ST代表储存操作数。这些计算出来的存储器地址会直接送到执行单元，用以存取数据。然而，此编译后的机器码需要较多的地址运算作业。
[0043]图1是依据本发明实施例的驱动执行单元的机器码产生装置的硬件架构图。编译器(compiler) 130搜集每一个存储器表面(memory surface)的属性(attributes)，存储器表面指一段存储器地址的数据。详细来说，编译器130扫描核心程序110的内容，用以获得关联于同一段存储器地址(亦即是同一个存储器表面)的数据存取指令。编译器130分析搜集到的数据存取指令的变量的数据形态，用以判断此存储器表面是否只读以及是否为专属形态缓存区或影像(type buffer or image)。编译器130另判断数据存取指令的变量的数据形态是否为可支持数据形态，若不是，则将原始指令改变为可支持数据形态的指令。编译器130更分析搜集到的数据存取指令的地址计算过程，用以产生以全域识别码地址(global-1d addresses)。最后，编译器130根据此存储器表面是否只读以及是否为专属形态缓存区或影像的信息、原始的以及/或改变后的指令、以及原始的或新产生的存储器地址产生机器码(又称为二进制代码)150。驱动器(driver) 170依据加载的机器码150设定适当的寄存器(register)，用以驱动执行单元(EU, Execut1n Unit) 200完成特定的运算。执行单元200可为顶点着色器(VS, Vertex Shader)、外壳着色器(HS，Hull Shader)、区域着色器(DS, Domain Shader)、几何着色器(GS, Geometry Shader)、画素着色器(PS, PixelShader)，或在三维图形管道中的其它硬件。驱动器170可以专用硬件电路或通用硬件(例如，单一处理器、具并行处理能力的多处理器或其它具运算能力的处理器)实施。
[0044]图2是依据本发明实施例的驱动执行单元的机器码产生方法流程图。此方法由编译器130完成。编译器130反复执行一个循环，用以产生关联于每一个存储器表面的加载、储存、缩减及采样指令的机器码150，并且指令中的存储器地址为全域识别码地址。在每一回合中，编译器130搜集核心程序110中关联于一个存储器表面的多个数据存取指令(步骤S210)。存储器表面可指一段存储器地址中储存的数据。接着，分析数据存取指令关联的地址样式，用以产生全域识别码地址(步骤S230)，以及产生包含全域识别码地址的机器码150(步骤S250)。此后，流程判断是否分析完核心程序(步骤S270) ?如果不是，则返回步骤S210循环操作；如果是，则结束该方法。
[0045]图3A至3C是依据本发明实施例的驱动执行单元的机器码产生方法流程图。此方法由编译器130完成。编译器130反复执行一个循环，用以产生关联于每一个存储器表面的加载、储存、缩减及采样指令的机器码150，并且指令中的存储器地址为全域识别码地址。在每一回合中，编译器130可包含三个阶段的分析。第一阶段分析关联于一个存储器表面的所有数据存取指令的类型，第二阶段分析数据存取指令中的变量的数据形态，以及第三阶段分析数据存取指令关联的地址样式(address pattern)。编译器130根据分析结果在需要的情况下改变原始指令为新指令，以及产生全域识别码地址。最后，编译器130根据原始的以及/或更新后的指令以及全域识别码地址产生机器码150。
[0046]详细来说，首先，编译器130搜集核心程序110中关联于一个存储器表面的多个数据存取指令(步骤S311)。存储器表面可指一段存储器地址中储存的数据。在第一阶段的分析中，编译器130判断数据存取指令是否关联于采样指令(步骤S312)。若是，判断此存储器表面为专属形态影像(type image)(步骤S314)。否则，编译器130更判断此存储器表面是否为只读(步骤S313)。在步骤S313，详细来说，当关联于此存储器表面的指令只有加载指令时，代表此存储器表面为只读。在第二阶段的分析中，编译器130判断数据存取指令中的所有变量的数据形态是否相同(步骤S315以及S317)。通过第一及第二阶段的分析结果，编译器130可判断此存储器表面的类型。详细来说，当存储器表面为只读及数据存取指令的所有变量的数据形态都相同时(步骤S313及步骤S315中“是”的路径)，判断此存储器表面为专属形态缓存区(type buffer)(步骤S321)。当存储器表面为只读及数据存取指令的变量的数据形态不同时(步骤S313中“是”的路径接步骤S315中“否”的路径)，判断此存储器表面为原始缓存区(raw buffer)(步骤S391)。当存储器表面非只读及数据存取指令中的变量的数据形态不同时(步骤S313及步骤S315中“否”的路径)，判断此存储器表面为非专属形态随机存取视图(un-Type UAV, Un-order Accessed View)(步骤S341)。当存储器表面非只读及数据存取指令的所有变量的数据形态相同时(步骤S313中“否”的路径接步骤S315中“是”的路径)，判断此存储器表面为专属形态随机存取视图(Type UAV)(步骤S331)。此外，当判断此存储器表面为专属形态缓存区或专属形态随机存取视图后(步骤S331或S321)，判断数据存取指令的变量的数据形态是否为可支持数据形态(步骤S351)。若否，编译器130改变原始的数据存取指令成为包含可支持数据形态的新指令(步骤S353)。例如，一个数据存取指令包含数据形态为CHAR16的变量，用以一次加载四