CN111712793B - 线程处理方法和图形处理器 - Google Patents

Abstract

一种应用于图形处理器的方法，该方法包括如下的步骤：第一线程处理器获取需要处理的第一待处理数据，确定第一待处理数据满足第一分支语句，将计数器中的数值加一步长。第一线程处理器根据计数器的数值确定M*N个线程中需运行第一分支语句的线程的数量。第一线程处理器在确认数量大于阈值的情况下，执行线程同步以及线程数据重映射，在有在第一分支的线程的数量较多的情况下，才使用线程数据重映射，可节约时间和运算资源。

Description

线程处理方法和图形处理器

技术领域

本发明实施例涉及数据处理领域，尤其涉及一种线程处理方法和图形处理器。

背景技术

图形处理器(graphics processing unit，GPU)在接收到由内核代码编译成的控制指令时，会创建一个任务，并针对该任务创建大量的线程进行并行处理。举例而言，GPU根据内核代码创建一个工作组(workgroup)，每个工作组包括多个线程束，一个线程束包括64个线程，在开放运算语言(Open Computing Language，OpenCL)中线程束称之为Wavefront(简称wave)，在计算统一设备架构(Compute Unified Device Architecture，CUDA)中线程束称之为warp。

一个工作组的多个线程束被分配至GPU的一个流多处理器，一个流多处理器包括多个流处理器，每个流处理器可以运行一个线程，流多处理器运行多个线程束。

具体而言，一个线程束中的线程绑定在一起在一个流多处理器上运行，每一时刻都执行统一的指令，工作组的多个线程束中，部分线程束可处于活跃状态，部分线程束可处于等待状态，当处于活跃状态的线程束在该流多处理器运行完毕，该流多处理器马上执行处于等待状态的线程束。举例而言，假设一个流多处理器包括64个流处理器，一个工作组包括4个线程束，一个线程束包括32个线程，则该流多处理器同一时刻可以运行两个线程束，这两个线程束处于活跃状态，另外两个线程束处于等待状态，当处于活跃状态的任一线程束运行完毕，流多处理器马上运行处于等待状态的线程束。

当同一线程束中的线程遇到分支且判断条件不唯一的时候，由于指令的统一性，该线程束需要串行执行其成员线程对应的分支，这被称为分支分歧问题。

举例而言，请参见以下的内核代码：

一个工作组的每个线程都执行该内核代码，其中，A[tid]为待处理数据，tid为线程标识，在每个线程均具有唯一的一个tid，在该内核代码中，假设针对多个线程而言，A[tid]＜5发生的概率小于A[tid]≥5发生的概率。但A[tid]＜5一旦发生，则要执行代码A(执行第一分支语句)，反之则执行代码B(即执行第二分支语句)。

在同一线程束中，判断到A[tid]＜5的线程需要执行第一分支语句，判断到A[tid]≥5的线程执行第二分支语句(或先执行第二分支语句然后执行第一分支语句)，流多处理器针对该线程束需串行执行第一分支语句和第二分支语句，从而降低了并行度以及执行的效率。

线程数据重映射(Thread-Data Remapping，TDR)是现有的解决分支分歧的软件技术。TDR通过改变分配给所有线程的待处理数据的排列，使得相同判断条件的待处理数据分配到相同的线程束中，同一线程束中的线程得到的判断条件一致，从而消除分支分歧。

例如，通过TDR重新调配待处理数据，可将需执行第一分支语句的线程均设置在同一个线程束中，并使得需执行第二分支语句的线程设置在其他线程束中，因此，流多处理器无需在同一个线程束中串行地执行第一分支语句和第二分支语句。

现有技术为解决分支分歧的问题，在线程执行分支语句之前先进行TDR，但是，若在工作组中需执行第一分支语句的线程的数量较少，则进行TDR是毫无意义的。

举例而言，假设整个工作组只有1个需执行第一分支语句的线程，其他线程均执行第二分支语句，那么，无论将需执行第一分支语句的线程的待处理数据集中分配到哪一个线程束，流多处理器针对该线程束均需要串行执行第一分支语句和第二分支语句。

因此，在工作组中需执行第一分支语句的线程的数量较少的情况下，进行TDR并没有产生作用，反而白白浪费了进行TDR的时间和运算资源。

发明内容

本发明实施例提供了一种线程处理方法和图形处理器，只有在需执行第一分支语句的线程的数量多于阈值的情况下，才使用线程数据重映射，可节约时间和运算资源。

第一方面，本申请提供一种线程处理方法，该方法应用于图形处理器，具体的，图形处理器用于处理M个线程束，每个线程束包括N个线程，图形处理器还包括至少一个线程束处理器，至少一个线程束处理器中的第一线程束处理器包括N的整数倍个线程处理器，第一线程束处理器包括第一线程处理器，第一线程处理器用以运行N个线程中的一个以处理满足第一分支语句的判断条件或满足第二分支语句的判断条件的待处理数据，图形处理器中设置有计数器，该方法包括如下的步骤：第一线程处理器获取需要处理的第一待处理数据，确定第一待处理数据满足第一分支语句，将计数器中的数值加一步长。第一线程处理器根据计数器的数值确定M*N个线程中需运行第一分支语句的线程的数量。第一线程处理器在确认数量大于阈值的情况下，执行线程同步以及线程数据重映射。

其中，第一线程处理器通过判断运行第一分支语句的线程的数量，在确认数量大于阈值的情况下，才执行线程同步以及线程数据重映射，可避免在该数量小于或等于阈值的情况下也执行线程同步以及线程数据重映射，可提高GPU的运行速度和运行效率。其中，步长可以是任意正整数，举例而言，可设置为1。

在第一方面的第一种可能的实现方式中，图形处理器还设置有标志位，标志位的值设置为第一标志值，第一标志值用于指示不执行重映射，该方法还包括如下步骤：第一线程处理器在确定数量大于阈值的情况之前，读取标志位。且，第一线程处理器在确定数量大于阈值的情况之后及执行线程同步之前，将第一标志值设置为第二标志值，第二标志值用于指示需要执行重映射。

第一个发现计数器的数值超过阈值的线程处理器将标志位设置为第二标志值，其他线程处理运行其他线程时，读取到第二标志值被设置就直接执行重映射，而不再去判断计数器的数值是否超过阈值，避免得出不一样的结论，采用这个设计后，只要M*N个线程中的一个线程作出执行线程同步的决定，其他线程必定跟随，规避了个别线程单独执行线程同步的问题。

根据第一方面或第一方面的第一种可能的实现方式，在第二种可能的实现方式中，第一线程处理器在执行线程同步之后且执行线程数据重映射之前，该方法还包括以下步骤：第一线程处理器将计数器中的数值清零。

在本线程已经确定需要执行线程数据重映射的情况下，将计数器清零，使得其他线程仅需根据标志位设置为第二标志值即可确认需要执行线程数据重映射，而无需再根据计数器的数值判断是否需要执行线程数据重映射。

根据第一方面、第一方面的第一种可能的实现方式以及第一方面的第二种可能的实现方式中的任一者，在第三种可能的实现方式中，第一线程束处理器包括第二线程处理器，第二线程处理器用以在运行N个线程中的一个以处理满足第一分支语句或满足第二分支语句的待处理数据，该方法还包括以下步骤：第二线程处理器读取标志位，在确认标志位的值为第二标志值时，执行线程同步以及线程数据重映射。第二线程处理器在确认标志位的值为第一标志值时，根据计数器的数值确定M*N个线程中需运行第一分支语句的线程的数量，在确认数量大于阈值的情况下，执行线程同步以及线程数据重映射。

其他线程可根据标志位直接判断是否需进行重映射，可规避个别线程单独执行线程同步的问题。

根据第一方面的第三种可能的实现方式，在第四种可能的实现方式中，第一线程处理器用以运行N个线程中的第一线程以处理满足第一分支语句的判断条件的的待处理数据，第二线程处理器用以运行N个线程中的第二线程以处理满足第二分支语句的判断条件的待处理数据，图形处理器还设置有一维数组、第一变量以及第二变量，其中一维数组的长度是M*N，第一变量的初始值是0，第二变量的初始值是M*N-1，第一线程处理器执行线程数据重映射，包括：第一线程处理器读取第二变量的数值，并将第二线程的线程标识写入一维数组中以第二变量的数值作为下标的位置，将第二变量的数值减一，并执行线程同步。第二线程处理器读取第一变量的数值，并将第二线程的线程标识写入一维数组中以第一变量的数值作为下标的位置，将第一变量的数值加一，并执行线程同步。第一线程处理器在线程同步结束后，读取一维数组中以第一线程的线程标识作为下标的位置上的数值，并将读取的数值作为线程数据重映射产生的第一线程的更新的线程标识。第二线程处理器在线程同步结束后，读取一维数组中以第二线程的线程标识作为下标的位置上的数值，并将读取的数值作为线程数据重映射产生的第二线程的更新的线程标识。

该重映射的实现方式仅涉及一维数组的指针变换，避免了在重映射过程中直接调用待处理数据，可有效提高运算速度。

根第一方面的第四种可能的实现方式，在第五种可能的实现方式中，第一线程处理器在执行线程同步之后且执行线程数据重映射之前，该方法还包括：第一线程处理器以第一线程的线程标识作为索引将第一待处理数据记录在索引表，其中第一线程的线程标识与第一待处理数据具有一一对应关系，索引表记录有M*N个线程的线程标识与待处理数据之间的一一对应关系。进一步地，第一线程处理器在执行线程数据重映射之后，该方法还包括以下步骤：第一线程处理器以执行线程数据重映射后产生的第一线程的更新的线程标识作为索引在索引表中读取与第一线程的更新的线程标识对应的第三待处理数据。第一线程处理器在第三待处理数据满足第一分支语句的判断条件时，执行第一分支语句，第一线程处理器在第三待处理数据满足第二分支语句的判断条件时，执行第二分支语句。

所有线程以自身的线程标识通过将待处理数据保存在索引表，并在获得重映射分配的更新的线程标识之后，以更新的线程标识作为索引从索引表获取更新的线性标识对应的待处理数据，并根据该待处理数据判断是执行第一分支语句抑或是第二分支语句，可实现线程之间的数据交换，并保证内核代码的正常运行。

在第一方面的一种可能的实现方式中，阈值为1。

在第一方面的一种可能的实现方式中，阈值为大于等于2且小于等于5的正整数。

在第一方面的一种可能的实现方式中，第一线程处理器执行第一分支语句的概率小于第一线程处理器执行第二分支语句的概率。

由于执行第一分支语句的概率较小，因此通过重映射将需执行第一分支语句的线程尽量集中在一个或多个线程束中，使得尽量多的线程束避免串行地执行第一分支语句和第二分支语句。

在第一方面的一种可能的实现方式中，计数器、索引表设置在图形处理器的共享存储器。

在第一方面的一种可能的实现方式中，待处理数据设置在图形处理器的全局存储器。

在第一方面的一种可能的实现方式中，第一线程处理器通过对计数器中的数值进行原子加一以实现加一步长的操作，第一线程处理器通过对计数器中的数值进行原子减一以实现减一步长的操作。

在第一方面的一种可能的实现方式中，待处理数据经由与图形处理器连接的中央处理器发送到图形处理器的全局存储器。

第二方面，本申请提供一种线程处理方法，该方法应用于图形处理器，具体的，图形处理器用于处理M个线程束，每个线程束包括N个线程，的图形处理器还包括至少一个线程束处理器，至少一个线程束处理器中的第一线程束处理器包括N的整数倍个线程处理器，第一线程束处理器包括第一线程处理器，第一线程处理器运行循环语句，用以在一个循环中运行N个线程中的一个以处理满足第一分支语句的判断条件或满足第二分支语句的判断条件的待处理数据，图形处理器中设置有计数器，该方法包括以下步骤：第一线程处理器在第一循环中获取需要处理的第一待处理数据，确定第一待处理数据满足第一分支语句，将计数器中的数值加一步长。第一线程处理器根据计数器的数值确定M*N个线程中需运行第一分支语句的线程的数量。第一线程处理器在确认数量大于阈值的情况下，执行线程同步并将计数器中的数值清零。第一线程处理器执行线程数据重映射。

综上，第一线程处理器在确认数量大于阈值的情况下，执行线程同步并将计数器中的数值清零，并执行线程数据重映射，使得第一线程处理器在下一循环读取到的计数器的数值不会受到第一循环的影响，并且，由于在确认数量大于阈值的情况下，才执行线程同步以及线程数据重映射，因此可避免在该数量小于或等于阈值的情况下也执行线程同步以及线程数据重映射，可在每一循环中动态地判断是否有执行线程数据重映射的需要，可避免进行无用的线程数据重映射，可提高GPU的运行速度和运行效率。

在第二方面的第一种可能的实现方式中，该方法还包括以下步骤：第一线程处理器获取在第二循环中需要处理的第二待处理数据，确定第二待处理数据满足第二分支语句的判断条件，将计数器中的数值减一步长。

将计数器中的数值减一步长，可与加一步长的操作抵消，可避免对其他线程的判断造成干扰，并且可避免对下一循环进行干扰。

