CN101436121A - 用于使用并行处理器架构执行扫描操作的方法及设备 - Google Patents

用于使用并行处理器架构执行扫描操作的方法及设备 Download PDF

Info

Publication number
CN101436121A
CN101436121A CNA2008101727200A CN200810172720A CN101436121A CN 101436121 A CN101436121 A CN 101436121A CN A2008101727200 A CNA2008101727200 A CN A2008101727200A CN 200810172720 A CN200810172720 A CN 200810172720A CN 101436121 A CN101436121 A CN 101436121A
Authority
CN
China
Prior art keywords
scan operation
processor architecture
parallel processor
instruction
treatment element
Prior art date
Legal status (The legal status is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the status listed.)
Granted
Application number
CNA2008101727200A
Other languages
English (en)
Other versions
CN101436121B (zh
Inventor
迈克尔·J·加兰
萨姆拉·M·莱内
蒂莫·O·艾拉
戴维·帕特里克·吕布克
Current Assignee (The listed assignees may be inaccurate. Google has not performed a legal analysis and makes no representation or warranty as to the accuracy of the list.)
Nvidia Corp
Original Assignee
Nvidia Corp
Priority date (The priority date is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the date listed.)
Filing date
Publication date
Application filed by Nvidia Corp filed Critical Nvidia Corp
Publication of CN101436121A publication Critical patent/CN101436121A/zh
Application granted granted Critical
Publication of CN101436121B publication Critical patent/CN101436121B/zh
Expired - Fee Related legal-status Critical Current
Anticipated expiration legal-status Critical

Links

Images

Classifications

    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/30—Arrangements for executing machine instructions, e.g. instruction decode
    • G06F9/30003—Arrangements for executing specific machine instructions
    • G06F9/30007—Arrangements for executing specific machine instructions to perform operations on data operands
    • G06F9/3001—Arithmetic instructions
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F7/00—Methods or arrangements for processing data by operating upon the order or content of the data handled
    • G06F7/38—Methods or arrangements for performing computations using exclusively denominational number representation, e.g. using binary, ternary, decimal representation
    • G06F7/48—Methods or arrangements for performing computations using exclusively denominational number representation, e.g. using binary, ternary, decimal representation using non-contact-making devices, e.g. tube, solid state device; using unspecified devices
    • G06F7/50—Adding; Subtracting
    • G06F7/505—Adding; Subtracting in bit-parallel fashion, i.e. having a different digit-handling circuit for each denomination
    • G06F7/506—Adding; Subtracting in bit-parallel fashion, i.e. having a different digit-handling circuit for each denomination with simultaneous carry generation for, or propagation over, two or more stages
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/30—Arrangements for executing machine instructions, e.g. instruction decode
    • G06F9/30003—Arrangements for executing specific machine instructions
    • G06F9/30007—Arrangements for executing specific machine instructions to perform operations on data operands
    • G06F9/30018—Bit or string instructions
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/30—Arrangements for executing machine instructions, e.g. instruction decode
    • G06F9/30003—Arrangements for executing specific machine instructions
    • G06F9/30007—Arrangements for executing specific machine instructions to perform operations on data operands
    • G06F9/30036—Instructions to perform operations on packed data, e.g. vector, tile or matrix operations
    • G06F9/30038—Instructions to perform operations on packed data, e.g. vector, tile or matrix operations using a mask
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/30—Arrangements for executing machine instructions, e.g. instruction decode
    • G06F9/38—Concurrent instruction execution, e.g. pipeline or look ahead
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F2207/00—Indexing scheme relating to methods or arrangements for processing data by operating upon the order or content of the data handled
    • G06F2207/506—Indexing scheme relating to groups G06F7/506 - G06F7/508
    • G06F2207/5063—2-input gates, i.e. only using 2-input logical gates, e.g. binary carry look-ahead, e.g. Kogge-Stone or Ladner-Fischer adder

Landscapes

  • Engineering & Computer Science (AREA)
  • Theoretical Computer Science (AREA)
  • Software Systems (AREA)
  • Physics & Mathematics (AREA)
  • General Physics & Mathematics (AREA)
  • General Engineering & Computer Science (AREA)
  • Pure & Applied Mathematics (AREA)
  • Mathematical Analysis (AREA)
  • Computational Mathematics (AREA)
  • Mathematical Optimization (AREA)
  • Computing Systems (AREA)
  • Mathematical Physics (AREA)
  • Advance Control (AREA)
  • Multi Processors (AREA)
  • Executing Machine-Instructions (AREA)
  • Complex Calculations (AREA)
  • Image Processing (AREA)

Abstract

本发明提供一种用于使用并行处理架构对一序列的一位值执行扫描操作的***、方法和计算机程序产品。在操作中，接收扫描操作指令。另外，响应于所述扫描操作指令，使用具有多个处理元件的并行处理器架构对一序列的一位值执行扫描操作。

Description

用于使用并行处理器架构执行扫描操作的方法及设备
技术领域
本发明涉及扫描操作，且更特定而言，涉及使用并行处理架构执行扫描操作。
背景技术
并行处理器架构通常用来执行大量不同的计算算法。通常使用此架构执行的算法的实例是扫描操作(例如，“所有前缀和”操作等)。在表格1中定义了一个此扫描操作。
表格1
[ I , a 0 , ( a 0 ⊕ a 1 ) , . . . , ( a 0 ⊕ a 1 ⊕ . . . ⊕ a n - 1 ) ] ,
具体来说，在给出阵列[a0，a1，...，an-1]和运算符
Figure A200810172720D00042
(“I”为运算符
Figure A200810172720D00043
的单位元素)的情况下，返回表格1的阵列。例如，如果运算符
Figure A200810172720D00044
是加法运算符时，对阵列[3 1 7 0 4 1 6 3]执行所述扫描操作将返回[0 3 4 11 11 15 16 22]等。尽管在以上实例中阐述加法运算符，但此运算符可以是两个运算对象的任何结合运算符。
此外，所述扫描操作可以是互斥扫描操作(如在表格1中所显示)或相容扫描操作。所述互斥扫描是指结果的每一元素j是输入阵列中一直到(但不包含)元素j的所有元素的和。另一方面，在相容扫描中，求包含元素j的所有元素的和。
到目前为止，继续存在对使用并行处理器架构更有效地执行计算算法(例如扫描操作)的需要。
发明内容
提供一种用于使用并行处理架构对一序列的一位值执行扫描操作的***、方法和计算机程序产品。在操作中，接收扫描操作指令。另外，响应于所述扫描操作指令，使用具有多个处理元件的并行处理器架构对一序列的一位值执行扫描操作。
附图说明
图1显示根据本发明的一个实施例用于使用并行处理架构对一序列的一位值执行扫描操作的方法。
图2显示根据本发明的一个实施例用于对一序列的一位值执行扫描操作的***。
图3显示根据本发明的一个实施例用于对一序列的一位值执行扫描操作的***的结果。
图4显示根据本发明的一个实施例用于使用并行处理架构在硬件中执行扫描操作的***。
图5显示根据本发明的又一实施例用于使用并行处理架构在硬件中执行扫描操作的***。
图6显示根据本发明的另一实施例用于使用并行处理架构在硬件中执行扫描操作的***。
图7图解说明其中可实施各种先前实施例的各种架构及/或功能性的实例性***。
具体实施方式
图1显示根据本发明的一个实施例用于使用并行处理架构对一位值执行扫描操作的方法100。如所显示，接收扫描操作指令。参见操作102。在本说明的背景中，扫描操作指令是指对应于扫描操作的任何指令或命令。
另外，响应于所述扫描操作指令，使用具有多个处理元件的并行处理器架构对一序列的一位值执行扫描操作。参见操作104。在本说明的背景下，处理元件是指并行处理器架构的任何组件。另外，所述序列的一位值可包含任何序列的一位值。通过此设计，在某些实施例中，可更有效地对一位输入执行计算算法，例如扫描操作。
此外，在本说明的背景下，所述扫描操作可以指涉及阵列的当前元素和至少一个先前元素的任何操作。例如，在各种实施例中，扫描操作可包含前缀和扫描操作、互斥扫描操作、相容扫描操作、及/或任何其他扫描操作(例如，涉及更多或更少的元素及/或其他运算符等)。
而且，在本说明的背景下，所述并行处理器架构可包含任何包含并行操作的两个或两个以上处理元件的架构。在一个实施例中，此并行处理器架构可采取图形处理器(例如，图形处理单元(GPU)等)、或具有图形处理能力的任何其他集成电路(例如，呈芯片组、芯片上***(SOC)、与CPU、离散处理器等整合在一起的核心的形式)的形式。在又一实施例中，前述并行处理架构可包含向量处理器。
现在将陈述关于各种可选架构和特征的更多说明性信息，按照用户的期望，可用或可不用所述各种可选架构和特征来实施前述框架。应极其注意，出于说明目的阐述以下信息，而不应将其视为以任何方式加以限制。可视需要在排除或不排除所述的其它特征的情况下并入任一以下特征。
图2显示根据本发明的一个实施例用于对一序列的一位值执行扫描操作的***200。作为一种选择，可实施本***来执行图1的方法。然而，本***当然可在任一所需环境中实施。还应注意，前述定义可应用于本说明期间。
如所显示，提供并行处理架构202。此并行处理架构包含多个并行处理器204。尽管未显示，但是此并行处理器可能够对预定数目的线程进行操作。为此目的，所述并行处理器中的每一者可并行地操作，同时对应的线程也可并行地操作。
在一个实施例中，所述并行处理架构可包含一个或一个以上单指令多数据(SIMD)处理元件。在此***中，可将处理器正执行的线程聚集为群组以便在任何时刻单个群组内的所有线程均精确地执行相同指令但在可能不同的数据上。在一个实施例中，以此方式操作的此线程群组可以称为“卷绕”。此外，在此群组中线程的预定数目可称为对应处理器的“卷绕大小”。
在另一实施例中，前述并行处理架构可包含图形处理器或具有图形处理能力的任何其他集成电路[例如，呈芯片组、芯片上***(SOC)、与CPU、离散处理器等整合在一起的核心的形式]。在又一实施例中，前述并行处理架构可包含具有一个或一个以上向量处理元件的处理器，例如，单元(Cell)处理器，是指由
Figure A200810172720D00061
和联合开发的单元宽带引擎(Cell Broadband Engine)微处理器架构。
继续参照图2，所述并行处理架构可包含本地共享存储器206。并行处理架构的并行处理器中的每一者可对其自身的本地共享存储器进行读取及/或写入操作。此共享存储器可由与每一处理器相关联的物理上分开的存储器组成或其可由在处理器之间共享的一个或一个以上存储器的分开分配的区域组成。此外，在所图解说明的实施例中，共享存储器可包含在并行处理架构的处理器包含在其上的集成电路上。
而且，显示包含全局存储器208。使用中，并行处理架构的所有处理器可存取此全局存储器。如所显示，此全局存储器可包含在集成电路上，所述集成电路与上述并行处理架构的处理器包含在其上的集成电路是分开的。尽管显示并行处理架构以特定的方式包含在图2的各种集成电路上，但是应注意如所需要，***组件可以或可以不包含在同一集成电路上。
而且，如所需，图2的本***可进一步包含驱动器210用来控制所述并行处理架构。在一个实施例中，驱动器可包含程序库用于促进此控制。例如，此程序库可包含可例示本文所阐述的功能性的程序库调用。
此外，在另一实施例中，驱动器可能够利用并行处理架构(例如，图形处理器等)提供通用计算能力。可结合由NVIDIA公司提供的CUDATM架构来提供此驱动器的实例。使用中，所述驱动器可用来控制并行处理架构以根据图1的方法进行操作。
图3显示根据本发明的一个实施例用于使用并行处理架构对一位输入执行扫描操作的***300的结果。作为一种选择，可以图1-2中的细节为背景来实施本***。然而，本***当然可在任一所需环境中实施。还应注意，前述定义可应用于本说明期间。
如所显示，提供包含为并行处理器架构的一部分的多个处理元件302。所述处理元件(例如，线程)各自处理1位值304。在一个实施例中，所述1位值可从对逻辑表述的评估中导出。在此情况中，1位值可以称为谓词位。
在操作中，可由并行处理器架构来接收扫描操作指令。在此情况下，扫描可包含前缀和扫描操作指令。响应于所述扫描操作指令，可使用具有多个处理元件的并行处理器架构执行前缀和扫描操作指令。
跨越N个处理元件的群组(即，卷绕)对谓词位输入执行前缀和扫描操作(在图式的实例中为互斥扫描)的结果导致对数(N)位的整数。图3显示N＝16处理元件(例如，线程)的卷绕的扫描结果306。当然，在各种实施例中可利用任何数目的处理元件。应注意，传送给处理元件“i”的值是具有其给定谓词位为1的较小索引的处理元件(例如，线程)数。在各种实施例中，此操作可用作许多计算核(例如串流压缩和基数分类)的基础。
在某些情况下，完全一般的扫描操作可不适合直接硬件实施方案。例如，扫描操作可涉及处理任意长度的序列和许多可能的数的类型(例如，整数型、浮点型、短型等)。相反，固定长度的小序列的二进制扫描原语可在硬件中实施且可作为机器指令提供。多处理器中的处理元件的数目是已知的架构常数，且数的类型可保持恒定为1位值。
图4显示根据本发明的一个实施例用于使用并行处理架构在硬件中执行扫描操作的***400。作为一种选择，可以图1-3的细节为背景实施本***。然而，本***当然可在任一所需环境中实施。而且，前述定义可应用于本说明期间。
如所显示，提供包含为并行处理器架构的一部分的多个处理元件402。另外，包含多个加法器404。此加法器可包含能够相加数的任何电路或装置。
在操作中，处理元件(例如，线程)可各自持有1位值。因此，当多个处理元件接收到扫描操作指令时，可使用具有多个处理元件的并行处理器架构来执行所述扫描操作指令。在此情况中，加法器404的聚集形成加法网络(例如，电路)，其接受来自处理元件402的每一者的1位输入值并将扫描操作的结果传送到处理元件406的每一者。
尽管以16个处理元件来图解说明图4，但应注意可利用任何数目的处理元件。另外，将图4中的***图解说明为执行互斥扫描的***。在另一实施例中，***可经配置以执行相容扫描。
此外，图4的***经配置具有等于处理元件的数目(N)的深度。在各种其他实施例中，所述***可经配置以最小化所述深度。可利用任何数目的技术实现此最小化。
图5显示根据本发明的另一实施例用于使用并行处理架构在硬件中执行扫描操作的***500。作为一种选择，可以图1-4的细节为背景来实施本***。然而，本***当然可在任一所需环境中实施。还应注意，前述定义可应用于本说明期间。
如所显示，提供包含为并行处理器架构的一部分的多个处理元件502。另外，包含加法器504树。在操作中，每一处理元件502贡献1位输入。
作为一选择，此1位输入可从指定的谓词寄存器获得。可通过所述加法器树来馈送这些输入，从而将前缀和值506作为输出传送到对应的处理元件。在一个实施例中，每一输出可存放在每一处理元件的指定数据寄存器中。
如所显示，由加法器504树形成的加法***具有深度值对数(N)，其中N是处理元件的数目。然而，在某些情况下，可需要减少***中加法器的数目。因此，可利用具有减少数目的加法器和增加的算法长度的***。
图6显示根据本发明的又一实施例用于使用并行处理架构在硬件中执行扫描操作的***600。作为一选择，可以图1-5的细节为背景来实施本***。然而，本***当然可在任一所需环境中实施。还应注意，前述定义可应用于本说明期间。
如所显示，提供包含为并行处理器架构的一部分的多个处理元件602。另外，包含多个加法器604。操作中，每一处理元件贡献1位输入。
应注意***的长度直接与***的延时相关。因此，如果***的总面积比总延时更重要，那么可需要具有少数目的加法器的***(例如，图6的***)。另一方面，如果延时比总面积更重要，那么可需要具有较多数目的加法器和较低长度的***(例如，图5的***)。
利用任一实施方案，扫描1位输入可比扫描一般数容易得多。例如，如果对全部32位整数求和，那么***中执行求和的加法器中的每一者将必须是32位加法器。然而，在1位输入的情况下，每一加法器的宽度为最多对数(N)，其中N是***中处理元件的数目。在本说明的背景中，加法器的宽度是指所述加法器能够处理的输入数可包含的位的最大数目。
在图6的特定情况和背景中，每一加法器将遇见每输入最多4个位。在一个实施例中，可在加法器树的不同层级处利用具有不同宽度的加法器。例如，在树的第一层级606(即，就在输入下面)中的加法器可包含仅1位输入。另外，第二层级608可包含仅2位输入。
在给出如以图2-6为背景所述的数据路径的情况下，跨越SIMD多处理器的处理元件的二进制扫描可作为机器指令暴露于程序。在一个实施例中，可利用谓词扫描指令(“PSCAN”)，其从每一处理元件提取寄存器(“Rpred”)中的1位谓词作为输入且返回另一寄存器(“Rsum”)中的适合前缀和到每一处理元件。在以下表格2中显示此指令。
表格2
PSCAN Rsum，Rpred
此指令的操作直接对应于图2-6的***。处理元件中的每一者向***的并行前缀加法网络的输入贡献谓词位，且每一者接收单个输出值。
大多数多处理器硬件并入用于在计算期间有选择地停用处理元件的机制。通常进行此操作来允许名义上的SIMD处理器阵列执行程序的发散路径。在此情形下，当活动的处理元件执行“PSCAN”指令时可假设停用的处理元件向并行前缀计算贡献“0”。然而，在另一实施例中，可提供所述指令的变体，其中不活动处理元件贡献“1”。
此外，尽管以加法操作为背景描述了图2-6，但其他操作可同样适用。例如，可对扫描操作和加法器进行概括以使用除加法之外的任何结合操作。因此，可利用并行处理器架构的多个功能单元来执行扫描操作。
在此情况下，功能单元可包含加法器、布尔(Boolean)逻辑运算符、算术和逻辑运算符以及各种其他功能单元。此外，如所显示，并行处理器架构可包含多个功能单元层级。在此情况下，层级的数目可以少于处理元件的数目。此外，层级的数目可通常少于处理元件的数目的对数。
在机器指令的背景下，可类似于加法指令来利用例如“与”、“或”及“异或”的指令。另外，对于1位输入，可减少例如最小、最大和相乘等的操作到所述3个前述1位操作。如以上所述，此指令的数据路径看起来与针对图3-6显示的数据路径相同，其中构成的加法器区块由适当的“与”/“或”/“异或”门代替。另外，在一个实例性实施例中，以图3-6为背景所述的***可以管道配置来实施。在此情况下，可利用锁存器来实施此管道配置。
应注意，可利用各种计算机编程语言(例如，C、C++等)来实施对应于扫描操作指令的机器指令。在一个实施例中，利用例如统一计算装置架构(Compute UnifiedDevice Architecture)(CUDATM)C作为简单内在的语言来执行所述指令。例如，表格3显示CUDATM C中的指令，其中“i”表示线程索引。
表格3
int sum_i＝PSCAN(A[i]<pivot)；
暴露此功能性的另一方法是对处理元件的“活动”位而非由程序明确地计算的谓词暗中地执行二进制前缀和。在以下表格4中显示此构造的实例。
表格4
if(A[i]<pivot)
{
   sum_i＝PSCAN_active()；
}
在此情况下，基础处理器机制可存在以供编译器利用来存取多处理器的“活动”状态。
当然，此仅是暴露较高级语言形式的原语的一个可能方法且具体来说，其与CUDATM C相关。考虑到暴露原语机器支持的其他方法。应注意具有大致不同设计的语言(例如，数据并行C等)将利用不同语言层级的实施例。
在一个实施例中，可在协作线程阵列(Cooperative Thread Array，CTA)中一起执行一个或一个以上处理元件或线程群组(例如，卷绕)。因此，并行处理器架构可提供处理元件之间的协调。在此情况下，协调可包含关于所写入的结果的目的地的协调。在一个实施例中，多个处理元件可能够经由芯片上共享存储器彼此通信且经由势垒同步化。
当跨越由多个线程组成的CTA执行扫描时，可执行两个扫描层级。第一扫描可发生在每一卷绕内。作为一选择，第一扫描可用如上所述的“PSCAN”原语来执行。第二扫描可从每一卷绕接收单个值，且对所述部分和执行扫描。应注意在卷绕宽度为32的情况下，其全部为5位整数。
在一个实施例中，可利用1位扫描原语来通过对每一二进制数字独立地执行所述扫描且然后对结果求和来计算多位数的前缀和。换句话说，并行处理器架构可通过对多位值中的个别位个别地执行扫描且在个别扫描的结果进行位移位之后对所述结果求和来对多位值执行所述扫描操作。例如，假定卷绕中的每一线程被赋予5位值“x_i”。可如表格5中所示计算所述值的前缀和。
表格5
int sum_i＝PSCAN(x_i & 1)；
   sum_i+＝PSCAN(x_i & 2)<<1；
   sum_i+＝PSCAN(x_i & 4)<<2；
   sum_i+＝PSCAN(x_i & 8)<<3；
   sum_i+＝PSCAN(x_i & 16)<<4；
此实施方案的结果将与具有全部扫描核的实施方案相同。然而，假设“PSCAN”利用单个指令来执行，那么当输入值中的位数目较小时此可比全部核更有效。关于扫描核的更多信息可在序列号为11/862,938标题为“用于执行扫描操作的***、方法和计算机程序产品(SYSTEM，METHOD AND COMPUTER PROGRAM PRODUCT FORPERFORMING A SCAN OPERATION)”的专利申请案中找到，所述申请案的全文以引用的方式并入本文中。
应注意可在任何所需环境(包含并行处理架构)中利用以上功能性且可在需要构造有效并行核的各种情形中实施。例如，假定维持对应于数据的项目对列且一卷绕线程将每线程最多1项目写入到所述队列中。如果每个线程总是写入1个项目，那么每一线程将始终提前知晓应将队列指针的多少偏移作为值写入。
然而，如果每一个别线程选择是否写入一值，那么卷绕中的所有线程必须计算写入其值的适当偏移。可使用对确定是否每一线程希望写入的谓词的扫描来实施计算此偏移。可使用如表格6中所图解说明的二进制扫描原语来简单且有效地表述此计算。
表格6
__device__void maybe_write(int*queue，int x，bool should_write)
{
   unsigned int i＝PSCAN(should_write)；
   if(should_write) queue[i]＝x；
}
可通过跨越卷绕暗中地扫描处理器“活动”位产生更压缩的变体。例如，在以下表格7中显示一个此变体。
表格7
__device__void maybe_write(int*queue，int x，bool should_write)
{
   if(should_write) queue[PSCAN_active()]＝x；
}
作为另一实例，线程的CTA可用每线程一个值来控制一序列的数。在此实例中，可选择“主元”值且可重新混洗一阵列以使得所述阵列中小于主元的所有值在所有其他数之前。例如，此是在例如快速分类(Quicksort)算法中的步骤。
为实施此操作，可定义接受谓词“p”的“秩()”原语。谓词为真的线程将接收具有谓词为真的较低线程索引的线程数目的计数。谓词为假的线程将接收具有谓词为假的较低线程索引的线程数目加上真谓词的总数目的计数。表格8显示CUDATM中代表函数的实例，其中函数“cta_prefix_sum()”以在07年9月27日提出申请的标题为“用于执行扫描操作的***、方法和计算机程序产品(SYSTEM，METHOD ANDCOMPUTER PROGRAM PRODUCT FOR PERFORMING A SCAN OPERATION)”的专利申请案号11/862,938中阐述的方式构建于内部卷绕扫描之上
表格8
__device__unsigned int rank(bool p)
{
  __shared__bool smem[ctasize]；
  smem[threadIdx.x]＝p；
  __syncthreads()；
  bool last_p＝smem[ctasize-1]；//Everybody gets last value
  __syncthreads()；
  cta_prefix_sum(smem)；        //Uses PSCAN.See also P003535.
  //(1)total number of True threads
  unsigned int ntrue＝last_p+smem[ctasize-1]；
  //(2)Compute this thread′s rank within ordering
  unsigned int r＝(p)？smem[threadIdx.x]
  :ntrue+threadIdx.x-smem[threadIdx.x]；
  return r；
}
在给出此原语的情况下，可写入分割函数。例如，表格9显示一个此分割函数的实例。
表格9
__global__void partition(unsigned int*v，const unsigned int pivot)
{
  unsigned int v_i＝v[threadIdx.x]；
  __syncthreads()；//make sure everyone is ready to write
  unsigned int j＝rank(v_i<pivot)；
  v[j]＝v_i；
}
类似于分割，分类数字序列是可用于许多应用中的另一操作。其也可容易地按照以上定义的“秩()”原语来实施。每一遍基数分类均仅是以基于数据值的一位的值而非基于比较谓词的“分割()”的方式进行重新混洗。在本说明的背景下，基数分类是通过处理个别数字分类整数的分类算法。在表格10中显示利用基数分类的实施方案的一个实例。
表格10
__device__void cta_radix_sort(unsigned int*v)
{
   for(unsigned int shift＝0；shift<32；++shift)
   {
       unsigned int v_i＝v[threadIdx.x]；
       __syncthreads()；
       unsigned int lsb＝(v_i>>shift)& 0x1；
       unsigned int r＝rank(!lsb)；
       v[r]＝v_i；
       __syncthreads()；//make sure everyone wrote
   }
}
尽管上文已描述各种实施例，但应了解，所述实施例仅以实例的方式而非限制的方式呈现。例如，在各种其他实施例中，可在前述图式的背景和细节中利用并实施任何数目的扫描算法。
图7图解说明其中可实施各种先前实施例的各种构架及/或功能性的实例性***700。如所显示，提供包含至少一个主处理器701的***，其中主处理器701连接到通信总线702。***还包含主存储器704。控制逻辑(软件)及数据存储在主存储器中，所述主存储器可采取随机存取存储器(RAM)的形式。
所述***还包含图形处理器706和显示器708，即计算机监视器。在一个实施例中，图形处理器可包含多个着色器模块、光栅化模块等。每一前述模块甚至可布置在单个半导体平台上以形成图形处理单元(GPU)。
在本说明中，单个半导体平台可指单独整体式基于半导体的集成电路或芯片。应注意，术语单个半导体平台还可指具有增强连接性的多芯片模块，其模拟芯片上操作且对利用常规中央处理单元(CPU)和总线实施方案做出显著改进。当然，还可按照用户的需要，单独地或者以半导体平台的各种组合形式布置各种模块。
所述***可还包含辅助存储装置710。辅助存储装置包含(例如)硬盘驱动及/或可装卸存储装置驱动(其代表软盘驱动、磁带驱动、光盘驱动等)。可装卸存储装置驱动以众所周知的方式从可装卸存储单元读取及/或向可装卸存储单元写入。
计算机程序或计算机控制逻辑算法可存储在主存储器及/或辅助存储装置中。在执行时，此计算机程序使***能够执行各种功能。存储器、存储装置及/或任何其它存储装置均为计算机可读媒体的可能实例。
在一个实施例中，各种先前图示的架构及/或功能性可以主处理器、图形处理器、能够实现主处理器和图形处理器二者的能力的至少一部分的集成电路(未显示)、芯片组(即设计用于作为执行相关功能的单元等来工作及出售的集成电路群组)及/或任何其他用于所述事项的集成电路为背景来实施。此外，在一个可能实施例中，各种先前图式的元件指派功能性可在驱动器712的控制下实施于前述集成电路中的任一者中。
而且，各种先前图示的架构及/或功能性可以通用计算机***、电路板***、专用于娱乐目的的游戏控制台***、专用***、及/或任何其他所需***为背景来实施。例如，所述***可采取桌上型计算机、膝上型计算机、及/或任何其他类型的逻辑的形式。而且，所述***可采取各种其他装置的形式，其包含但不限于个人数字助理(PDA)装置、移动电话装置、电视等。
此外，尽管未显示，但所述***可出于通信目的而耦合到网络(例如，电信网络、局域网络(LAN)、无线网络、例如因特网等的广域网络(WAN)、点对点网络、电缆网络等)。
尽管上文已描述各种实施例，但应了解，所述实施例仅以实例的方式而非限制的方式呈现。因此，优选的实施例的广度和范围不应受限于上文所述实例性实施例的任一者，而应仅根据所附权利要求书及其等效内容来界定。

Claims (20)

1、一种方法，其包括：
接收扫描操作指令；及
响应于所述扫描操作指令，使用具有多个处理元件的并行处理器架构对一序列的一位值执行扫描操作。
2、如权利要求1所述的方法，其中所述扫描操作包含前缀和扫描操作。
3、如权利要求1所述的方法，其中所述扫描操作包含相容扫描操作。
4、如权利要求1所述的方法，其中所述扫描操作包含互斥扫描操作。
5、如权利要求1所述的方法，其中所述并行处理器架构提供所述处理元件之间的协调。
6、如权利要求5所述的方法，其中所述协调包含关于所写入的结果的目的地的协调。
7、如权利要求1所述的方法，其中所述处理元件各自并行执行多个线程。
8、如权利要求1所述的方法，其中利用所述并行处理器架构的多个功能单元执行所述扫描操作。
9、如权利要求8所述的方法，其中所述功能单元包含加法器。
10、如权利要求8所述的方法，其中所述功能单元包含布尔逻辑运算符。
11、如权利要求8所述的方法，其中所述功能单元包含算术和逻辑运算符。
12、如权利要求8所述的方法，其中所述并行处理器架构包含多个层级的功能单元。
13、如权利要求12所述的方法，其中所述层级的数目少于所述处理元件的数目。
14、如权利要求12所述的方法，其中所述层级的数目少于所述处理元件的数目的对数。
15、如权利要求1所述的方法，其中所述并行处理器架构通过个别地执行对多位值的个别位的扫描且在所述个别扫描的结果进行位移位之后对所述结果求和来执行对多位值的所述扫描操作。
16、如权利要求1所述的方法，其中所述并行处理器架构包含一个或一个以上单指令多数据处理器。
17、如权利要求1所述的方法，其中所述并行处理器架构包含图形处理器。
18、一种包含在计算机可读媒体上的计算机程序产品，所述计算机程序产品包括：
用于响应于扫描操作指令使用具有多个处理元件的并行处理器架构对一序列的一位值执行扫描操作的计算机代码。
19、一种设备，其包括：
并行处理器架构，其包含多个处理元件；及
指令，其用于使用所述并行处理器架构对一序列的一位值执行扫描操作。
20、如权利要求19所述的设备，其中所述并行处理器架构经由总线保持与存储器和显示器的通信。
CN200810172720.0A 2007-11-15 2008-11-11 用于使用并行处理器架构执行扫描操作的方法及设备 Expired - Fee Related CN101436121B (zh)

Applications Claiming Priority (2)

Application Number Priority Date Filing Date Title
US11/940,782 US8661226B2 (en) 2007-11-15 2007-11-15 System, method, and computer program product for performing a scan operation on a sequence of single-bit values using a parallel processor architecture
US11/940,782 2007-11-15

Publications (2)

Publication Number Publication Date
CN101436121A true CN101436121A (zh) 2009-05-20
CN101436121B CN101436121B (zh) 2014-12-17

Family

ID=40643242

Family Applications (1)

Application Number Title Priority Date Filing Date
CN200810172720.0A Expired - Fee Related CN101436121B (zh) 2007-11-15 2008-11-11 用于使用并行处理器架构执行扫描操作的方法及设备

Country Status (5)

Country Link
US (1) US8661226B2 (zh)
JP (1) JP2009169935A (zh)
KR (1) KR20090050977A (zh)
CN (1) CN101436121B (zh)
TW (1) TWI398779B (zh)

Cited By (11)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN104040482A (zh) * 2011-12-28 2014-09-10 英特尔公司 用于在打包数据元素上执行增量解码的***、装置和方法
CN104081337A (zh) * 2011-12-23 2014-10-01 英特尔公司 用于响应于单个指令来执行横向部分求和的***、装置和方法
CN104813279A (zh) * 2012-12-28 2015-07-29 英特尔公司 用于减少具有步幅式访问模式的向量寄存器中的元素的指令
WO2015161826A1 (en) * 2014-04-24 2015-10-29 Mediatek Inc. Cpu control method, electronic system control method and electronic system
CN105453028A (zh) * 2013-08-14 2016-03-30 高通股份有限公司 向量积累方法及设备
CN107003988A (zh) * 2014-12-19 2017-08-01 英特尔公司 用于执行卷积运算的存储设备和方法
US9965282B2 (en) 2011-12-28 2018-05-08 Intel Corporation Systems, apparatuses, and methods for performing delta encoding on packed data elements
CN112256580A (zh) * 2020-10-23 2021-01-22 济南浪潮数据技术有限公司 一种代码扫描方法、装置、设备及存储介质
WO2022089421A1 (en) * 2020-10-30 2022-05-05 International Business Machines Corporation ReLU COMPRESSION TO REDUCE GPU MEMORY
CN114510271A (zh) * 2022-02-09 2022-05-17 海飞科(南京)信息技术有限公司 用于在单指令多线程计算***中加载数据的方法和装置
CN115718724A (zh) * 2023-01-09 2023-02-28 阿里巴巴（中国）有限公司 Gpu、数据选择方法及芯片

Families Citing this family (29)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US8996846B2 (en) 2007-09-27 2015-03-31 Nvidia Corporation System, method and computer program product for performing a scan operation
US8264484B1 (en) 2007-10-29 2012-09-11 Nvidia Corporation System, method, and computer program product for organizing a plurality of rays utilizing a bounding volume
US8284188B1 (en) 2007-10-29 2012-10-09 Nvidia Corporation Ray tracing system, method, and computer program product for simultaneously traversing a hierarchy of rays and a hierarchy of objects
US8773422B1 (en) 2007-12-04 2014-07-08 Nvidia Corporation System, method, and computer program product for grouping linearly ordered primitives
US8243083B1 (en) 2007-12-04 2012-08-14 Nvidia Corporation System, method, and computer program product for converting a scan algorithm to a segmented scan algorithm in an operator-independent manner
US8417735B1 (en) * 2007-12-12 2013-04-09 Nvidia Corporation Instruction-efficient algorithm for parallel scan using initialized memory regions to replace conditional statements
US8502819B1 (en) 2007-12-17 2013-08-06 Nvidia Corporation System and method for performing ray tracing node traversal in image rendering
US8289324B1 (en) 2007-12-17 2012-10-16 Nvidia Corporation System, method, and computer program product for spatial hierarchy traversal
US8200947B1 (en) * 2008-03-24 2012-06-12 Nvidia Corporation Systems and methods for voting among parallel threads
US8527742B2 (en) * 2008-08-15 2013-09-03 Apple Inc. Processing vectors using wrapping add and subtract instructions in the macroscalar architecture
US8356159B2 (en) * 2008-08-15 2013-01-15 Apple Inc. Break, pre-break, and remaining instructions for processing vectors
US9335997B2 (en) 2008-08-15 2016-05-10 Apple Inc. Processing vectors using a wrapping rotate previous instruction in the macroscalar architecture
US9342304B2 (en) 2008-08-15 2016-05-17 Apple Inc. Processing vectors using wrapping increment and decrement instructions in the macroscalar architecture
US9335980B2 (en) 2008-08-15 2016-05-10 Apple Inc. Processing vectors using wrapping propagate instructions in the macroscalar architecture
US8321492B1 (en) 2008-12-11 2012-11-27 Nvidia Corporation System, method, and computer program product for converting a reduction algorithm to a segmented reduction algorithm
US8327119B2 (en) * 2009-07-15 2012-12-04 Via Technologies, Inc. Apparatus and method for executing fast bit scan forward/reverse (BSR/BSF) instructions
US8564616B1 (en) 2009-07-17 2013-10-22 Nvidia Corporation Cull before vertex attribute fetch and vertex lighting
US8542247B1 (en) 2009-07-17 2013-09-24 Nvidia Corporation Cull before vertex attribute fetch and vertex lighting
KR101109717B1 (ko) * 2009-10-12 2012-01-31 주식회사 씨네스티 원격 제어를 위한 영상 처리 방법 및 이를 실행하기 위한 원격 제어 프로그램이 기록된 컴퓨터가 판독 가능한 기록매체
US8976195B1 (en) 2009-10-14 2015-03-10 Nvidia Corporation Generating clip state for a batch of vertices
US8384736B1 (en) 2009-10-14 2013-02-26 Nvidia Corporation Generating clip state for a batch of vertices
US9389860B2 (en) 2012-04-02 2016-07-12 Apple Inc. Prediction optimizations for Macroscalar vector partitioning loops
US9442755B2 (en) * 2013-03-15 2016-09-13 Nvidia Corporation System and method for hardware scheduling of indexed barriers
US9817663B2 (en) 2013-03-19 2017-11-14 Apple Inc. Enhanced Macroscalar predicate operations
US9348589B2 (en) 2013-03-19 2016-05-24 Apple Inc. Enhanced predicate registers having predicates corresponding to element widths
US9652284B2 (en) * 2013-10-01 2017-05-16 Qualcomm Incorporated GPU divergence barrier
US20180005346A1 (en) * 2016-07-01 2018-01-04 Google Inc. Core Processes For Block Operations On An Image Processor Having A Two-Dimensional Execution Lane Array and A Two-Dimensional Shift Register
US20180007302A1 (en) 2016-07-01 2018-01-04 Google Inc. Block Operations For An Image Processor Having A Two-Dimensional Execution Lane Array and A Two-Dimensional Shift Register
US10296342B2 (en) * 2016-07-02 2019-05-21 Intel Corporation Systems, apparatuses, and methods for cumulative summation

Citations (3)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN1202651A (zh) * 1996-12-20 1998-12-23 索尼公司 算术机逻辑单元的运算方法,存贮介质和算术和逻辑单元
US20030081833A1 (en) * 2001-04-23 2003-05-01 National Aeronautics And Space Administration Method for implementation of recursive hierarchical segmentation on parallel computers
CN1431588A (zh) * 2002-01-08 2003-07-23 北京南思达科技发展有限公司 一种逻辑可重组电路

Family Cites Families (52)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US4628483A (en) 1982-06-03 1986-12-09 Nelson Raymond J One level sorting network
US4855937A (en) 1984-08-08 1989-08-08 General Electric Company Data block processing for fast image generation
US5193207A (en) 1990-05-31 1993-03-09 Hughes Aircraft Company Link sorted memory
US5274718A (en) 1991-09-30 1993-12-28 At&T Bell Laboratories Image representation using tree-like structures
US5361385A (en) 1992-08-26 1994-11-01 Reuven Bakalash Parallel computing system for volumetric modeling, data processing and visualization
JP3202074B2 (ja) 1992-10-21 2001-08-27 富士通株式会社 並列ソート方式
JPH06223198A (ja) 1993-01-26 1994-08-12 Hitachi Ltd 光線追跡による画像生成装置及び方法
US5650862A (en) 1994-07-25 1997-07-22 Canon Kabushiki Kaisha Image signal processing apparatus with improved pixel parallel/line parallel conversion
US5793379A (en) * 1995-04-03 1998-08-11 Nvidia Corporation Method and apparatus for scaling images having a plurality of scan lines of pixel data
US6065005A (en) 1997-12-17 2000-05-16 International Business Machines Corporation Data sorting
WO2000004494A1 (en) 1998-07-17 2000-01-27 Intergraph Corporation Graphics processing system with multiple strip breakers
WO2000011607A1 (en) 1998-08-20 2000-03-02 Apple Computer, Inc. Deferred shading graphics pipeline processor
US6549907B1 (en) 1999-04-22 2003-04-15 Microsoft Corporation Multi-dimensional database and data cube compression for aggregate query support on numeric dimensions
US6489955B1 (en) 1999-06-07 2002-12-03 Intel Corporation Ray intersection reduction using directionally classified target lists
SG93211A1 (en) 1999-07-28 2002-12-17 Univ Singapore Method and apparatus for generating atomic parts of graphic representation through skeletonization for interactive visualization applications
US6556200B1 (en) 1999-09-01 2003-04-29 Mitsubishi Electric Research Laboratories, Inc. Temporal and spatial coherent ray tracing for rendering scenes with sampled and geometry data
US6738518B1 (en) 2000-05-12 2004-05-18 Xerox Corporation Document image decoding using text line column-based heuristic scoring
US7495664B2 (en) 2000-06-19 2009-02-24 Mental Images Gmbh Instant ray tracing
US6633882B1 (en) 2000-06-29 2003-10-14 Microsoft Corporation Multi-dimensional database record compression utilizing optimized cluster models
US7580927B1 (en) 2001-05-29 2009-08-25 Oracle International Corporation Quadtree center tile/boundary tile optimization
US6879980B1 (en) 2001-06-29 2005-04-12 Oracle International Corporation Nearest neighbor query processing in a linear quadtree spatial index
US7024414B2 (en) 2001-08-06 2006-04-04 Sensage, Inc. Storage of row-column data
US6943796B2 (en) 2002-07-22 2005-09-13 Sun Microsystems, Inc. Method of maintaining continuity of sample jitter pattern across clustered graphics accelerators
US7194125B2 (en) 2002-12-13 2007-03-20 Mitsubishi Electric Research Laboratories, Inc. System and method for interactively rendering objects with surface light fields and view-dependent opacity
US7146486B1 (en) 2003-01-29 2006-12-05 S3 Graphics Co., Ltd. SIMD processor with scalar arithmetic logic units
US20050177564A1 (en) 2003-03-13 2005-08-11 Fujitsu Limited Server, method, computer product, and terminal device for searching item data
US7403944B2 (en) 2003-04-21 2008-07-22 Gerald Budd Reduced comparison coordinate-value sorting process
US7418585B2 (en) 2003-08-28 2008-08-26 Mips Technologies, Inc. Symmetric multiprocessor operating system for execution on non-independent lightweight thread contexts
DE102004007835A1 (de) 2004-02-17 2005-09-15 Universität des Saarlandes Vorrichtung zur Darstellung von dynamischen komplexen Szenen
US7580977B1 (en) * 2004-04-29 2009-08-25 Cisco Technology, Inc. System for using text terminal for the deaf (TTD) devices for internet instant messaging and chat sessions
US7616782B2 (en) 2004-05-07 2009-11-10 Intelliview Technologies Inc. Mesh based frame processing and applications
US7348975B2 (en) 2004-12-28 2008-03-25 Intel Corporation Applications of interval arithmetic for reduction of number of computations in ray tracing problems
EP1783604A3 (en) 2005-11-07 2007-10-03 Slawomir Adam Janczewski Object-oriented, parallel language, method of programming and multi-processor computer
US7728841B1 (en) 2005-12-19 2010-06-01 Nvidia Corporation Coherent shader output for multiple targets
JP2009523292A (ja) 2006-01-10 2009-06-18 ブライトスケール インコーポレイテッド 並列処理システムにおけるマルチメディア・データ処理をスケジューリングするための方法及び装置
US8024394B2 (en) * 2006-02-06 2011-09-20 Via Technologies, Inc. Dual mode floating point multiply accumulate unit
US7903125B1 (en) 2006-02-07 2011-03-08 Adobe Systems Incorporated Compact clustered 2-D layout
US7933940B2 (en) 2006-04-20 2011-04-26 International Business Machines Corporation Cyclic segmented prefix circuits for mesh networks
US20070264023A1 (en) 2006-04-26 2007-11-15 Virgin Islands Microsystems, Inc. Free space interchip communications
US8261270B2 (en) 2006-06-20 2012-09-04 Google Inc. Systems and methods for generating reference results using a parallel-processing computer system
CN101657795B (zh) 2007-04-11 2013-10-23 苹果公司 多处理器上的数据并行计算
US8341611B2 (en) 2007-04-11 2012-12-25 Apple Inc. Application interface on multiple processors
US8286196B2 (en) 2007-05-03 2012-10-09 Apple Inc. Parallel runtime execution on multiple processors
US8081181B2 (en) 2007-06-20 2011-12-20 Microsoft Corporation Prefix sum pass to linearize A-buffer storage
US8996846B2 (en) * 2007-09-27 2015-03-31 Nvidia Corporation System, method and computer program product for performing a scan operation
US8072460B2 (en) 2007-10-17 2011-12-06 Nvidia Corporation System, method, and computer program product for generating a ray tracing data structure utilizing a parallel processor architecture
US8284188B1 (en) 2007-10-29 2012-10-09 Nvidia Corporation Ray tracing system, method, and computer program product for simultaneously traversing a hierarchy of rays and a hierarchy of objects
US8264484B1 (en) 2007-10-29 2012-09-11 Nvidia Corporation System, method, and computer program product for organizing a plurality of rays utilizing a bounding volume
US8065288B1 (en) 2007-11-09 2011-11-22 Nvidia Corporation System, method, and computer program product for testing a query against multiple sets of objects utilizing a single instruction multiple data (SIMD) processing architecture
US8243083B1 (en) 2007-12-04 2012-08-14 Nvidia Corporation System, method, and computer program product for converting a scan algorithm to a segmented scan algorithm in an operator-independent manner
US20100076941A1 (en) 2008-09-09 2010-03-25 Microsoft Corporation Matrix-based scans on parallel processors
US8321492B1 (en) 2008-12-11 2012-11-27 Nvidia Corporation System, method, and computer program product for converting a reduction algorithm to a segmented reduction algorithm

Patent Citations (3)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN1202651A (zh) * 1996-12-20 1998-12-23 索尼公司 算术机逻辑单元的运算方法,存贮介质和算术和逻辑单元
US20030081833A1 (en) * 2001-04-23 2003-05-01 National Aeronautics And Space Administration Method for implementation of recursive hierarchical segmentation on parallel computers
CN1431588A (zh) * 2002-01-08 2003-07-23 北京南思达科技发展有限公司 一种逻辑可重组电路

Cited By (25)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US9678751B2 (en) 2011-12-23 2017-06-13 Intel Corporation Systems, apparatuses, and methods for performing a horizontal partial sum in response to a single instruction
CN104081337A (zh) * 2011-12-23 2014-10-01 英特尔公司 用于响应于单个指令来执行横向部分求和的***、装置和方法
CN104081337B (zh) * 2011-12-23 2017-11-07 英特尔公司 用于响应于单个指令来执行横向部分求和的***、装置和方法
US9965282B2 (en) 2011-12-28 2018-05-08 Intel Corporation Systems, apparatuses, and methods for performing delta encoding on packed data elements
CN104040482B (zh) * 2011-12-28 2018-02-16 英特尔公司 用于在打包数据元素上执行增量解码的***、装置和方法
US10671392B2 (en) 2011-12-28 2020-06-02 Intel Corporation Systems, apparatuses, and methods for performing delta decoding on packed data elements
US10037209B2 (en) 2011-12-28 2018-07-31 Intel Corporation Systems, apparatuses, and methods for performing delta decoding on packed data elements
CN104040482A (zh) * 2011-12-28 2014-09-10 英特尔公司 用于在打包数据元素上执行增量解码的***、装置和方法
CN104813279B (zh) * 2012-12-28 2018-12-18 英特尔公司 用于减少具有步幅式访问模式的向量寄存器中的元素的指令
US9921832B2 (en) 2012-12-28 2018-03-20 Intel Corporation Instruction to reduce elements in a vector register with strided access pattern
CN104813279A (zh) * 2012-12-28 2015-07-29 英特尔公司 用于减少具有步幅式访问模式的向量寄存器中的元素的指令
CN105453028A (zh) * 2013-08-14 2016-03-30 高通股份有限公司 向量积累方法及设备
CN105453028B (zh) * 2013-08-14 2019-04-09 高通股份有限公司 向量积累方法及设备
US9740660B2 (en) 2014-04-24 2017-08-22 Mediatek Inc. CPU control method, electronic system control method and electronic system for improved CPU utilization in executing functions
WO2015161826A1 (en) * 2014-04-24 2015-10-29 Mediatek Inc. Cpu control method, electronic system control method and electronic system
CN105940376A (zh) * 2014-04-24 2016-09-14 联发科技股份有限公司 中央处理单元控制方法、电子***控制方法及电子***
CN107003988B (zh) * 2014-12-19 2021-04-02 英特尔公司 用于执行卷积运算的存储设备和方法
CN107003988A (zh) * 2014-12-19 2017-08-01 英特尔公司 用于执行卷积运算的存储设备和方法
CN112256580A (zh) * 2020-10-23 2021-01-22 济南浪潮数据技术有限公司 一种代码扫描方法、装置、设备及存储介质
CN112256580B (zh) * 2020-10-23 2024-02-13 济南浪潮数据技术有限公司 一种代码扫描方法、装置、设备及存储介质
WO2022089421A1 (en) * 2020-10-30 2022-05-05 International Business Machines Corporation ReLU COMPRESSION TO REDUCE GPU MEMORY
GB2615267A (en) * 2020-10-30 2023-08-02 Ibm ReLU compression to reduce GPU memory
CN114510271A (zh) * 2022-02-09 2022-05-17 海飞科(南京)信息技术有限公司 用于在单指令多线程计算***中加载数据的方法和装置
CN114510271B (zh) * 2022-02-09 2023-08-15 海飞科(南京)信息技术有限公司 用于在单指令多线程计算***中加载数据的方法和装置
CN115718724A (zh) * 2023-01-09 2023-02-28 阿里巴巴（中国）有限公司 Gpu、数据选择方法及芯片

Also Published As

Publication number Publication date
JP2009169935A (ja) 2009-07-30
US8661226B2 (en) 2014-02-25
US20090132878A1 (en) 2009-05-21
TWI398779B (zh) 2013-06-11
CN101436121B (zh) 2014-12-17
KR20090050977A (ko) 2009-05-20
TW200928780A (en) 2009-07-01

Similar Documents

Publication Publication Date Title
CN101436121B (zh) 用于使用并行处理器架构执行扫描操作的方法及设备
CN107608715B (zh) 用于执行人工神经网络正向运算的装置及方法
CN101398753A (zh) 用于执行扫描运算的***、方法及计算机程序产品
Rios et al. Exploring parallel multi-GPU local search strategies in a metaheuristic framework
CN110333946A (zh) 一种基于人工智能cpu数据处理***及方法
Watanabe et al. SIMD vectorization for the Lennard-Jones potential with AVX2 and AVX-512 instructions
Cococcioni et al. A novel posit-based fast approximation of elu activation function for deep neural networks
Leitao et al. GPU acceleration of the stochastic grid bundling method for early-exercise options
Ranjan et al. Upper and lower I/O bounds for pebbling r-pyramids
CN103955443A (zh) 一种基于gpu加速的蚁群算法优化方法
Nishimura et al. Accelerating the Smith-waterman algorithm using bitwise parallel bulk computation technique on GPU
Lim et al. R High Performance Programming
Mahony et al. A parallel and pipelined implementation of a pascal-simplex based multi-asset option pricer on FPGA using OpenCL
US20080313612A1 (en) Hysteresis for mixed representation of java bigdecimal objects
US5257394A (en) Logical expression processing pipeline using pushdown stacks for a vector computer
US5261111A (en) Pipelined processor for vector data classification according to vector attributes using coupled comparator chain and logic gate to respective bin chain
JP3259622B2 (ja) ベクトル演算装置を用いたバケットソート処理方式
Fujita et al. An efficient GPU implementation of CKY parsing using the bitwise parallel bulk computation technique
US20230118082A1 (en) Apparatus, method and system for matrix multiplication reusing multiply accumulate operation
US20230004788A1 (en) Hardware architecture for processing tensors with activation sparsity
US20230385025A1 (en) Method and apparatus with repeated multiplication
Panova et al. Black-Scholes Option Pricing on Intel CPUs and GPUs: Implementation on SYCL and Optimization Techniques
Sikorski et al. A Generalized Parallel Prefix Sums Algorithm for Arbitrary Size Arrays
Srivastava et al. A Review on Hardware Implementations of Signal Processing Algorithms
Gustavson et al. High performance Cholesky factorization via blocking and recursion that uses minimal storage

Legal Events

Date Code Title Description
C06 Publication
PB01 Publication
C10 Entry into substantive examination
SE01 Entry into force of request for substantive examination
C14 Grant of patent or utility model
GR01 Patent grant
CF01 Termination of patent right due to non-payment of annual fee

Granted publication date: 20141217

Termination date: 20151111

EXPY Termination of patent right or utility model