【技术实现步骤摘要】
本公开涉及计算机架构，尤其涉及一种在单指令多线程(single instructionmultiple threads，simt)计算系统中使用标量存储器地址来加载和存储多线程的向量数据的多线程计算机架构。
技术介绍
1、图形处理单元(graphics processing unit，gpu)架构提供了一种以单指令多线程(single instruction multiple threads，simt)方式执行并行线程的方法。诸如gpu的simt处理器具有多个用于同时执行多个线程的核，并且特别适合于大规模并行计算应用。为利用上述用于并行执行的多个核，计算机程序通常需要通过调用专门为多核工作而设计的应用编程接口(application programming interface，api)的函数来为多核架构进行定制。最近，利用gpu在传统上由中央处理器(central processing unit，cpu)处理的应用中执行计算的通用计算gpu(general-purpose computing on gpu，gpgpu)变得更加实用和流行。然而，为多个并发线程加载和存储数据需要每个线程的数据地址。上述数据地址通常需要事先通过向量运算来准备，并且需要使用向量地址总线。因此，本领域需要更有效地为多个线程加载和存储数据的处理器。
技术实现思路
1、本公开描述了使用标量存储器地址有效地加载向量数据的装置、方法和系统。在一示例性实施例中，一种方法可以包括：序列发生器向存储器端口发送序列发生器的
2、在另一示例性实施例中，处理器可以包括序列发生器、被耦合至该序列发生器的存储器端口、以及经由存储器接口耦合至该存储器端口的存储器单元。序列发生器可以用于向存储器端口发送序列发生器的标量寄存器中存储的线程块大小信息和寻址参数，并且存储器端口可以用于基于线程块大小信息和寻址参数生成标量存储器地址，所述标量存储器地址指向数据warp的第一线程的数据的存储位置。存储器单元和存储器接口可以用于：对于数据加载过程，从起始于标量存储器地址的k个连续存储器地址将warp的数据的k个字加载到向量寄存器，或者对于数据存储过程，从向量寄存器将warp的数据的k个字存储到起始于标量存储器地址的k个连续存储器地址。k可以是warp大小，warp大小表示要通过执行向量指令来同时处理的线程数。该数据的大小可以小于或等于字大小。
3、在另一实施例中，一种方法可以包括：序列发生器向存储器端口发送序列发生器的标量寄存器中存储的线程块大小信息和寻址参数；存储器端口基于线程块大小信息和寻址参数生成标量存储器地址；存储器接口和存储器单元从向量寄存器将warp的数据的k个字存储到起始于所述标量存储器地址的k个连续存储器地址。标量存储器地址可以指向数据warp的第一线程的数据的存储位置，并且k可以是warp大小，warp大小表示要通过执行向量指令来同时处理的线程数。该数据的大小可以小于或等于字大小。
本文档来自技高网...【技术保护点】
1.一种方法，包括：
2.根据权利要求1所述的方法，还包括：
3.根据权利要求2所述的方法，其中，所述线程块大小信息包括x轴上的线程大小的第一值X、y轴上的线程大小的第二值Y、以及X乘以Y和Z的乘法结果XYZ，其中Z是z轴上的线程大小，并且当所述存储器单元是字可寻址存储器单元时，所述寻址参数包括BASE、STRIDE_Y、以及STRIDE_Z，或者当所述存储器单元是字节可寻址存储器单元时，所述寻址参数包括BASE、STRIDE_Y、STRIDE_Z、以及DS，其中BASE是所述内核程序要处理的数据的起始存储器地址，对于所述字可寻址存储器单元，STRIDE_Y和STRIDE_Z分别是y轴和z轴上的存储器地址的间隙，或者对于所述字节可寻址存储器单元，STRIDE_Y和STRIDE_Z分别是y轴和z轴上的存储器地址的间隙除以DS，并且DS是以字节为单位的所述存储器单元的数据大小。
4.根据权利要求3所述的方法，其中，所述存储器单元包括多个存储器组，所述标量存储器地址包括作为存储器组索引的组选择位，所述存储器组索引指向所述多个存储器组中的存储器组，并且
5.根据权利要求4所述的方法，其中，所述标量存储器地址由所述存储器端口计算，对于所述字可寻址存储器单元，计算为ADDR＝BASE+threadIdx.x+threadIdx.y*STRIDE_Y+threadIdx.z*STRIDE_Z，或者对于所述字节可寻址存储器单元，计算为ADDR＝BASE+(threadIdx.x+threadIdx.y*STRIDE_Y+threadIdx.z*STRIDE_Z)*DS，其中threadIdx.x是所述内核程序的线程块中x轴上的线程索引，并且是K的倍数，threadIdx.y是所述线程块中y轴上的线程索引，threadIdx.z是所述线程块中z轴上的线程索引。
6.根据权利要求5所述的方法，其中，所述内核程序的所述线程块被映射到图像的数据集中的矩形窗口，并且BASE是所述矩形窗口左上角的存储器地址。
7.根据权利要求1所述的方法，其中，所述存储器端口通过包括标量存储器地址总线的私有存储器地址端口经由所述存储器接口耦合至所述存储器单元，所述标量存储器地址经由所述标量存储器地址总线传送到所述存储器接口，并且如果所述存储器单元是字节可寻址存储器单元，则数据大小DS经由所述标量存储器地址总线的数据大小位传送到所述存储器接口。
8.根据权利要求1所述的方法，其中，所述向量寄存器被映射到交换盒的输出数据缓冲器，或是被映射到从所述存储器端口到所述交换盒的读数据连接的虚拟向量寄存器。
9.一种处理器，包括：
10.根据权利要求9所述的处理器，所述序列发生器还用于从所述序列发生器的任务缓冲器加载线程块索引信息，并执行内核程序中的标量指令，以：
11.根据权利要求10所述的处理器，所述线程块大小信息包括x轴上的线程大小的第一值X、y轴上的线程大小的第二值Y、以及X乘以Y和Z的乘法结果XYZ，其中Z是z轴上的线程大小，并且当所述存储器单元是字可寻址存储器单元时，所述寻址参数包括BASE、STRIDE_Y、以及STRIDE_Z，或者当所述存储器单元是字节可寻址存储器单元时，所述寻址参数包括BASE、STRIDE_Y、STRIDE_Z、以及DS，其中BASE是所述内核程序要处理的数据的起始存储器地址，对于所述字可寻址存储器单元，STRIDE_Y和STRIDE_Z分别是y轴和z轴上的存储器地址的间隙，或者对于所述字节可寻址存储器单元，STRIDE_Y和STRIDE_Z分别是y轴和z轴上的存储器地址的间隙除以DS，并且DS是以字节为单位的所述存储器单元的数据大小。
12.根据权利要求11所述的处理器，其中，所述存储器单元包括多个存储器组，所述标量存储器地址包括作为存储器组索引的组选择位，所述存储器组索引指向所述多个存储器组中的存储器组，并且其中，所述存储器接口用于使用从通过所述组选择位选择的所述存储器组开始的所述标量存储器地址，生成多个组地址，其中每个存储器组一个组地址，以及
13.根据权利要求12所述的处理器，其中，所述标量存储器地址由所述存储器端口计算，对于所述字可寻址存储器单元，计算为ADDR＝BASE+threadIdx.x+threadIdx.y*STRIDE_Y+threadIdx.z*STRIDE_Z，或者对于...
【技术特征摘要】
1.一种方法，包括：
2.根据权利要求1所述的方法，还包括：
3.根据权利要求2所述的方法，其中，所述线程块大小信息包括x轴上的线程大小的第一值x、y轴上的线程大小的第二值y、以及x乘以y和z的乘法结果xyz，其中z是z轴上的线程大小，并且当所述存储器单元是字可寻址存储器单元时，所述寻址参数包括base、stride_y、以及stride_z，或者当所述存储器单元是字节可寻址存储器单元时，所述寻址参数包括base、stride_y、stride_z、以及ds，其中base是所述内核程序要处理的数据的起始存储器地址，对于所述字可寻址存储器单元，stride_y和stride_z分别是y轴和z轴上的存储器地址的间隙，或者对于所述字节可寻址存储器单元，stride_y和stride_z分别是y轴和z轴上的存储器地址的间隙除以ds，并且ds是以字节为单位的所述存储器单元的数据大小。
4.根据权利要求3所述的方法，其中，所述存储器单元包括多个存储器组，所述标量存储器地址包括作为存储器组索引的组选择位，所述存储器组索引指向所述多个存储器组中的存储器组，并且其中，所述存储器接口用于使用从通过所述组选择位选择的所述存储器组开始的所述标量存储器地址，生成多个组地址，其中每个存储器组一个组地址，并且按照相应线程的顺序，将从所述多个存储器组接收的数据的k个字重新排列。
5.根据权利要求4所述的方法，其中，所述标量存储器地址由所述存储器端口计算，对于所述字可寻址存储器单元，计算为addr＝base+threadidx.x+threadidx.y*stride_y+threadidx.z*stride_z，或者对于所述字节可寻址存储器单元，计算为addr＝base+(threadidx.x+threadidx.y*stride_y+threadidx.z*stride_z)*ds，其中threadidx.x是所述内核程序的线程块中x轴上的线程索引，并且是k的倍数，threadidx.y是所述线程块中y轴上的线程索引，threadidx.z是所述线程块中z轴上的线程索引。
6.根据权利要求5所述的方法，其中，所述内核程序的所述线程块被映射到图像的数据集中的矩形窗口，并且base是所述矩形窗口左上角的存储器地址。
7.根据权利要求1所述的方法，其中，所述存储器端口通过包括标量存储器地址总线的私有存储器地址端口经由所述存储器接口耦合至所述存储器单元，所述标量存储器地址经由所述标量存储器地址总线传送到所述存储器接口，并且如果所述存储器单元是字节可寻址存储器单元，则数据大小ds经由所述标量存储器地址总线的数据大小位传送到所述存储器接口。
8.根据权利要求1所述的方法，其中，所述向量寄存器被映射到交换盒的输出数据缓冲器，或是被映射到从所述存储器端口到所述交换盒的读数据连接的虚拟向量寄存器。
9.一种处理器，包括：
10.根据权利要求9所述的处理器，所述序列发生器还用于从所述序列发生器的任务缓冲器加载线程块索引信息，并执行内核程序中的标量指令，以：
11.根据权利要求10所述的处理器，所述线程块大小信息包括x轴上的线程大小的第一值x、y轴上的线程大小的第二值y、以及x乘以y和z的乘法结果xyz，其中z是z轴上的线程大小，并且当所述存储器单元是字可寻址存储器单元时，所述寻址参数包括base、stride_y、以及stride_z，或者当所述存储器单元是字节可寻址存储器单元时，所述寻址参数包括base、stride_y、stride_z、以及ds，其中base是所述内核程序要处理的数据的起始存储器地址，对于所述字可寻址存储器单元，stride_y和stride_z分别是y轴和z轴上的存储器地址的间隙，或者对于所述字节可寻址存...
【专利技术属性】
技术研发人员：永田敏雄，李原，朱建斌，
申请(专利权)人：珠海市芯动力科技有限公司，
类型：发明
国别省市：
还没有人留言评论。发表了对其他浏览者有用的留言会获得科技券。