结合第二方面或第二方面的第一种可能的实现方式，在第二种可能的实现方式中，图形处理器还设置有标志位，标志位的值设置为第一标志值，第一标志值用于指示不执行重映射，该方法还包括以下步骤：第一线程处理器在确定数量大于阈值的情况之前，读取标志位。且，第一线程处理器在确定数量大于阈值的情况之后及执行线程同步之前，将第一标志值设置为第二标志值，第二标志值用于指示需要执行重映射。

第一个发现计数器的数值超过阈值的线程处理器将标志位设置为第二标志值，其他线程处理运行其他线程时，读取到第二标志值被设置就直接执行重映射，而不再去判断计数器的数值是否超过阈值，避免得出不一样的结论，采用这个设计后，只要M*N个线程中的一个线程作出执行线程同步的决定，其他线程必定跟随，规避了个别线程单独执行线程同步的问题。

结合第二方面的第二种可能的实现方式，在第三种可能的实现方式中，第一线程束处理器包括第二线程处理器，第二线程处理器用以在运行N个线程中的一个以处理满足第一分支语句或满足第二分支语句的待处理数据，方法还包括：第二线程处理器读取标志位，在确认标志位的值为第二标志值时，执行线程同步以及线程数据重映射；第二线程处理器在确认标志位的值为第一标志值时，根据计数器的数值确定M*N个线程中需运行第一分支语句的线程的数量，在确认数量大于阈值的情况下，执行线程同步以及线程数据重映射。

由于其他线程可根据标志位直接判断是否需进行重映射，可规避个别线程单独执行线程同步的问题。

根据第二方面的第三种可能的实现方式，在第四种可能的实现方式中，第一线程处理器用以运行N个线程中的第一线程以处理满足第一分支语句的判断条件的的待处理数据，第二线程处理器用以运行N个线程中的第二线程以处理满足第二分支语句的判断条件的待处理数据，图形处理器还设置有一维数组、第一变量以及第二变量，其中一维数组的长度是M*N，第一变量的初始值是0，第二变量的初始值是M*N-1，第一线程处理器执行线程数据重映射，包括：第一线程处理器读取第二变量的数值，并将第二线程的线程标识写入一维数组中以第二变量的数值作为下标的位置，将第二变量的数值减一，并执行线程同步。第二线程处理器读取第一变量的数值，并将第二线程的线程标识写入一维数组中以第一变量的数值作为下标的位置，将第一变量的数值加一，并执行线程同步。第一线程处理器在线程同步结束后，读取一维数组中以第一线程的线程标识作为下标的位置上的数值，并将读取的数值作为线程数据重映射产生的第一线程的更新的线程标识。第二线程处理器在线程同步结束后，读取一维数组中以第二线程的线程标识作为下标的位置上的数值，并将读取的数值作为线程数据重映射产生的第二线程的更新的线程标识。

该重映射的实现方式仅涉及一维数组的指针变换，避免了在重映射过程中直接调用待处理数据，可有效提高运算速度。

结合第二方面的四种可能的实现方式中的任一者，在第五种可能的实现方式中，第一线程处理器运行第一线程，第一线程处理器在执行线程同步之后且执行线程数据重映射之前，该方法还包括以下步骤：第一线程处理器以第一线程的线程标识作为索引将第一待处理数据和第一循环变量记录在索引表，其中，第一线程的线程标识与第一待处理数据具有一一对应关系。第一线程处理器在执行线程数据重映射之后，该方法还包括以下步骤：第一线程处理器以执行线程数据重映射后产生的第一线程的更新的线程标识作为索引在索引表中读取与第一线程的更新的线程标识对应的第三待处理数据。第一线程处理器在第三待处理数据满足第一分支语句的判断条件时，执行第一分支语句，第一线程处理器在第三待处理数据满足第二分支语句的判断条件时，执行第二分支语句。

所有线程以自身的线程标识通过将待处理数据保存在索引表，并在获得重映射分配的更新的线程标识之后，以更新的线程标识作为索引从索引表获取更新的线性标识对应的待处理数据，并根据该待处理数据判断是执行第一分支语句抑或是第二分支语句，可在线程之间交换待处理数据，并保证内核代码的正常运行。

结合第二方面的第五种可能的实现方式，在第六种可能的实现方式中，图形处理器中还记录有每个线程的循环变量，循环变量用于指示线程当前所在的循环的序号，索引表中记录有第一线程的循环变量与第一线程的线程标识、第一线程在循环变量所指示的循环中的待处理数据的对应关系，第一线程处理器在执行线程数据重映射之后，该方法还包括：第一线程处理器以执行线程数据重映射后产生的第一线程的更新的线程标识作为索引在索引表中读取与第一线程的更新的线程标识对应的循环变量。第一线程处理器在执行第一分支语句或第二分支语句之后，将第一线程的更新的线程标识对应的循环变量加一以获取更新的循环变量，且在更新的循环变量不符合循环语句规定的循环条件时，结束第一线程，在更新的循环变量符合循环语句规定的循环条件时，运行第一线程的第二循环。

综上，可实现循环变量在线程之间交换，使得同一线程束可存在处于不同循环的线程。

在第二方面的一种可能的实现方式中，阈值为1。

在第二方面的一种可能的实现方式中，阈值为大于等于2且小于等于5的正整数。

在第二方面的一种可能的实现方式中，第一线程处理器执行第一分支语句的概率小于第一线程处理器执行第二分支语句的概率。

由于执行第一分支语句的概率较小，因此通过重映射将需执行第一分支语句的线程尽量集中在一个或多个线程束中，使得尽量多的线程束避免串行地执行第一分支语句和第二分支语句。

在第二方面的一种可能的实现方式中，计数器、索引表设置在图形处理器的共享存储器。

在第二方面的一种可能的实现方式中，待处理数据设置在图形处理器的全局存储器。

在第二方面的一种可能的实现方式中，第一线程处理器通过对计数器中的数值进行原子加一以实现加一步长的操作，第一线程处理器通过对计数器中的数值进行原子减一以实现减一步长的操作。

在第二方面的一种可能的实现方式中，待处理数据经由与图形处理器连接的中央处理器发送到图形处理器的全局存储器。

第三方面，本申请提供一种图形处理器，图形处理器用于处理M个线程束，每个线程束包括N个线程，的图形处理器还包括至少一个线程束处理器，至少一个线程束处理器中的第一线程束处理器包括N的整数倍个线程处理器，第一线程束处理器包括第一线程处理器，第一线程处理器用以运行N个线程中的一个以处理满足第一分支语句的判断条件或满足第二分支语句的判断条件的待处理数据，图形处理器中设置有计数器，其中，第一线程处理器，用于获取需要处理的第一待处理数据，确定第一待处理数据满足第一分支语句，将计数器中的数值加一步长。第一线程处理器，用于根据计数器的数值确定M*N个线程中需运行第一分支语句的线程的数量。第一线程处理器，用于在确认数量大于阈值的情况下，执行线程同步以及线程数据重映射。

第三方面或第三方面任意一种实现方式是第一方面或第一方面任意一种实现方式对应的装置实现，第一方面或第一方面任意一种实现方式中的描述适用于第三方面或第三方面任意一种实现方式，在此不再赘述。

第四方面，本申请提供一种图形处理器，图形处理器用于处理M个线程束，每个线程束包括N个线程，的图形处理器还包括至少一个线程束处理器，至少一个线程束处理器中的第一线程束处理器包括N的整数倍个线程处理器，第一线程束处理器包括第一线程处理器，第一线程处理器运行循环语句，用以在一个循环中运行N个线程中的一个以处理满足第一分支语句的判断条件或满足第二分支语句的判断条件的待处理数据，图形处理器中设置有计数器，其中，第一线程处理器，用于在第一循环中获取需要处理的第一待处理数据，确定第一待处理数据满足第一分支语句，将计数器中的数值加一步长。第一线程处理器，用于根据计数器的数值确定M*N个线程中需运行第一分支语句的线程的数量。第一线程处理器，用于在确认数量大于阈值的情况下，执行线程同步并将计数器中的数值清零。第一线程处理器，用于执行线程数据重映射。

第四方面或第四方面任意一种实现方式是第二方面或第二方面任意一种实现方式对应的装置实现，第二方面或第二方面任意一种实现方式中的描述适用于第四方面或第四方面任意一种实现方式，在此不再赘述。

第五方面，本申请一种线程处理方法，该方法应用于图形处理器，具体的，图形处理器用于处理M个线程束，每个线程束包括N个线程，该方法包括以下步骤：检测M*N个线程中需运行第一分支语句的线程的数量。在确认数量大于阈值的情况下，对M*N个线程进行线程数据重映射。

其中，在确认数量大于阈值的情况下，才执行线程同步以及线程数据重映射，可避免在该数量小于或等于阈值的情况下也执行线程同步以及线程数据重映射，可提高GPU的运行速度和运行效率。

在第五方面的第一种可能的实现方式中，M*N个线程分别设置有线程标识和待处理数据，待处理数据与线程标识具有一一对应关系，对M*N个线程进行线程数据重映射，包括：获取M*N个线程的待处理数据。在M*N个线程的任一线程的待处理数据满足第一分支语句的判断条件的情况下，将满足第一分支语句的判断条件的待处理数据依次映射至M*N个线程中相邻的部分线程。在M*N个线程的任一线程的待处理数据满足第二分支语句的判断条件的情况下，将满足第二分支语句的判断条件的待处理数据依次映射至M*N个线程中相邻的部分线程。将M*N个线程的线程标识分别更新为所映射的待处理数据对应的线程标识。

因此通过重映射将需执行第一分支语句的线程尽量集中在一个或多个线程束中，通过重映射将需执行第二分支语句的线程尽量集中在一个或多个线程束中，使得尽量多的线程束避免串行地执行第一分支语句和第二分支语句。

结合第五方面和第五方面的第一种可能的实现方式，在第二种可能的实现方式中，对M*N个线程进行线程数据重映射之前，该方法还包括：控制M*N个线程以自身的线程标识作为索引记录自身的待处理数据至索引表。对M*N个线程进行线程数据重映射之后，该方法还包括：控制M*N个线程以更新的线程标识作为索引在索引表获取与更新的线程标识对应的待处理数据。

所有线程以自身的线程标识通过将待处理数据保存在索引表，并在获得重映射分配的更新的线程标识之后，以更新的线程标识作为索引从索引表获取更新的线性标识对应的待处理数据，并根据该待处理数据判断是执行第一分支语句抑或是第二分支语句，可实现线程之间的数据交换，并保证内核代码的正常运行。

结合第五方面的第一种可能的实现方式，在第三种可能的实现方式中，M*N个线程分别运行循环语句，检测M*N个线程中需运行第一分支语句的线程的数量，包括：获取M*N个线程需要处理的待处理数据，在M*N个线程中的任一线程的待处理数据满足第一分支的判断条件的情况下，将计数器的数值加一。

结合第五方面的第三种可能的实现方式，在第四种可能的实现方式中，对M*N个线程进行线程数据重映射之前，控制M*N个线程以自身的线程标识作为索引记录自身的待处理数据和循环变量至索引表。且在对M*N个线程进行线程数据重映射之后，控制M*N个线程以更新的线程标识作为索引在索引表获取与更新的线程标识对应的待处理数据和循环变量。

通过将循环变量存储在索引表，可以实现循环变量在线程之间的交换。

结合第五方面的第三种可能的实现方式，在第五种可能的实现方式中，对M*N个线程进行线程数据重映射，还包括以下子步骤：控制确认计算变量大于阈值的线程以自身的线程标识作为索引记录自身的待处理数据和循环变量到索引表，将标志位设置为第二标志值，执行线程同步。控制其他线程在检测到标志位为第二标志值的情况下，以自身的线程标识作为索引记录自身的待处理数据和循环变量到索引表，执行线程同步。控制M*N个线程执行线程数据重映射以获取更新的线程标识。

发现计数器的数值超过阈值的线程处理器将标志位设置为第二标志值，其他线程处理运行其他线程时，读取到第二标志值被设置就直接执行重映射，而不再去判断计数器的数值是否超过阈值，避免得出不一样的结论，采用这个设计后，只要M*N个线程中的一个线程作出执行线程同步的决定，其他线程必定跟随，规避了个别线程单独执行线程同步的问题。

在执行线程同步的步骤之后，结合第五方面的第四或第五种可能的实现方式，在第六种可能的实现方式中，该方法还包括：将标志位设置为第一标志值，并将计数器的数值设置为0。

将计数器中的数值清零，线程在下一循环读取到的计数器的数值不会受到本循环的影响。

第六方面，本申请提供一种图形处理器，该图形处理器包括线程束处理器，该线程束处理器用于处理M个线程束，每个线程束包括N个线程，该线程束处理器，用于检测M*N个线程中需运行第一分支语句的线程的数量。该线程束处理器，用于在确认数量大于阈值的情况下，对M*N个线程进行线程数据重映射。

第六方面或第六方面任意一种实现方式是第五方面或第五方面任意一种实现方式对应的装置实现，第一方面或第一方面任意一种实现方式中的描述适用于第三方面或第三方面任意一种实现方式，在此不再赘述。

附图说明

图1是根据本发明实施例的图形处理器和中央处理器的连接关系示意图；

图2是根据本发明实施例的工作组投放至第一线程束处理器的示意图；

图3是根据本发明实施例的线程数据重映射的流程图；

图4是根据本发明实施例的线程数据重映射的数据流向示意图；

图5是根据本发明实施例的线程数据重映射的前后时间开销示意图；

图6是根据本发明实施例的线程数据重映射的另一数据流向图；

图7是是根据本发明实施例的线程数据重映射的另一前后时间开销示意图；

图8是根据本发明实施例的线程处理方法的流程图；

图9是根据本发明实施例的线程处理方法的另一流程图；

图10是根据本发明实施例的线程数据重映射的前后时间开销示意图；

图11是根据本发明实施例的线程数据重映射的另一前后时间开销示意图；

图12是根据本发明实施例的线程数据重映射的原理图；

图13是根据本发明实施例的线程处理方法的另一流程图；

图14是根据本发明实施例的线程处理方法的另一流程图；

图15是根据本发明实施例的异构系统的装置结构示意图。

具体实施方式

为了方便理解本发明的各实施例，下面先对本发明各实施例涉及到的一些技术术语进行介绍，后文的各实施例可以参考下面的技术术语介绍。

1、图形处理器

图形处理器(Graphics Processing Unit，GPU)是一种在个人电脑、工作站、游戏机和一些移动设备(如平板电脑、智能手机等)等设备上进行图像运算工作的微处理器。GPU的用途是将计算机系统所需要的显示信息进行转换驱动，并向显示器提供行扫描信号，控制显示器的正确显示。

GPU作为一个大规模并行计算元件，因其日益强大的计算能力，已经被广泛应用于通用计算中。不同领域中的大量程序都用GPU进行加速，如传统计算密集型的科学计算、文件系统、网络系统、数据库系统和云计算等。

如图1所示，GPU包括全局存储器、调度器以及多个线程束处理器，其中图1是根据本发明实施例的图形处理器和中央处理器的连接关系示意图。

全局存储器存储有接收自CPU的主机端代码(host code)、内核代码(kernelcode)以及待处理数据。

调度器用于根据主机代码设置工作组(workgroup)，该工作组包括M个线程束，每个线程束包括N个线程，调度器选择一个空闲的线程束处理器，并将M个线程束投放到该线程束处理器，该线程束处理器用于处理该工作组的M*N个线程，其中M≥1，N≥1。

每个线程束包括预定数量的线程，一个线程束处理器可同一时间运行该预定数量的线程。在不同的产品中，所述线程束处理器可能有别的名称。比如，美国超微半导体公司(Advanced Micro Devices，AMD)会将该线程束处理器称作计算单元(Computing Unit，CU)，英伟达公司(NVIDIA Corporation，NVIDIA)会将线程束处理器称作流多处理器(Stream Multiprocessor，SM)。

每个线程束处理器包括共享存储器、N个线程处理器以及N个私有存储器，每个线程处理器用于在同一时刻运行一个线程，私有存储器用于存储线程运行过程中涉及的待处理数据以及过程数据，其中过程数据包括计算结果、循环变量、以及计算过程中涉及的中间值等。

其中，一个线程束处理器包括的线程处理器的数量为N的整数倍。

每个线程处理器分别对应设置有一个私有存储器，线程处理器只能访问与自身对应的私有存储器，其中私有存储器可例如为寄存器组或内存，同一个线程束处理器内的线程处理器均可访问同一个线程束处理器内的共享存储器，但不能访问不同线程束处理器内的共享存储器。举例而言，在图1所示的图形处理器20中，第一线程处理器211、第二线程处理器212、第三线程处理器213以及第三第四线程处理器214均可访问共享存储器210，但不能访问共享存储器220。第一线程处理器211可访问第一线程处理器211的私有存储器215，但不能访问第二线程处理器212的私有存储器216。

进一步，每个线程束处理器中的线程处理器均可访问全局存储器，举例而言，在图1中，第一线程处理器211和第五线程处理器221均可访问全局存储器24。

2、GPU程序

一个GPU程序可以分为两个部分：如上所述的主机端代码和内核代码。在CPU上运行的代码编辑器可编辑GPU程序，设置待处理数据，GPU程序可通过在CPU上运行的编译器编译成二进制格式的可GPU可执行代码。CPU发送待处理数据和编译好的GPU程序至GPU的全局存储器。

GPU的调度器读取全局存储器中的主机端代码对内核代码上下文进行初始化以创建工作组。GPU的调度器将待处理数据分配至工作组中的线程，并通知工作组中的每一线程执行内核代码以处理各自分配到的待处理数据。

在另外一些示例中，待处理数据可由GPU通过执行内核代码产生。

在主机端代码设置好的待处理数据的情况下，内核代码上下文的初始化可将主机端代码设置好的待处理数据以数组的方式分配至线程，其中该数组以线程标识为下标，一个GPU的内核代码描述的是一个线程的行为，对应线程标识的线程可读取该数组中的待处理数据，并根据待处理数据执行该行为。在另一些示例中，在主机端代码没有设置好的待处理数据的情况下，内核代码描述的是一个线程的行为，该线程可根据内核代码产生待处理数据，并根据待处理数据执行该行为。

3、线程束

线程束为由工作组成的集合体。GPU作为协处理器，在从CPU接收到内核代码启动调用时，会通过调度器创建大量的线程。这些线程会分层的组织到一起。举例而言，一个线程束可包括32个线程，即N＝32，又或者，一个线程束可包括64个线程或其他数量的线程，即N＝64。

线程束中的线程在线程束处理器上执行时，绑定在一起执行，每一时刻都执行统一的指令，但处理不同的待处理数据。一些线程束组成一个工作组。内核代码对应分配有一个或多个工作组，每一工作组包括M个线程束，每一线性束包括N个线程，其中，M为CPU执行主机端代码进行内核代码上下文的初始化后控制GPU创建的工作组所包括的线程束的数量，每个线程束所包括的线程的数量N为所创建的该工作组的任意一个线程束所包括的线程的数量。

M的具体数值可在主机端代码中设置，或由GPU预先设定为一个固定数值，在一些示例中，M的取值范围可以为4≤M≤32，在另一些示例中，M的取值范围可以为1≤M。

N与GPU的一个线程束处理器所包括的线程处理器的数量E之间通常具有倍数关系，E＝n*N，n为大于或等于1的正整数。

可结合图2进行参考，图2是根据本发明实施例的工作组投放至第一线程束处理器的示意图。在图2的示例中，调度器23创建的工作组包括3个线程束11、12、13，即M＝3，线程束11包括线程1至4，线程束12包括线程5至8，线程束13包括线程9至13，即N＝4。

并且，在图2中，第一线程束处理器21仅设置有4个第一线程处理器211至214，而第一线程束处理器21在同一时刻仅可运行一个线程束，例如线程束11，当线程束处理器运行线程束11时，第一线程处理器211运行线程1，第二线程处理器212运行线程2，第三线程处理器213运行线程3，第四线程处理器214运行线程4。在线程11的所有线程运行完毕之后，线程束处理器21运行线程束12或线程束13，其中，若线程束处理器21运行线程束12，第一线程处理器211运行线程5，第二线程处理器212运行线程6，第三线程处理器213运行线程7，第四线程处理器214运行线程8。之后，线程束处理器21可运行线程束13，在线程束处理器21运行线程束13时，第一线程处理器211运行线程9，第二线程处理器212运行线程10，第三线程处理器213运行线程11，第四线程处理器214运行线程12。

其中，第一线程束处理器21运行工作组的线程束的先后顺序由调度器23决定，调度器23会让第一线程束处理器21优先运行没有发生读延迟的线程束，其中读延迟是指线程处理器在从全局存储器24读取待处理数据时产生的延迟。

举例而言，在图2中，若调度器23首先控制第一线程束处理器21运行线程束11，而在此过程中，第一线程处理器211从全局存储器24读取待处理数据A[0]至共享存储器210时产生延迟，此时调度器23可通知第一线程束处理器21停止运行线程束11，转而运行线程束12或线程束13，从而避免等待该延迟。

值得注意的是，第一线程束处理器21运行线程束12或线程束13的过程中，调度器23可通知存储器控制器(memory controller，图未示出)将待处理数据A[0]从共享存储器210读取至全局存储器24中。因此，在第一线程束处理器21运行线程束12或线程束13完毕之后，可在调度器23的控制下无需等待待处理数据A[0]，继续运行线程束11，从而实现掩盖延迟的功能。

值得注意的是，在实际应用中，线程束处理器21可包括更多数量的线程处理器，因此，线程束处理器21在同一时刻可运行不止1个线程束，而在本发明实施例中，为了便于说明，将线程束处理器21设置为包括4个线程处理器，使得线程束处理器21在同一时刻运行一个线程束。

4、分支分歧

分支分歧是GPU计算中常见的导致性能损耗的因素。

每个线程的线程标识(thread identification，TID)以及读取到的待处理数据都不尽相同，因此遭遇分支的时候会得出不同的判断条件。当同一个线程束中的线程需要执行不同的分支时，由于执行指令的统一性，该线程束将串行执行成员线程需要执行的所有分支，这被称为分支分歧问题。每条分支都是全部线程一起执行，但是无关的线程运行的结果会被舍弃，这降低了并行度以及执行的效率。例如，单层的分支可以将效率降至50％，而循环语句中的多层嵌套分支更可造成指数级增长的减速。

5、循环语句

举例而言，循环语句可包括for语句：

for语句的一般形式如下：

“for(循环变量的初始值；循环变量的范围；循环变量加1)”

循环变量可例如为i，初始值设置为0，循环变量的范围可限定为小于1000次，即i＜1000，并且，循环变量在每个循环过程进行加一操作i++，即i＝i+1，使得循环变量在每个循环执行时进行自加1的操作。

其中，for(i＝0；i＜1000；i++)语句1，表示需要执行语句1一千次。

在另外一些示例中，循环语句也可包括while语句，其中while语句并没有限定循环变量，即其循环的次数没有限定，while语句的一般形式如下：

“while(待处理数据是否满足判断条件)”

当待处理数据满足判断条件时，继续循环，当待处理不满足判断条件时，则退出循环。

值得注意的是循环语句还可以包括其他语句，如select语句，于此不作赘述。

6、分支语句

举例而言，分支语句可包括if语句：

if语句是选择语句，if语句用于实现两个分支的选择。if语句的一般形式如下：

“if(条件判断式)第一分支语句

[else第二分支语句]”

其中，方括号内的部分(即else子句)为可选的，即可以有，也可以没有。条件执行语句1和条件执行语句2可以是一个简单的语句，也可以是一个复合语句，还可以是另一个if语句(即在一个if语句中又包括另一个或多个内嵌的if语句)。条件判断式也可以称之为表达式(expression)，条件执行语句也可以称之为语句(statement)。

在本发明实施例中，条件判断式设置为使得需执行第一分支语句的次数小于需执行第二分支语句的次数。

举例而言，条件判断式可设置为

其中，temp＝rand()*1000。

rand()函数为伪随机函数，可用于产生0至1之间的任意数值，temp作为待处理数据分配给工作组的每个线程，每个线程分到的temp不相同，在该内核代码中，temp＜5发生的概率小于temp＞＝5发生的概率。但temp＜5一旦发生，则要执行代码A从而执行第一分支语句，反之则执行代码B从而执行第二分支语句。

在if语句中又包括一个或多个if语句称为if语句的嵌套。在上述的(3)的形式属于if语句的嵌套，其一般形式如下：

在其他示例中，分支语句还包括switch语句，于此不作赘述。

7、线程同步

同一工作组的线程间进行数据交换时需要进行同步，而GPU提供了软件接口，举例而言线程可通过调用障碍函数barrier()实现线程同步，对于一个调用了障碍函数的线程来说，除非同一个工作组内其他线程都执行了障碍函数，否则该线程将被阻止执行障碍函数之后的内核代码，且在线程处理器执行了障碍函数之后，线程处理器为该线程设置停顿点(break point)，该停顿点记录了障碍函数的下一语句在内核代码中的位置，将该位置记录在该线程的私有存储器中，并暂停运行该线程。

当同一个工作组内其他线程都执行了障碍函数之后，线程处理器从私有存储器中读取停顿点，从而可执行障碍函数的下一语句，使得线程继续运行。

因此，针对执行了障碍函数的线程而言，线程同步结束的条件是一个工作组内M*N个线程都执行了障碍函数。

8、线程数据重映射

线程数据重映射是解决分支分歧的软件技术，其通过调整线程与待处理数据间的映射关系，使同一线程束中的线程得到的判断条件一致，从而消除分支分歧。

为了进一步清楚说明，请结合图3和图4一并参考，图3是根据本发明实施例的线程数据重映射的流程图，而图4是根据本发明实施例的线程数据重映射的数据流向示意图，在本发明实施例中，假设图2所示的3个线程束11、12、13由调度器23投放至第一线程束处理器21，因此图3所示的方法由第一线程束处理器21执行，线程数据重映射包括以下步骤：

步骤S1041：获取待处理数据。

举例而言，如图4所示，线程1至12的线程标识tid分别为0至12，待处理数据为数组A[tid]＝{13，6，0，1，2，7，8，9，10，11，3，12}中的数据，数组A[tid]由CPU30发送至GPU的全局存储器24，并被第一线程束处理器21读取至共享存储器210，线程处理器在运行线程时可根据该线程的线程标识tid从共享存储器210读取到待处理数据A[tid]，第一线程处理器211运行线程1时，可根据线程1的线程标识0，在共享存储器210的数组A[tid]读取待处理数据A[tid]＝13。

待处理数据与待处理数据所在线程的线程标识具有一一对应关系，例如待处理数据13对应线程标识0，待处理数据6对应线程标识1。

步骤S1042：判断待处理数据满足何种分支的判断条件。

举例而言，判断条件为分支语句的判断条件，如上文的代码段中的A[tid]＜5为第一分支语句的判断条件，A[tid]≥5为第二分支语句的判断条件，即在满足该条件时，执行第一分支语句，在不满足该条件时，执行第二分支语句。

结合图4，A[2]＝0，A[3]＝1，A[4]＝2，A[10]＝3，均小于5，因此，第一线程处理器211运行线程3、线程5和线程9是需执行第一分支语句，第四线程处理器214运行线程4时需执行第二分支语句。

步骤S1043：将满足第二分支语句的判断条件的待处理数据依次映射至M*N个线程中相邻的部分线程。

举例而言，如图4所示，将满足第二分支语句的判断条件的待处理数据13，6，7，8，9，10，11，12分配至M个线程束中从第一个线程开始且相邻的多个线程1至8。

步骤S1044：将满足第二分支语句的判断条件的待处理数据依次映射至M*N个线程中相邻的其他部分线程。

举例而言，将满足第一分支语句的判断条件的待处理数据0，1，2，3分配至M个线程束中从最后一个线程开始且相邻的多个线程9至12。

步骤S1045：将M*N个线程的线程标识分别更新为所映射的待处理数据对应的线程标识。

对应地，在图4中，将线程1的线程标识0修改为待处理数据13的线程标识0，将线程2的线程标识1修改为待处理数据6的线程标识1，将线程3的线程标识2修改为待处理数据7的线程标识，将线程4的线程标识3修改为待处理数据8的线程标识6，将线程5的线程标识4修改为待处理数据9的线程标识7，将线程6的线程标识5修改为待处理数据10的线程标识8，将线程7的线程标识6修改为待处理数据11的线程标识9，将线程8的线程标识7修改为待处理数据12的线程标识11，将线程9的线程标识8修改为待处理数据3的线程标识10，将线程10的线程标识9修改为待处理数据2的线程标识4，将线程11的线程标识10修改为待处理数据1的线程标识3，将线程12的线程标识11修改为待处理数据0的线程标识2。

为了进一步清楚说明线程数据重映射，以下请结合图5，图5是根据本发明实施例的线程数据重映射的前后时间开销示意图。

图5的上半部分示出的是未做线程数据重映射时第一线程束处理器21分别运行线程束11至13所需的时间，于此假设线程执行第一分支语句的时间为T2，线程执行第二分支语句的时间为T1，由于每个线程束均存在执行不同分支的线程，因此每个线程束均需要串行地执行第一分支语句和第二分支语句，所需总时间为T1+T2，第一线程束处理器21运行三个线程束共需时间t1＝3T1+3T2。

图5的下半部分示出的是线程数据重映射之后线程束处理器21分别运行线程束11至13所需的时间，由于引起线程执行第一分支语句的待处理均分配到同一线程束13中，而线程束11和12的线程均只需执行第二分支语句，因此线程束11所需时间为T1，线程束12所需时间为T1，线程束13所需时间为T2，第一线程束处理器21运行三个线程束共需时间时间t2＝2T1+T2。

综上可知，运行线程数据重映射之后，可节约t1-t2＝T1+2T2的时间，可以预见到，当线程束数量不止3个，分支数量不止2个时，所能节约的时间更多。

但是，并不是每种情况下都适合做线程数据重映射，具体可参见图6，图6是根据本发明实施例的线程数据重映射的另一数据流向图，在图6的示例中，只有线程3需执行第一分支语句。

并请结合图7进行参考，图7是是根据本发明实施例的线程数据重映射的另一前后时间开销示意图。

如图7所示，由于只有1个线程束11具有需执行第一分支语句的一个线程3，因此，即便将线程3的待处理数据0转移到最后一个线程12，线程束13也需要串行执行第一分支语句和第二分支语句，因此，在图7的示例中，当需执行第一分支语句的线程的数量为1时，无论执行线程数据重映射与否，均需时3T1+T2。故在这种情况下，进行线程数据重映射是没有效果的，反而会浪费第一线程束处理器21有限的运算资源，并造成无用的时间开销。

值得注意的是，在另外一些示例中，当需执行第一分支语句的线程的数量可以为2、3、4、5等正整数时，进行线程数据重映射也会在一定程度上浪费线程束处理器21有限的运算资源，并造成无用的时间开销。

另外，进行线程数据重映射之前，往往需要进行线程同步，每次线程同步都会造成M*N个线程的暂停，如此一来将会极大地影响GPU的运行速度和运行效率，并且，在需要执行第一分支语句的线程的数量为1或其他较小的数值如为2、3、4以及5时，进行线程数据重映射并不能有效解决分支分歧的问题，反而引入了线程同步，造成延时。

有鉴于此，本发明实施例提供了一种线程处理方法，该方法应用于图形处理器，图形处理器用于处理M个线程束，每个线程束包括N个线程，图形处理器还包括至少一个线程束处理器，至少一个线程束处理器中的第一线程束处理器包括N的整数倍个线程处理器，第一线程束处理器包括第一线程处理器，第一线程处理器用以运行N个线程中的一个以处理满足第一分支语句的判断条件或满足第二分支语句的判断条件的待处理数据，图形处理器中设置有计数器，该方法包括：

第一线程处理器获取需要处理的第一待处理数据，确定第一待处理数据满足第一分支语句，将计数器中的数值加一步长；

第一线程处理器根据计数器的数值确定M*N个线程中需运行第一分支语句的线程的数量；

第一线程处理器在确认数量大于阈值的情况下，执行线程同步以及线程数据重映射。

本发明实施例通过第一线程处理器判断运行第一分支语句的线程的数量，在确认数量大于阈值的情况下，才执行线程同步以及线程数据重映射，可避免在该数量小于或等于阈值的情况下也执行线程同步以及线程数据重映射，可提高GPU的运行速度和运行效率。

为了对上述方法作出清楚说明，具体可参见图8，图8是根据本发明实施例的线程处理方法的流程图，该方法应用于第一线程束处理器21的每一线程处理器，值得注意的是，在本实施例中，假设全局存储器24存储了待处理数据A[tid]，A[tid]是由CPU设置好并发送至GPU的一维数组，其以线程的标识tid作为下标，工作组包括线程束11、12、13，且线程束11、12、13被调度器23投放到第一线程束处理器21。

并且，在本发明实施例中，A[tid]的数值小于5的概率小于A[tid]的数值大于或等于5的概率。

其中，A[tid]的数值小于5为第一分支语句的判断条件，A[tid]的数值大于或等于5为第二分支语句的判断条件，即线程执行第一分支语句的次数小于线程执行第二分支语句的次数。

举例而言，A[tid]＝[13，6，0，1，2，7，8，9，10，11，3，12]。

在一些示例中，A[tid]可在CPU端根据temp＝rand()*1000生成，其中rand()是伪随机数，其会产生0-1之间的任意数值。

值得注意的是，在图8所示的方法被线程处理器运行之前，需先在共享存储器210声明计数器的初始数值为0，标志位的初始数值为第一标志值。

举例而言，第一标志值的初始数值可设置为0。

其中，在共享存储器210声明计数器的初始数值为0，标志位的初始数值为第一标志值的步骤可由调度器23在初始化过程中，在建立工作组并将工作组的线程束11、12、13投放至第一线程束处理器21之前，从第一线程束处理器21选择任一线程处理器执行。

并且，计数器设置在共享存储器210中，第一线程束处理器21的每一线程处理器在执行图8所示的方法以运行线程时，均可访问计数器，具体而言，第一线程束处理器21的每一线程处理器在运行线程时，均可读取计数器的数值，并可修改计数器的数值。

值得注意的是，在本实施例中，第一线程束处理器21的线程处理器在运行线程时，可原子读取计数器的数值，并可对计数器的数值进行原子加一步长。

其中，“原子读取”表示一个线程处理器在运行线程过程中读取计数器的数值时，位于同一线程束处理器的其他线程处理器不能读取计数器的数值，只有前面的线程处理器读取完毕之后，位于同一线程束处理器的其他线程处理器才读取计数器的数值，即，同一时间仅允许同一工作组的M*N个线程中的一个线程读取计数器的数值。

同理，“原子加一步长”是指同一时间仅允许同一工作组的M*N个线程中的一个线程对计数器的数值加一步长。

其中，步长可以是任意正整数，举例而言，可设置为1。

该线程处理方法包括以下步骤：

步骤S401：流程开始。

步骤S402：判断需要处理的待处理数据是否满足第一分支语句的判断条件，如果是，执行步骤S403，如果否，执行步骤S404。

举例而言，第一线程处理器211可根据正在第一线程处理器211上运行的线程1的线程标识tid＝0从全局存储器24读取需要处理的待处理数据A[0]＝13。

步骤S403：对计数器的数值加一步长。

由于本方法应用于同一线程束处理器21中的每一线程处理器，因此，第一线程处理器211、第二线程处理器212、第三线程处理器213以及第四线程处理器214在运行各自的线程时，在各自需要处理的待处理数据满足第一分支语句的判断条件时，可分别对处于共享存储器210的计数器的数值加一步长，即，计数器可累计同一工作组的线程束11、12以及13中需执行第一分支语句的线程的数量。

步骤S404：判断标志位的值是否为第二标志值，如果是执行步骤S408，如果否，执行步骤S405。

第二标志值是与第一标志值不同的数值，当第一标志值是0时，第二标志值是1，当第一标志值是1时，第二标志值可为0。

步骤S405：读取计数器的数值。

具体而言，在本步骤中，线程处理器原子读取计数器的数值，由于本方法应用于同一线程束处理器21中的每一线程处理器，因此，在步骤S403至本步骤之间，可累计第一线程处理器211、第二线程处理器212、第三线程处理器213以及第四线程处理器214分别运行线程束11、线程束12、以及线程束13的线程时，需运行第一分支语句的线程的数量。

步骤S406：判断计数器的数值是否大于阈值，如果是，执行步骤S409，如果否，执行步骤S413。

此处以阈值为1为例。如上所述，具体可根据实际需要设置该阈值，比如设置为大于等于2且小于等于5的正整数。

步骤S407：将标志位的值设置为第二标志值，第二标志值用于指示需要执行重映射。

在本步骤中，第一个发现计数器的数值超过阈值的线程处理器将标志位设置为第二标志值，其他线程处理运行其他线程时，发现第二标志位被设置就会跳至步骤S408，直接执行线程同步，而不再去判断计数器的数值是否超过阈值，避免得出不一样的结论。采用这个设计后，只要M*N个线程中的一个线程作出执行线程同步的决定，其他线程必定跟随，规避了个别线程单独执行线程同步的问题。

其中，个别线程单独执行线程同步的问题是因为以下原因引起的：在步骤S405中，线程处理器运行线程时读取到的计数器的数值是该线程处理器执行原子读取的指令之后从共享存储器210读取到的数值，然而从执行指令到读取数值之间会存在时间间隔，而这个间隔中其他线程可能对计数器的数值进行原子加一的操作，因此不同线程处理器运行各自线程读取到的计数值可能存在差异，从而使得个别线程单独执行线程同步，而其他线程不执行线程同步。

步骤S408：执行线程同步。

正如以上所述，线程处理器运行线程时可通过调用障碍函数barrier()实现线程同步，对于一个调用了障碍函数的线程来说，除非同一个工作组内其他线程都执行了障碍函数，否则该线程将被阻止执行障碍函数之后的内核代码，且在线程处理器执行了障碍函数之后，线程处理器为该线程设置停顿点(break point)，该停顿点记录了障碍函数的下一语句在内核代码中的位置，将该位置记录在该线程的私有存储器中，并暂停运行该线程。

当同一个工作组内其他线程都执行了障碍函数之后，线程处理器从私有存储器中读取停顿点，从而可执行障碍函数的下一语句，使得线程继续运行。

因此，在图2所示的第一线程束处理器21中，只要任一线程处理器在运行线程1至12中任一者时调用了障碍函数，则线程同步结束的条件是线程1至12在被各自对应的线程处理器运行过程中都调用了障碍函数。

步骤S409：在线程同步结束后，将计数器的数值设置为0，将标志位的值设置为第一标志值，以线程标识为索引将第一线程的待处理数据记录在索引表。

其中，索引表设置在共享存储器210中，线程1至12均可访问索引表，线程标识为线程处理器正在运行的线程的线程标识，举例而言，线程处理器可通过调用get_global_id()函数获得线程标识。

举例而言，假设第一线程为线程1，其待处理数据如图4所示为A[0]＝13，则线程1在索引表记录线程标识为0对应待处理数据13。

由于本方法应用于同一线程束处理器21中的每一线程处理器，因此，第一线程可为线程1至12中的任一者。故在线程1至12均运行完本方法之后，产生的索引表如表1所示：

线程标识(索引)	待处理数据
		0	13
1	6
		2	0
3	1
		4	2
5	7
		6	8
7	9
		8	10
9	11
		10	3
11	12

在表1中，线程标识与待处理数据之间具有一一对应关系。

步骤S410：执行线程数据重映射，获取线程数据重映射产生的更新的线程标识，根据更新的线程标识从索引表获取与更新的线程标识对应的待处理数据。

其中，线程数据重映射的原理可参见图3及其对应描述，于此不作赘述。

在本步骤中，举例而言，根据图4，线程数据重映射之后，线程3的tid从2更新为5，因此，第三线程处理器213在运行线程3并执行线程数据重映射之后，获取线程数据重映射产生的更新的线程标识5，根据更新的线程标识5从索引表获取与更新的线程标识5对应的待处理数据7。

步骤S411：判断待处理数据满足第一分支语句的判断条件还是第二分支语句的判断条件，如果满足第一分支语句的判断条件，则执行步骤S412，如果满足第二分支语句的判断条件，则执行步骤S413。

举例而言，第一分支语句的判断条件可例如为A[tid]＜5，第二分支语句的判断条件可例如为A[tid]≥5。

步骤S412：执行第一分支语句。

步骤S413：执行第二分支语句。

步骤S414：流程结束。

在本步骤中，流程结束是指线程处理器结束当前运行的线程，举例而言，第一线程处理器211结束线程1，值得注意的是，第一线程处理器211可在结束线程1后，在线程1所在的线程束11的其他线程均结束时，转而运行其他线程束的线程，如线程5或线程9。

在线程处理器运行每一线程以实现本方法之后，由于每个线程在运行过程中均对计数器的数值与阈值进行比较，在计数器的数值不大于阈值时，不执行线程同步和线程数据重映射，因此可避免出现图7所示的情况，从而防止因无效的线程数据重映射造成延时。

由于GPU并行运算的场景中，涉及多循环的场景较多，因此本发明实施例进一步提供一种线程处理方法，以在多循环的场景实现无用的线程数据重映射的识别。具体而言，该线程处理方法应用于图形处理器，图形处理器用于处理M个线程束，每个线程束包括N个线程，每个线程运行循环语句，图形处理器还包括至少一个线程束处理器，至少一个线程束处理器中的第一线程束处理器包括N的整数倍个线程处理器，第一线程束处理器包括第一线程处理器，第一线程处理器用以在一个循环中运行N个线程中的一个以处理满足第一分支语句的判断条件或满足第二分支语句的判断条件的待处理数据，图形处理器中设置有计数器，该方法包括：

第一线程处理器在第一循环中获取需要处理的第一待处理数据，确定第一待处理数据满足第一分支语句，将计数器中的数值加一步长；

第一线程处理器根据计数器的数值确定M*N个线程中需运行第一分支语句的线程的数量：

第一线程处理器在确认数量大于阈值的情况下，执行线程同步并将计数器中的数值清零；

第一线程处理器执行线程数据重映射。

本发明实施例通过第一线程处理器在确认数量大于阈值的情况下，执行线程同步并将计数器中的数值清零，并执行线程数据重映射，使得第一线程处理器在下一循环读取到的计数器的数值不会受到第一循环的影响，并且，由于在确认数量大于阈值的情况下，才执行线程同步以及线程数据重映射，因此可避免在该数量小于或等于阈值的情况下也执行线程同步以及线程数据重映射，可在每一循环中动态地判断是否有执行线程数据重映射的需要，可避免进行无用的线程数据重映射，可提高GPU的运行速度和运行效率。

为了对上述方法作出清楚说明，具体可参见图9，图9是根据本发明实施例的线程处理方法的另一流程图，该方法应用于第一线程束处理器21的每一线程处理器，值得注意的是，在本实施例中，假设共享存储器24存储了待处理数据A[i，tid]，A[i，tid]是由CPU设置好并发送至GPU的一维数组，其以线程标识tid作为行下标，以循环变量i作为列下标，循环变量i用于指示线程当前所在的循环的序号。工作组包括线程束11、12、13，且线程束11、12、13被调度器23投放到第一线程束处理器21。并且，线程执行第一分支语句的概率小于线程执行第二分支语句的概率。

在另一些示例中，CPU也可以将待处理数据设置为A[tid，i]。

在本实施例中，以循环语句for(i＝0；i＜1000；i++)为例进行说明，且tid为0至11，因此A[i，tid]具有1000行，12列。由于二维数组涉及的数值较多，为便于说明，仅示出A[i，tid]之部分，具体而言，在本实施例中，假设线程束11在当前时刻运行至i＝11的循环，线程束12在当前时刻运行至i＝8的循环，线程束13在当前时刻运行至i＝10的循环。

并且，假设在当前时刻，线程1需要处理的待处理数据是A[11，0]＝13，线程2需要处理的待处理数据是A[11，1]＝6，线程3需要处理的待处理数据是A[11，2]＝0，线程4需要处理的待处理数据是A[11，3]＝1，线程5需要处理的待处理数据是A[8，4]＝2，线程6需要处理的待处理数据是A[8，5]＝7，线程7需要处理的待处理数据是A[8，6]＝8，线程8需要处理的待处理数据是A[8，7]＝9，线程9需要处理的待处理数据是A[10，8]＝10，线程10需要处理的待处理数据是A[10，9]＝11，线程11需要处理的待处理数据是A[10，10]＝3，线程12在需要处理的待处理数据是A[10，11]＝12。具体可结合图10的上半部分进行参考，其中，图10是根据本发明实施例的线程数据重映射的前后时间开销示意图。

值得注意的是，在本实施例中，A[i，tid]数组设置为在i为固定值时A[i，tid]的数值小于5的概率小于A[i，tid]的数值大于或等于5的概率，即在工作组的每次循环中，线程执行第一分支语句的次数小于执行第二分支语句的次数。

在一些示例中，A[i，tid]可由CPU根据temp＝rand()*1000生成，其中rand()是伪随机数，其会产生0-1之间的任意数值。

并且，与上一实施例类似，在图9所示的方法被线程处理器运行之前，需先在共享存储器210声明计数器的初始数值为0，标志位的初始数值为第一标志值。

举例而言，第一标志值可为0。

其中，在共享存储器210声明计数器的初始数值为0，标志位的初始数值为第一标志值的步骤可由调度器23在初始化过程中，在建立工作组并将工作组的线程束11、12、13投放至第一线程束处理器21之前，从第一线程束处理器21选择任一线程处理器执行。

并且，计数器设置在共享存储器210中，第一线程束处理器21的每一线程处理器在运行线程时，均可访问计数器，具体而言，第一线程束处理器21的每一线程处理器在运行线程时，均可读取计数器的数值，并可修改计数器的数值。

值得注意的是，在本实施例中，第一线程束处理器21的线程处理器在运行线程时，可原子读取计数器的数值，并可对计数器的数值进行原子加一步长。

其中，“原子读取”表示一个线程处理器在运行线程过程中读取计数器的数值时，位于同一线程束处理器的其他线程处理器不能读取计数器的数值，只有前面的线程处理器读取完毕之后，位于同一线程束处理器的其他线程处理器才读取计数器的数值，即，同一时间仅允许同一工作组的M*N个线程中的一个线程读取计数器的数值。

同理，“原子加一步长”是指同一时间仅允许同一工作组的M*N个线程中的一个线程对计数器的数值加一步长。

其中，步长可以是任意正整数，举例而言，可设置为1。

该流程处理方法包括以下步骤：

步骤S501：流程开始。

步骤S502：判断循环是否结束，如果是，执行步骤S503，如果否，执行步骤S504。

可选地，在本实施例中，循环语句例如为for(i＝0；i＜1000；i++)，其中，i为循环变量，循环变量用于指示线程已完成的循环的次数。

其中，i的初始值是0，且线程在完成一次循环时，对i加1，当i的值累加至1000时，线程可跳出循环语句，此时循环结束。

步骤503：流程结束。

在本步骤中，流程结束是指线程处理器结束当前运行的线程，举例而言，第一线程处理器211结束线程1，值得注意的是，第一线程处理器211可在结束线程1后，在线程1所在的线程束11的其他线程均结束时，转而运行其他线程束的线程，如线程5或线程9。

步骤S504：判断需要处理的待处理数据是否满足第一分支语句的判断条件，如果是，执行步骤S505，如果否，执行步骤S506。

其中，全局存储器24记录有需要处理的待处理数据，需要处理的待处理数据为A[i，tid]，线程处理器可根据在线程处理器上运行的线程的线程标识和循环变量的数值，从全局存储器24获取A[i，tid]。

举例而言，第一分支语句的判断条件为A[i，tid]＜5。

步骤S505：对计数器的数值加一步长。

由于本方法应用于同一线程束处理器21中的每一线程处理器，因此，第一线程处理器211、第二线程处理器212、第三线程处理器213以及第四线程处理器214在运行各自的线程时，在各自需要处理的待处理数据满足第一分支语句的判断条件时，可分别对处于共享存储器210的计数器的数值加一步长，即，计数器可累计同一工作组的线程束11、12以及13中需执行第一分支语句的线程的数量。

步骤S506：判断标志位的值是否为第二标志值，如果是执行步骤S510，如果否，执行步骤S507。

第二标志值是与第一标志值不同的数值，第二标志值用于指示需要执行重映射，在一些示例中，第一标志值是0，第二标志值是1，在另一些示例中，第一标志值是1，第二标志值是0。

步骤S507：读取计数器的数值。

具体而言，在本步骤中，线程处理器原子读取计数器的数值，由于本方法应用于同一线程束处理器21中的每一线程处理器，因此，在步骤S505至本步骤之间，可累计第一线程处理器211、第二线程处理器212、第三线程处理器213以及第四线程处理器214分别运行线程束11、线程束12、以及线程束13的线程时，需运行第一分支语句的线程的数量。

步骤S508：判断计数器的数值是否大于阈值，如果是，执行步骤S509，如果否，执行步骤S513。

在本实施例中，阈值为1。

在另外一些示例中，阈值可为2-5之间的正整数。

步骤S509：将标志位的值设置为第二标志值，第二标志值用于指示需要执行重映射。

在本步骤中，第一个发现计数器的数值超过阈值的线程处理器将标志位设置为第二标志值，其他线程处理运行其他线程时，在步骤S506读取到标志位设置为第二标志值时会跳至步骤S510，直接执行线程同步，而不再去判断计数器的数值是否超过阈值(步骤508)，避免得出不一样的结论。采用这个设计后，只要M*N个线程中的一个线程作出执行线程同步的决定，其他线程必定跟随，规避了个别线程单独执行线程同步的问题。

其中，个别线程单独执行线程同步的问题是因为以下原因引起的：在步骤S507中，线程处理器运行线程时读取到的计数器的数值是该线程处理器执行原子读取的指令之后从共享存储器210读取到的数值，然而从执行指令到读取数值之间会存在时间间隔，而这个间隔中其他线程在各自的循环中可能对计数器的数值进行原子加一的操作，因此不同线程处理器运行各自线程读取到的计数值可能存在差异，从而使得个别线程单独执行线程同步，而其他线程不执行线程同步。

步骤S510：执行线程同步。

正如以上所述，线程处理器运行线程时可通过调用障碍函数barrier()实现线程同步，对于一个调用了障碍函数的线程来说，除非同一个工作组内其他线程都执行了障碍函数，否则该线程将被阻止执行障碍函数之后的内核代码，且在线程处理器执行了障碍函数之后，线程处理器为该线程设置停顿点(break point)，该停顿点记录了障碍函数的下一语句在内核代码中的位置，将该位置记录在该线程的私有存储器中，并暂停运行该线程。

当同一个工作组内其他线程都执行了障碍函数之后，线程处理器从私有存储器中读取停顿点，从而可执行障碍函数的下一语句，使得线程继续运行。

因此，针对图2所示的第一线程束处理器21，只要任一线程处理器在运行线程1至12中任一者时调用了障碍函数，则线程同步结束的条件是线程1至12在被各自的线程处理器运行过程中都调用了障碍函数。

步骤S511：在线程同步结束后，将计数器的数值设置为0，将标志位设置为第一标志值，以线程标识为索引将第一线程的待处理数据和循环变量记录在索引表。

其中，索引表设置在共享存储器210中，线程1至12均可访问索引表。

由于本方法应用于同一线程束处理器21中的每一线程处理器，因此，第一线程可为线程1至12中的任一者。故在线程1至12均运行完本方法之后，产生的索引表如表2所示：

其中，线程标识、待处理数据以及循环变量三者之间具有一一对应关系。

步骤S512：执行线程数据重映射，获取线程数据重映射产生的更新的线程标识，根据更新的线程标识从索引表获取与更新的线程标识对应的待处理数据。

其中，线程数据重映射的原理可参见图3及其对应描述。

举例而言，请参见图10的下半部分，线程数据重映射之后，线程3的tid从2更新为5，因此，第三线程处理器213在运行线程3并执行线程数据重映射之后，获取线程数据重映射产生的更新的线程标识5，根据更新的线程标识5从索引表获取与更新的线程标识5对应的待处理数据7。

其他线程亦可根据类似方式获取待处理数据，具体可参见图10，于此不作赘述。

步骤S513：对计数器的数值减一步长。

举例而言，步长可为任意正整数，在本实施例中，步长设置为1，且对计数器的数值减一步长具体为对计数器的数值进行原子减一操作。

在本步骤中，在待处理数据不满足第一分支语句的判断条件时对计数器的数值减一步长，与步骤S505中加一步长的操作抵消，可避免对其他线程的判断造成干扰。

步骤S514：判断待处理数据满足第一分支语句的判断条件还是第二分支语句的判断条件，如果满足第一分支语句的判断条件，则执行步骤S515，如果满足第二分支语句的判断条件，则执行步骤S516。

步骤S515：执行第一分支语句。

步骤S516：执行第二分支语句。

步骤S5517：对循环变量加一，并跳转至步骤S502。

举例而言，并请结合图11，图11是根据本发明实施例的线程数据重映射的另一前后时间开销示意图，图11的上部和中部示出对各个线程分别对自身的循环变量加1进入下一循环的过程，在每一线程分别执行本步骤之后，线程1的i＝12，tid＝0，线程2的i＝12，tid＝1，线程3的i＝9，tid＝5，线程4的i＝9，tid＝6，线程5的i＝9，tid＝7，线程6的i＝11，tid＝8，线程7的i＝11，tid＝9，线程8的i＝11，tid＝11，线程9的i＝11，tid＝10，线程10的i＝9，tid＝4，线程11的i＝12，tid＝3，线程12的i＝12，tid＝2。

并且，在第一线程处理器211运行线程1并进入线程1的下一循环时，从全局存储器24读取A[12，0]＝100，第二线程处理器212运行线程2并进入线程2的下一循环时，从全局存储器24读取A[12，1]＝2，第三线程处理器213运行线程3并进入线程3的下一循环时，从全局存储器24读取A[9，5]＝4，第四线程处理器214运行线程4并进入线程4的下一循环时，从全局存储器24读取A[9，6]＝101。

在第一线程处理器211运行线程5并进入线程1的下一循环时，从全局存储器24读取A[11，9]＝666，第二线程处理器212运行线程6并进入线程6的下一循环时，从全局存储器24读取A[11，11]＝410，第三线程处理器213运行线程7并进入线程7的下一循环时，从全局存储器24读取A[11，10]＝510，第四线程处理器214运行线程8并进入线程8的下一循环时，从全局存储器24读取A[11，11]＝410。

在第一线程处理器211运行线程9并进入线程9的下一循环时，从全局存储器24读取A[11，10]＝510，第二线程处理器212运行线程10并进入线程10的下一循环时，从全局存储器24读取A[9，4]＝777，第三线程处理器213运行线程11并进入线程11的下一循环时，从全局存储器24读取A[12，3]＝63，第四线程处理器214运行线程12并进入线程12的下一循环时，从全局存储器24读取A[12，2]＝1。

各线程处理器根据当前循环需要处理的待处理执行图9中步骤S504及其以下的步骤，值得注意的是，在本轮循环中，由于计数器的数值是3，大于阈值1，因此会进行图11下半部分所示的线程数据重映射，其中，线程数据重映射的具体方式可参见图12，于此不作赘述，其中图12是根据本发明实施例的线程数据重映射的原理图。

综上，线程处理器在确认数量大于阈值的情况下，执行线程同步并将计数器中的数值清零，并执行线程数据重映射，使得第一线程处理器在下一循环读取到的计数器的数值不会受到第一循环的影响，并且，由于在确认数量大于阈值的情况下，才执行线程同步以及线程数据重映射，因此可避免在该数量小于或等于阈值的情况下也执行线程同步以及线程数据重映射，可在每一循环中动态地判断是否有执行线程数据重映射的需要，可避免进行无用的线程数据重映射，可提高GPU的运行速度和运行效率。

下文将从单个线程的角度进一步清楚说明图4、图6以及图12涉及的线程数据重映射的原理，其中，图形处理器的共享存储器210设置有一维数组id_pood[]、第一变量H以及第二变量R，其中一维数组的长度是M*N，第一变量H的初始值是0，第二变量R的初始值是M*N-1，可通过以下方式执行线程数据重映射：

步骤1：在确认待处理数据满足第二分支语句的判断条件时，读取第一变量H的数值，并将线程标识写入一维数组id_pood[]中以第一变量H的数值作为下标的位置，将第一变量H的数值加一，并执行线程同步；在确认第一线程的第一待处理数据满足第一分支语句的判断条件时，读取第二变量R的数值，并将线程标识写入一维数组id_pood[]中以第二变量的数值作为下标的位置，将第二变量R的数值减一，并执行线程同步；

步骤2：在线程同步结束后，读取一维数组id_pood[]中以第一线程的线程标识作为下标的位置上的数值，并将读取的数值作为线程数据重映射产生的更新的线程标识。

值得注意的是，上述方法的执行主体分别是第一线程处理器211、第二线程处理器212、第三线程处理器213和第四线程处理器214。

并且，线程同步结束是指工作组内的所有线程，如线程1至12，均执行了线程同步。

举例而言，在图3所示的场景中，第一线程处理器211在运行线程1时，判断到线程1的待处理数据13满足第二分支语句的判断条件(大于或等于5)，读取第一变量H的数值0，将线程1的线程标识0写入id_pood[0]，并将第二变量R的数值原子加1，使之变成1。

第二线程处理器212在运行线程2时，判断到线程2的待处理数据6满足第二分支语句的判断条件(大于或等于5)，读取第一变量H的数值1，将线程2的线程标识1写入id_pood[1]，并将第一变量H的数值原子加1，使之变成2。

第三线程处理器213在运行线程3时，判断到线程3的待处理数据0满足第一分支语句的判断条件(小于5)，读取第二变量R的数值11(M＝3，N＝4，M*N-1＝12-1＝11)，将线程3的线程标识2写入id_pood[11]，并将第二变量R的数值原子加1，使之变成10。

第四线程处理器214在运行线程4时，判断到线程4的待处理数据1满足第一分支语句的判断条件(小于5)，读取第二变量R的数值10，将线程4的线程标识3写入id_pood[10]，并将第二变量R的数值原子加1，使之变成9。

同理，第一线程处理器211分别运行线程5和线程9，第二线程处理器212分别运行线程6和线程10，第三线程处理器213分别运行线程7和线程11，第四线程处理器214分别运行线程8和线程12时，作出类似处理，得到id_pood[]如下：

0

1

5

6

7

8

9

11

10

4

3

2

第一线程处理器211线程同步结束时，以线程1的线程标识0作为下标读取id_pood[0]＝0，以线程5的线程标识4作为下标读取id_pood[4]＝7，以线程9的线程标识8作为下标读取id_pood[8]＝10。

第二线程处理器212线程同步结束时，以线程2的线程标识1作为下标读取id_pood[1]＝1，以线程6的线程标识5作为下标读取id_pood[5]＝8，以线程10的线程标识9作为下标读取id_pood[9]＝10。

第三线程处理器213线程同步结束时，以线程3的线程标识2作为下标读取id_pood[2]＝5，以线程7的线程标识6作为下标读取id_pood[6]＝9，以线程11的线程标识10作为下标读取id_pood[10]＝3。

第四线程处理器214线程同步结束时，以线程4的线程标识3作为下标读取id_pood[3]＝6，以线程8的线程标识7作为下标读取id_pood[7]＝11，以线程12的线程标识11作为下标读取id_pood[8]＝12。

针对图6以及图12所示的场景，也可以类似的方式进行线程数据重映射，于此不作赘述。

本发明实施例进一步提供一种线程处理方法，该方法应用于图形处理器的线程束处理器，图形处理器包括M个线程束，每个线程束包括N个线程，M个线程束的M*N个线程中存在至少一个线程需运行第一分支语句，该方法包括：

检测M*N个线程中需运行第一分支语句的线程的数量；

在确认数量大于阈值的情况下，对M*N个线程进行线程数据重映射。

为了清楚说明，以下结合图13进行详细说明，图13是根据本发明实施例的线程处理方法的另一流程图，图13所示的方法应用于第一线程束处理器21，其图8所示的实施例相比，区别在于本实施例以第一线程束处理器21作为执行主体进行描述，该方法包括：

步骤S101：检测M*N个线程中需运行第一分支语句的线程的数量。

举例而言，针对图2所示，M＝3，N＝4。

针对以下代码：

第一分支语句的判断条件为A[tid]＜5。

步骤S102：判断该数量是否大于阈值，如果是，执行步骤S104，如果否，执行步骤S103。

阈值可例如为1。

步骤S103：不进行线程数据重映射。

步骤S104：进行线程数据重映射。

步骤S105：运行M*N个线程并根据待处理数据执行第一分支语句或第二分支语句。

在本发明实施例中，由于在步骤S102中引入阈值判断的步骤，因此可以过滤无效的线程数据重映射，从而节约图形处理器的运算资源和减少不必要的时间开销。

值得注意的是，当阈值＝1时，可杜绝在只有1个线程需运行第一分支语句而进行无效的线程数据重映射的情况发生，但是，在实际应用中，由于线程数量较多，也可以将阈值设置为其他数值，例如阈值＝N，或N的整数倍，可以将执行第一分支语句的线程集中设置在一个或多个线程束中，该些线程束的线程只需执行第一分支语句。或者，在另一些示例中，阈值也可以设置为经验值，通过实验可以对阈值进行取值，来达到时间开销与运算资源之间的平衡。

在上述实施例中，只对执行一个分支语句的情况进行说明，但是，在实际应用中，分支语句往往嵌套在循环语句中，因此分支语句需要循环执行多次，此时，若每层循环均执行线程数据重映射，将极大地浪费运算资源，并造成较高的时间开销。

举例而言，可参见以下的内核代码：

由于for循环的存在，每个线程需执行分支判断1000次，此时无效的线程数据重映射会造成较多的时间开销和浪费较多的运算资源。

为此，本发明另一实施例进一步提供一种运行在循环中的线程处理方法，以解决上述技术问题。

以下请参见图14，图14是根据本发明实施例的线程处理方法的另一流程图，该方法运行于第一线程束处理器21，其图10所示的实施例相比，区别在于本实施例以第一线程束处理器21作为执行主体进行描述，，该方法包括：

如图14所示，该方法具体包括以下步骤：

步骤S201：获取M*N个线程需要处理的待处理数据，在任一线程的待处理数据满足第一分支语句的判断条件的情况下，将计数器的数值加一。

并请结合图10一并参考，图10是根据本发明实施例的线程数据重映射的另一前后时间开销示意图。在图10的上半部分示出，线程束11当前处于第12个循环，循环变量i＝11，线程束12当前处于第9个循环，循环变量i＝8，线程束13当前处于第11个循环，循环变量i＝10。

根据图10，线程3、4、5、11的待处理数据满足第一分支语句的判断条件(小于5)，因此计数器的数值进行4次加一的操作，使得计数器的数值的值＝4。

步骤S202：读取计数器的数值。

步骤S203：判断计数器的数值是否大于阈值，如果是，执行步骤S204，如果否，执行步骤S213。

步骤S204：控制确认计算变量大于阈值的线程以自身的线程标识作为索引记录自身的待处理数据和循环变量，将标志位设置为第二标志值，执行线程同步。

步骤S205：控制其他线程在检测到标志位为第二标志值的情况下，以自身的线程标识作为索引记录自身的待处理数据和循环变量到索引表，执行线程同步。

步骤S206：控制M*N个线程执行线程数据重映射以获取更新的线程标识。

本步骤的线程数据重映射的具体过程可参见图10的下半部分所示，经线程数据重映射之后，线程1的线程标识为0，线程2的线程标识为1，线程3的线程标识为5，线程4的线程标识为6，线程5的线程标识为7，线程6的线程标识为8，线程7的线程标识为9，线程8的线程标识为11，线程9的线程标识为10，线程10的线程标识为4，线程11的线程标识为3，线程12的线程标识为2

步骤S207：将计数变量设置为0，将标志位设置为第一标志值。

在本步骤中，通过对标志位和计数器的数值清零，可确保本循环的数值不会对下一循环造成影响。

步骤S208：控制M*N个线程以更新的线程标识作为索引在索引表获取待处理数据和循环变量。

步骤S209：控制M*N个线程在各自的待处理数据在满足第一分支语句的判断条件的情况下，执行第一分支语句，满足第二分支语句的判断条件的情况下，执行第二分支语句。

根据图10，经线程数据重映射之后，线程1至8均执行第二分支语句，线程9至12均执行第一分支语句。

步骤S210：控制M*N个线程线程根据各自的循环变量判断是否进入下一循环，如果是，执行步骤S212，如果否，执行步骤S211。

具体而言，针对循环语句for，可将循环变量i进行加一，在加一之后的i小于1000的情况下，进入下一循环，反之，则结束循环。

步骤S211：控制线程退出循环。

步骤S212：控制线程进入下一循环。

步骤S213：将计数器的数值设置为0，获取M*N个线程的待处理数据和循环变量。

步骤S214：控制M*N个线程在各自的待处理数据满足第一分支语句的判断条件的情况下，执行第一分支语句，满足第二分支语句的判断条件的情况下，执行第二分支语句。

步骤S215：控制M*N个线程根据各自的循环变量判断是否进入下一循环，如果是，执行步骤S217，如果否，执行步骤S216。

步骤S216：控制线程退出循环。

步骤S217：控制线程进入下一循环。

在图10的示例中，由于控制执行线程数据重映射后的线程的循环变量在进行加一操作之后均小于1000，因此进入下一循环，执行步骤S201。

具体可结合图11和12进行参考，由于前述以线程处理器的角度对图11和图12进行过详细介绍，于此不做赘述。

示例地，为帮助理解，以下将列出内核代码的一种具体形式，其中，该内核代码以C语言编写，可结合该内核代码对本发明实施例进行理解，而在每句代码后的注释则说明了该代码的功能：

以上内核代码仅为本发明实施例中第一个示例，可由CPU端的代码编辑器编辑，线程运行该内核代码，可实现图10所述的方法。

可参见图15，图15是根据本发明实施例的异构系统的装置结构示意图，如图15所示，异构系统包括中央处理器30和图形处理器20，中央处理器30包括主机代码301、内核代码302、编译器304以及运行时刻库307，其中，主机代码301和内核代码302设置在代码编辑器300上。

内核代码302举例而言可为：

其中，代码编辑器300可将分支处理代码设置在内核代码302中，形成新的内核代码。

举例而言，代码编辑器300可将分支处理代码加入到内核代码302中，例如：将

加入到以下内核代码之前，

形成新的内核代码。

其中，主机代码301中设置了待处理数据A[i，tid]，例如为A[i，tid]＝rand()*1000。

代码编辑器300将二进制的主机代码301和内核代码302发送至编译器304，编译器304产生二进制的内核代码和主机代码。

CPU将二进制的内核代码、主机代码以及待处理数据A[i，tid]发送至图形处理器20。

本发明实施例进一步提供一种图形处理器，其包括第一线程处理器，第一线程处理器用于执行图8或图9所示的方法。

本发明实施例进一步提供一种图形处理器，其包括第一线程束处理器，第一线程束处理器用于执行图13或图14所示的方法。

需说明的是，以上描述的任意装置实施例都仅仅是示意性的，其中所述作为分离部件说明的单元可以是或者也可以不是物理上分开的，作为单元显示的部件可以是或者也可以不是物理单元，即可以位于一个地方，或者也可以分布到多个网络单元上。可以根据实际的需要选择其中的部分或者全部模块来实现本实施例方案的目的。另外，本发明提供的装置实施例附图中，模块之间的连接关系表示它们之间具有通信连接，具体可以实现为一条或多条通信总线或信号线。本领域普通技术人员在不付出创造性劳动的情况下，即可以理解并实施。

通过以上的实施方式的描述，所属领域的技术人员可以清楚地了解到本发明可借助软件加必需的通用硬件的方式来实现，当然也可以通过专用硬件包括专用集成电路、专用CPU、专用存储器、专用元器件等来实现。一般情况下，凡由计算机程序完成的功能都可以很容易地用相应的硬件来实现，而且，用来实现同一功能的具体硬件结构也可以是多种多样的，例如模拟电路、数字电路或专用电路等。但是，对本发明而言更多情况下软件程序实现是更佳的实施方式。基于这样的理解，本发明的技术方案本质上或者说对现有技术做出贡献的部分可以以软件产品的形式体现出来，该计算机软件产品存储在可读取的存储介质中，如计算机的软盘，U盘、移动硬盘、只读存储器(ROM，Read-Only Memory)、随机存取存储器(RAM，Random Access Memory)、磁碟或者光盘等，包括若干命令用以使得一台计算机设备(可以是个人计算机，服务器，或者网络设备等)执行本发明各个实施例所述的方法。

所属领域的技术人员可以清楚地了解到，上述描述的系统、装置或单元的具体工作过程，可以参考前述方法实施例中的对应过程，在此不再赘述。

以上所述，仅为本发明的具体实施方式，但本发明的保护范围并不局限于此，任何熟悉本技术领域的技术人员在本发明揭露的技术范围内，可轻易想到变化或替换，都应涵盖在本发明的保护范围之内。因此，本发明的保护范围应以所述权利要求的保护范围为准。

Claims (38)

1.一种线程处理方法，其特征在于，所述方法应用于图形处理器，所述图形处理器用于处理M个线程束，每个所述线程束包括N个线程，所述图形处理器还包括至少一个线程束处理器，所述至少一个线程束处理器中的第一线程束处理器包括N的整数倍个线程处理器，所述第一线程束处理器包括第一线程处理器，所述第一线程处理器用以运行N个线程中的一个以处理满足第一分支语句的判断条件或满足第二分支语句的判断条件的待处理数据，所述图形处理器中设置有计数器，所述方法包括：

所述第一线程处理器获取需要处理的第一待处理数据，确定所述第一待处理数据满足第一分支语句，将所述计数器中的数值加一步长；

所述第一线程处理器根据所述计数器的数值确定M*N个线程中需运行第一分支语句的线程的数量；

所述第一线程处理器在确认所述数量大于阈值的情况下，执行线程同步以及线程数据重映射。

2.根据权利要求1所述的方法，其特征在于，所述图形处理器还设置有标志位，所述标志位的值设置为第一标志值，所述第一标志值用于指示不执行重映射，所述方法包括：

所述第一线程处理器在确定所述数量大于阈值的情况之前，读取所述标志位；且，

所述第一线程处理器在确定所述数量大于阈值的情况之后及执行线程同步之前，将所述第一标志值设置为第二标志值，所述第二标志值用于指示需要执行重映射。

3.根据权利要求2所述的方法，其特征在于，所述第一线程处理器在执行所述线程同步之后且执行所述线程数据重映射之前，所述方法还包括：

所述第一线程处理器将所述计数器中的数值清零。

4.根据权利要求2或3所述的方法，其特征在于，所述第一线程束处理器包括第二线程处理器，所述第二线程处理器用以在运行N个线程中的一个以处理满足第一分支语句或满足第二分支语句的待处理数据，所述方法还包括：

第二线程处理器读取所述标志位，在确认所述标志位的值为所述第二标志值时，执行线程同步以及线程数据重映射；

所述第二线程处理器在确认所述标志位的值为所述第一标志值时，根据所述计数器的数值确定所述M*N个线程中需运行第一分支语句的线程的数量，在确认所述数量大于阈值的情况下，执行线程同步以及线程数据重映射。

5.根据权利要求4所述的方法，其特征在于，所述第一线程处理器用以运行N个线程中的第一线程以处理满足第一分支语句的判断条件的待处理数据，所述第二线程处理器用以运行N个线程中的第二线程以处理满足第二分支语句的判断条件的待处理数据，所述图形处理器还设置有一维数组、第一变量以及第二变量，其中所述一维数组的长度是M*N，所述第一变量的初始值是0，所述第二变量的初始值是M*N-1，所述第一线程处理器执行所述线程数据重映射，包括：

所述第一线程处理器读取所述第二变量的数值，并将所述第二线程的线程标识写入所述一维数组中以所述第二变量的数值作为下标的位置，将所述第二变量的数值减一，并执行所述线程同步；

所述第二线程处理器读取所述第一变量的数值，并将所述第二线程的线程标识写入所述一维数组中以所述第一变量的数值作为下标的位置，将所述第一变量的数值加一，并执行所述线程同步；

所述第一线程处理器在所述线程同步结束后，读取所述一维数组中以所述第一线程的线程标识作为下标的位置上的数值，并将读取的所述数值作为所述线程数据重映射产生的所述第一线程的更新的线程标识；

所述第二线程处理器在所述线程同步结束后，读取所述一维数组中以所述第二线程的线程标识作为下标的位置上的数值，并将读取的所述数值作为所述线程数据重映射产生的所述第二线程的更新的线程标识。

6.根据权利要求5所述的方法，其特征在于，所述第一线程处理器在执行所述线程同步之后且执行所述线程数据重映射之前，所述方法还包括：

所述第一线程处理器以所述第一线程的线程标识作为索引将所述第一待处理数据记录在索引表，其中所述第一线程的线程标识与所述第一待处理数据具有一一对应关系，所述索引表记录有所述M*N个线程的线程标识与待处理数据之间的一一对应关系；

所述第一线程处理器在执行所述线程数据重映射之后，所述方法还包括：

所述第一线程处理器以执行所述线程数据重映射后产生的所述第一线程的更新的线程标识作为索引在所述索引表中读取与所述第一线程的更新的线程标识对应的第三待处理数据；

所述第一线程处理器在所述第三待处理数据满足第一分支语句的判断条件时，执行所述第一分支语句，所述第一线程处理器在所述第三待处理数据满足第二分支语句的判断条件时，执行所述第二分支语句。

7.根据权利要求6所述的方法，其特征在于，所述阈值为1。

8.根据权利要求6所述的方法，其特征在于，所述阈值为大于等于2且小于等于5的正整数。

9.根据权利要求1至3任一项所述的方法，其特征在于，所述第一线程处理器执行所述第一分支语句的概率小于所述第一线程处理器执行所述第二分支语句的概率。

10.一种线程处理方法，其特征在于，所述方法应用于图形处理器，所述图形处理器用于处理M个线程束，每个所述线程束包括N个线程，所述图形处理器还包括至少一个线程束处理器，所述至少一个线程束处理器中的第一线程束处理器包括N的整数倍个线程处理器，所述第一线程束处理器包括第一线程处理器，所述第一线程处理器运行循环语句，用以在一个循环中运行N个线程中的一个以处理满足第一分支语句的判断条件或满足第二分支语句的判断条件的待处理数据，所述图形处理器中设置有计数器，所述方法包括：

所述第一线程处理器在第一循环中获取需要处理的第一待处理数据，确定所述第一待处理数据满足第一分支语句，将所述计数器中的数值加一步长；

所述第一线程处理器根据所述计数器的数值确定M*N个线程中需运行第一分支语句的线程的数量；

所述第一线程处理器在确认所述数量大于阈值的情况下，执行线程同步并将所述计数器中的数值清零；

所述第一线程处理器执行线程数据重映射。

11.根据权利要求10所述的方法，其特征在于，所述方法还包括：

所述第一线程处理器获取在第二循环中需要处理的第二待处理数据，确定所述第二待处理数据满足所述第二分支语句的判断条件，将所述计数器中的数值减一步长。

12.根据权利要求10或11所述的方法，其特征在于，所述图形处理器还设置有标志位，所述标志位的值设置为第一标志值，所述第一标志值用于指示不执行重映射，所述方法包括：

所述第一线程处理器在确定所述数量大于阈值的情况之前，读取所述标志位；且，

所述第一线程处理器在确定所述数量大于阈值的情况之后及执行线程同步之前，将所述第一标志值设置为第二标志值，所述第二标志值用于指示需要执行重映射。

13.根据权利要求12所述的方法，其特征在于，所述第一线程束处理器包括第二线程处理器，所述第二线程处理器用以在运行N个线程中的一个以处理满足第一分支语句或满足第二分支语句的待处理数据，所述方法还包括：

第二线程处理器读取所述标志位，在确认所述标志位的值为所述第二标志值时，执行线程同步以及线程数据重映射；

第二线程处理器在确认所述标志位的值为所述第一标志值时，根据所述计数器的数值确定所述M*N个线程中需运行第一分支语句的线程的数量，在确认所述数量大于阈值的情况下，执行线程同步以及线程数据重映射。

14.根据权利要求13所述的方法，其特征在于，所述第一线程处理器用以运行N个线程中的第一线程以处理满足第一分支语句的判断条件的的待处理数据，所述第二线程处理器用以运行N个线程中的第二线程以处理满足第二分支语句的判断条件的待处理数据，所述图形处理器还设置有一维数组、第一变量以及第二变量，其中所述一维数组的长度是M*N，所述第一变量的初始值是0，所述第二变量的初始值是M*N-1，所述第一线程处理器执行所述线程数据重映射，包括：

所述第一线程处理器读取所述第二变量的数值，并将所述第二线程的线程标识写入所述一维数组中以所述第二变量的数值作为下标的位置，将所述第二变量的数值减一，并执行所述线程同步；

所述第二线程处理器读取所述第一变量的数值，并将所述第二线程的线程标识写入所述一维数组中以所述第一变量的数值作为下标的位置，将所述第一变量的数值加一，并执行所述线程同步；

所述第一线程处理器在所述线程同步结束后，读取所述一维数组中以所述第一线程的线程标识作为下标的位置上的数值，并将读取的所述数值作为所述线程数据重映射产生的所述第一线程的更新的线程标识；

所述第二线程处理器在所述线程同步结束后，读取所述一维数组中以所述第二线程的线程标识作为下标的位置上的数值，并将读取的所述数值作为所述线程数据重映射产生的所述第二线程的更新的线程标识。

15.根据权利要求14所述的方法，其特征在于，所述第一线程处理器在执行所述线程同步之后且执行所述线程数据重映射之前，所述方法还包括：

所述第一线程处理器以所述第一线程的线程标识作为索引将所述第一待处理数据和第一循环变量记录在索引表，其中，所述第一线程的线程标识与所述第一待处理数据具有一一对应关系；

所述第一线程处理器在执行所述线程数据重映射之后，所述方法还包括：

所述第一线程处理器以执行所述线程数据重映射后产生的所述第一线程的更新的线程标识作为索引在所述索引表中读取与所述第一线程的更新的线程标识对应的第三待处理数据；

所述第一线程处理器在所述第三待处理数据满足第一分支语句的判断条件时，执行所述第一分支语句，所述第一线程处理器在所述第三待处理数据满足第二分支语句的判断条件时，执行所述第二分支语句。

16.根据权利要求15所述的方法，其特征在于，所述图形处理器中还记录有每个线程的循环变量，所述循环变量用于指示线程当前所在的循环的序号，所述索引表中记录有所述第一线程的循环变量与所述第一线程的线程标识、第一线程在所述循环变量所指示的循环中的待处理数据的对应关系，所述第一线程处理器在执行所述线程数据重映射之后，所述方法还包括：

所述第一线程处理器以执行所述线程数据重映射后产生的所述第一线程的更新的线程标识作为索引在所述索引表中读取与所述第一线程的更新的线程标识对应的循环变量；

所述第一线程处理器在执行所述第一分支语句或所述第二分支语句之后，将所述第一线程的更新的线程标识对应的循环变量加一以获取更新的循环变量，且在所述更新的循环变量不符合所述循环语句规定的循环条件时，结束所述第一线程，在所述更新的循环变量符合所述循环语句规定的循环条件时，运行所述第一线程的第二循环。

17.根据权利要求16所述的方法，其特征在于，所述阈值为1。

18.根据权利要求16所述的方法，其特征在于，所述阈值为大于或等于2且小于或等于5的正整数。

19.根据权利要求10或11所述的方法，其特征在于，所述第一线程处理器执行所述第一分支语句的概率小于所述第一线程处理器执行所述第二分支语句的概率。

20.一种图形处理器，所述图形处理器用于处理M个线程束，每个所述线程束包括N个线程，所述图形处理器还包括至少一个线程束处理器，所述至少一个线程束处理器中的第一线程束处理器包括N的整数倍个线程处理器，所述第一线程束处理器包括第一线程处理器，所述第一线程处理器用以运行N个线程中的一个以处理满足第一分支语句的判断条件或满足第二分支语句的判断条件的待处理数据，所述图形处理器中设置有计数器，其中，

所述第一线程处理器，用于获取需要处理的第一待处理数据，确定所述第一待处理数据满足第一分支语句，将所述计数器中的数值加一步长；

所述第一线程处理器，用于根据所述计数器的数值确定M*N个线程中需运行第一分支语句的线程的数量；

所述第一线程处理器，用于在确认所述数量大于阈值的情况下，执行线程同步以及线程数据重映射。

21.根据权利要求20所述的图形处理器，其特征在于，所述图形处理器还设置有标志位，所述标志位的值设置为第一标志值，所述第一标志值用于指示不执行重映射，其中，

所述第一线程处理器，用于在确定所述数量大于阈值的情况之前，读取所述标志位；且，

所述第一线程处理器，用于在确定所述数量大于阈值的情况之后及执行线程同步之前，将所述第一标志值设置为第二标志值，所述第二标志值用于指示需要执行重映射。

22.根据权利要求21所述的图形处理器，其特征在于，所述第一线程处理器，还用于在执行所述线程同步之后且执行所述线程数据重映射之前，将所述计数器中的数值清零。

23.根据权利要求21或22所述的图形处理器，其特征在于，所述第一线程束处理器包括第二线程处理器，所述第二线程处理器用以在运行N个线程中的一个以处理满足第一分支语句或满足第二分支语句的待处理数据，

第二线程处理器，用于读取所述标志位，在确认所述标志位的值为所述第二标志值时，执行线程同步以及线程数据重映射；

第二线程处理器，用于在确认所述标志位的值为所述第一标志值时，根据所述计数器的数值确定所述M*N个线程中需运行第一分支语句的线程的数量，在确认所述数量大于阈值的情况下，执行线程同步以及线程数据重映射。

24.根据权利要求23所述的图形处理器，其特征在于，所述第一线程处理器用以运行N个线程中的第一线程以处理满足第一分支语句的判断条件的的待处理数据，所述第二线程处理器用以运行N个线程中的第二线程以处理满足第二分支语句的判断条件的待处理数据，所述图形处理器还设置有一维数组、第一变量以及第二变量，其中所述一维数组的长度是M*N，所述第一变量的初始值是0，所述第二变量的初始值是M*N-1，所述第一线程处理器执行所述线程数据重映射，

所述第一线程处理器，用于读取所述第二变量的数值，并将所述第二线程的线程标识写入所述一维数组中以所述第二变量的数值作为下标的位置，将所述第二变量的数值减一，并执行所述线程同步；

所述第二线程处理器，用于读取所述第一变量的数值，并将所述第二线程的线程标识写入所述一维数组中以所述第一变量的数值作为下标的位置，将所述第一变量的数值加一，并执行所述线程同步；

所述第一线程处理器，用于在所述线程同步结束后，读取所述一维数组中以所述第一线程的线程标识作为下标的位置上的数值，并将读取的所述数值作为所述线程数据重映射产生的所述第一线程的更新的线程标识；

所述第二线程处理器，用于在所述线程同步结束后，读取所述一维数组中以所述第二线程的线程标识作为下标的位置上的数值，并将读取的所述数值作为所述线程数据重映射产生的所述第二线程的更新的线程标识。

25.根据权利要求24所述的图形处理器，其特征在于，所述第一线程处理器在执行所述线程同步之后且执行所述线程数据重映射之前，

所述第一线程处理器，用于以所述第一线程的线程标识作为索引将所述第一待处理数据记录在索引表，其中所述第一线程的线程标识与所述第一待处理数据具有一一对应关系，所述索引表记录有所述M*N个线程的线程标识与待处理数据之间的一一对应关系；

所述第一线程处理器，用于在执行所述线程数据重映射之后，以执行所述线程数据重映射后产生的所述第一线程的更新的线程标识作为索引在所述索引表中读取与所述第一线程的更新的线程标识对应的第三待处理数据；

所述第一线程处理器，用于在所述第三待处理数据满足第一分支语句的判断条件时，执行所述第一分支语句，在所述第三待处理数据满足第二分支语句的判断条件时，执行所述第二分支语句。

26.根据权利要求25所述的图形处理器，其特征在于，所述阈值为1。

27.根据权利要求25所述的图形处理器，其特征在于，所述阈值为大于等于2且小于等于5的正整数。

28.根据权利要求20至22任一项所述的图形处理器，其特征在于，所述第一线程处理器执行所述第一分支语句的概率小于所述第一线程处理器执行所述第二分支语句的概率。

29.一种图形处理器，所述图形处理器用于处理M个线程束，每个所述线程束包括N个线程，所述图形处理器还包括至少一个线程束处理器，所述至少一个线程束处理器中的第一线程束处理器包括N的整数倍个线程处理器，所述第一线程束处理器包括第一线程处理器，所述第一线程处理器运行循环语句，用以在一个循环中运行N个线程中的一个以处理满足第一分支语句的判断条件或满足第二分支语句的判断条件的待处理数据，所述图形处理器中设置有计数器，其中，

所述第一线程处理器，用于在第一循环中获取需要处理的第一待处理数据，确定所述第一待处理数据满足第一分支语句，将所述计数器中的数值加一步长；

所述第一线程处理器，用于根据所述计数器的数值确定M*N个线程中需运行第一分支语句的线程的数量；

所述第一线程处理器，用于在确认所述数量大于阈值的情况下，执行线程同步并将所述计数器中的数值清零；

所述第一线程处理器，用于执行线程数据重映射。

30.根据权利要求29所述的图形处理器，其特征在于，

所述第一线程处理器，用于获取在所述第一线程的第二循环中需要处理的第二待处理数据，确定所述第二待处理数据满足所述第二分支语句的判断条件，将所述计数器中的数值减一步长。

31.根据权利要求29或30所述的图形处理器，其特征在于，所述图形处理器还设置有标志位，所述标志位的值设置为第一标志值，所述第一标志值用于指示不执行重映射，

所述第一线程处理器，用于在确定所述数量大于阈值的情况之前，读取所述标志位；

所述第一线程处理器，用于在确定所述数量大于阈值的情况之后及执行线程同步之前，将所述第一标志值设置为第二标志值，所述第二标志值用于指示需要执行重映射。

32.根据权利要求31所述的图形处理器，其特征在于，所述第一线程束处理器包括第二线程处理器，所述第二线程处理器用以在运行N个线程中的一个以处理满足第一分支语句或满足第二分支语句的待处理数据，

第二线程处理器，用于读取所述标志位，在确认所述标志位的值为所述第二标志值时，执行线程同步以及线程数据重映射；

第二线程处理器，用于在确认所述标志位的值为所述第一标志值时，根据所述计数器的数值确定所述M*N个线程中需运行第一分支语句的线程的数量，在确认所述数量大于阈值的情况下，执行线程同步以及线程数据重映射。

33.根据权利要求32所述的图形处理器，其特征在于，所述第一线程处理器用以运行N个线程中的第一线程以处理满足第一分支语句的判断条件的的待处理数据，所述第二线程处理器用以运行N个线程中的第二线程以处理满足第二分支语句的判断条件的待处理数据，所述图形处理器还设置有一维数组、第一变量以及第二变量，其中所述一维数组的长度是M*N，所述第一变量的初始值是0，所述第二变量的初始值是M*N-1，所述第一线程处理器执行所述线程数据重映射，

所述第一线程处理器，用于读取所述第二变量的数值，并将所述第二线程的线程标识写入所述一维数组中以所述第二变量的数值作为下标的位置，将所述第二变量的数值减一，并执行所述线程同步；

所述第二线程处理器，用于读取所述第一变量的数值，并将所述第二线程的线程标识写入所述一维数组中以所述第一变量的数值作为下标的位置，将所述第一变量的数值加一，并执行所述线程同步；

所述第一线程处理器，用于在所述线程同步结束后，读取所述一维数组中以所述第一线程的线程标识作为下标的位置上的数值，并将读取的所述数值作为所述线程数据重映射产生的所述第一线程的更新的线程标识；

所述第二线程处理器，用于在所述线程同步结束后，读取所述一维数组中以所述第二线程的线程标识作为下标的位置上的数值，并将读取的所述数值作为所述线程数据重映射产生的所述第二线程的更新的线程标识。

34.根据权利要求33所述的图形处理器，其特征在于，所述第一线程处理器运行第一线程，

所述第一线程处理器，用于在执行所述线程同步之后且执行所述线程数据重映射之前，以所述第一线程的线程标识作为索引将所述第一待处理数据记录在索引表，其中，所述第一线程的线程标识与所述第一待处理数据具有一一对应关系；

所述第一线程处理器，用于在执行所述线程数据重映射之后，以执行所述线程数据重映射后产生的所述第一线程的更新的线程标识作为索引在所述索引表中读取与所述第一线程的更新的线程标识对应的第三待处理数据；

所述第一线程处理器，用于在所述第三待处理数据满足第一分支语句的判断条件时，执行所述第一分支语句，所述第一线程处理器在所述第三待处理数据满足第二分支语句的判断条件时，执行所述第二分支语句。

35.根据权利要求34所述的图形处理器，其特征在于，所述图形处理器中还记录有每个线程的循环变量，所述循环变量用于指示线程当前所在的循环的序号，所述索引表中记录有所述第一线程的循环变量与所述第一线程的线程标识、第一线程在所述循环变量所指示的循环中待处理数据的对应关系，

所述第一线程处理器，用于在执行所述线程数据重映射之后，执行所述线程数据重映射后产生的所述第一线程的更新的线程标识作为索引在所述索引表中读取与所述第一线程的更新的线程标识对应的循环变量；

所述第一线程处理器，用于在执行所述第一分支语句或所述第二分支语句之后，将所述第一线程的更新的线程标识对应的循环变量加一以获取更新的循环变量，且在所述更新的循环变量不符合所述循环语句规定的循环条件时，结束所述第一线程，在所述更新的循环变量符合所述循环语句规定的循环条件时，运行所述第一线程的第二循环。

36.根据权利要求35所述的图形处理器，其特征在于，所述阈值为1。

37.根据权利要求35所述的图形处理器，其特征在于，所述阈值为大于等于2且小于等于5的正整数。

38.根据权利要求29或30所述的图形处理器，其特征在于，所述第一线程处理器执行所述第一分支语句的概率小于所述第一线程处理器执行所述第二分支语句的概率。