2020-10-19

TI OpenCL 用户指南2

12.Execution Model
OpenCL是一个基于主机CPU的库，

1. 它允许您发现系统中可用的计算设备，
2. 定义在这些计算设备上运行的程序，
3. 定义可用于在主机程序和运行在计算设备上的程序之间通信数据的缓冲区，
4. 以及对计算设备排队。
   步骤1记录在设备发现中。步骤3记录在内存使用中。步骤2中定义的程序的编译记录在“编译”中。 其余项目： 步骤2中程序的表达式，以及 步骤4中用于排队工作的机制。

是执行模型的本质，因为它涉及如何在计算设备上完成工作，这是本章的主要主题。
术语
以下术语的定义是复制的，或者是OpenCL1.1规范第2章中定义的修改形式。
计算设备
计算设备是从主机CPU卸载计算的目标.命令队列是在OpenCL应用程序中创建的，并绑定到特定的计算设备上。在内部，计算设备是计算单元的集合。OpenCL计算设备通常对应于GPU、多核CPU或多核DSP.
计算单元
计算设备包含一个或多个计算单元。对于多核设备，计算单元通常对应于其中一个核.工作组在单个计算单元上执行，多个工作组可以在设备内的多个计算单元上并发执行。计算单元将具有只能由计算单元访问的本地内存。
命令队列
由主机应用程序中的OpenCL API创建的对象。在上下文中针对特定设备创建命令队列。命令队列保存将在该特定设备上执行的命令。命令队列的命令按顺序排队，但根据命令队列创建过程中指定的属性，可以按顺序或无序执行。计算设备可以具有与其相关联的许多命令队列，但是命令队列仅与一个计算设备相关联。
内核
内核是在OpenCL C程序中声明并在Compute设备上执行的函数。内核由内核或内核限定符标识。内核通过命令队列排队计算设备。
工作项
当内核排队到命令队列时，enQueue命令指定要完成的工作项数。对于进入队列的NDRangeKernel，工作项的数量由全局大小参数显式指定。对于进入队列任务，工作项的数目隐式地指定为1。
由命令在设备上调用的内核的并行执行的集合之一。由作为在计算单元上执行的工作组的一部分的一个或多个处理元件来执行工作项。工作项与集合中的其他执行区别在于其全局ID和本地ID。
工作组：
在单个计算单元上执行的相关工作项的集合。组中的工作项执行相同的内核并共享本地内存。
全局ID
全局ID用于唯一标识工作项，并从执行内核时指定的全局工作项数派生出来。全局ID是一个N维值，从所有维度中的0开始.
本地ID
本地ID指定正在执行内核的给定工作组中的唯一工作项ID。本地ID是一个N维值，从所有维度中的0开始.
13. Device Discovery
使用OpenCL C++绑定，设备发现是简单的样板代码，它将枚举C++向量中所有请求类型的设备。
Context

context(CL_DEVICE_TYPE_ACCELERATOR);
std::vector<Device> devices = context.getInfo<CL_CONTEXT_DEVICES>();

• 上面的代码将创建一个OpenCL上下文，其中包含CL_DEVICE_TYPE_ACCELERATOR类型系统中的所有设备。然后查询上下文，将这些设备的列表检索到C++向量对象中。TI OpenCL实现中的DSP具有CL_DEVICE_TYPE_ACCELERATOR的设备类型。其他OpenCL设备类型包括：CL_DEVICE_TYPE_DEFAULT
• CL_DEVICE_TYPE_CPU
• CL_DEVICE_TYPE_GPU
• CL_DEVICE_TYPE_ACCELERATOR
• CL_DEVICE_TYPE_ALL
TI OpenCL实现目前只支持DSP设备，因此设备类型CL_Device_type_DEFAULT、CL_Device_TYPE_All都是同义词。

为了最好地为快速执行构建OpenCL代码，必须清楚地理解OpenCL C内核、工作组、工作项、内核中的显式迭代以及这些概念之间的关系。
此外，了解这些概念如何映射到底层硬件体系结构非常重要，这样您的应用程序才能得到适当的结构。要运行OpenCL应用程序，不需要了解这些概念。OpenCL代码是可移植的表达式，这意味着一个符合的OpenCL应用程序应该在任何符合的OpenCL实现上正确运行。然而，为了高效运行OpenCL应用程序，需要了解这些概念，即OpenCL应用程序不一定是性能可移植的，简单地调整显式内核迭代与工作组中工作项的数量之间的平衡可能会对性能产生非常大的影响。
13.1 Enqueueing a Kernel
有两个OpenCLAPI用于在内核中排队。
enQueeNDRangeKernel()和enQueeTask()。
加入队列任务只是一个特殊的情况，在一个维度中，偏移量、全局大小和局部大小分别固定为0、1和1。这个特例在许多情况下非常有用，但是为了解释OpenCL C内核、工作项和工作组之间的关系，本节将重点介绍通用的API enQueeNDRangeKernel。本节还将假设enQueeNDRangeKernel的偏移量为0是所有维度。在相对于全局大小和本地大小的有限用例中，偏移量非常有用，有关使用偏移量的更多信息，请参阅OpenCL1.1规范。

OpenCL C内核的表达式实际上是一个工作项的算法表达式。一个工作项表达式也与内核名称相关联。当进行enQueeNDRangeKernelAPI调用时，API的三个关键参数是：以前与内核名称相关联的内核对象，以及希望执行的工作项(称为全局大小)的数量，以及希望分组到工作组中的工作项的数量(称为本地大小)。

For example, the following C++ code
Q.enqueueNDRangeKernel(K, NullRange, NDRange(1024), NDRange(128));

演示了一个enQueeNDRangeKernel命令，它将队列到名为Q的OpenCL队列，一个名为K的内核对象，第二个参数是偏移量，正如前面所述，我们将假设它在所有维度中为0。NullRange对象将满足该0规范。第三个参数是全局大小，它指定希望执行与内核对象K关联的内核源代码中指定的工作项的1024个实例。第四个参数是本地大小，它指定应该将多少工作项分组到一个工作组中。在这种情况下，指定为每个工作组128个工作项。由于有1024个总工作项和128个工作项/工作组，所以1024/128=8个工作组的简单划分。全局大小(Gsz)是工作项总数(WI)，本地大小(Lsz)是每个工作组(wi/wg)的工作项数，工作组数目是全局大小/本地大小，或gsz/lsz或wg。
还可以在2或3维中指定全局大小和局部大小。例如，2D内核队列可能如下所示：

Q.enqueueNDRangeKernel(K, NullRange, NDRange(640, 480), NDRange(640, 1))

此队列指定：全局大小为640个工作项(维度0中的项)和480个工作项(维度1中的项)，总共640*480=307，200项工作项(WI)。它还规定了0维的640 WI/WG和1维的1 WI/WG的局部大小。这使得在0维的640/640=1工作组和第1维的480/1工作组总共有480个工作组(WG)。
13.2 Mapping the OpenCL C work-item Built-in Functions

OpenCL C包含八个内置函数，可在内核的算法表达式中用于查询全局大小、本地大小等。这八个函数是：

13.3 OpenCL C Kernel Code
OpenCL C内核中的代码表示将应用于单个工作项的算法。工作项的粒度由实施者确定。如果我们采用元素的矢量加法示例，我们将两个一维向量作为输入，将它们相加，并将结果写回到第一个矢量中，我们可以用以下任一方法表达内核来实现此行为：

kernel vectorAdd(global int* A, global const int * B)
{
    int gid = get_global_id(0);
    A[gid] += B[gid];
}
Next example
#define ITER 16

kernel vectorAdd(global int* A, global const int * B)
{
    int gid = get_global_id(0) * ITER;
    int i;
    for (i = 0; i < ITER; ++i)
    {
        A[gid + i] += B[gid+ i];
    }
}

第一个内核将执行每个工作项的输入数组的一个元素的添加，对于长度为1024的数组，enQueeNDRangeKernel调用将需要指定1024作为全局大小。
第二个内核将对每个工作项执行16个元素添加，对于相同的1024个长度的输入数组，enQueeNDRangeKernel调用只需要指定一个64的全局大小，即每个工作项的1024个元素/16个元素。
13.4 NDRangeKernel Execution on DSP Devices
13.4.1 Number of Cores Available for NDRangeKernel Execution
德州仪器公司OpenCL实现中的DSP设备可以是单核DSP，也可以是多核DSP，具有不同数量的核心。无论哪种方式，所有DSP核心的分组将作为一个具有一定数量计算单元的虚拟DSP设备呈现给OpenCL开发人员。计算单元的数量将等于设备中可用的DSP核的数量。
若要查询DSP设备的计算单元(或核心)数量，请使用OpenCL设备查询功能。下面的代码说明了主机OpenCL应用程序如何确定DSP设备中的核数。

 Context context(CL_DEVICE_TYPE_ACCELERATOR);
 std::vector<Device> devices = context.getInfo<CL_CONTEXT_DEVICES>();

 int num;
 devices[0].getInfo(CL_DEVICE_MAX_COMPUTE_UNITS, &num);

如在上一节所述，第1行和第2行将枚举TexasInstruments“OpenCL平台”中的DSP设备的数量。在具有嵌入式ARMDSP设备的EVMS上，这将在设备[0]中返回1个DSP设备。有关PCIe卡上的TMS320C6678设备的“平台”，它将返回4或8个设备。第5行假定设备[0]中只有一个设备，并且它向该设备查询该设备的数量，该设备将在此示例中设置在变量NUM中。
13.4.2 NDRangeKernel Work-Group Execution
虚拟DSP设备中的DSP核心（计算单元）的行为类似于由主机上的入队NDRangeKernel调用创建的工作组的异构线程池。每个DSP核心将从工作组队列中提取一个工作组（如线程池队列）。它将执行工作组以完成，然后从队列中提取另一个工作组。这将继续，直到完成NDRangeKernel提交的所有工作组。
下图说明了在4核心DSP上执行8个工作组的NDRangeKernel的飞行执行情况。绿色方框表示DSP核心。圆圈表示工作组。蓝色工作组正在等待执行，粉红色工作组当前正在执行，黄色工作组已完成。

在本例中，工作组0和3已经完成，工作组1、2和4目前正在执行，工作组6和7正在等待，而Work-group 5刚刚从等待队列中被选中，即将分配给空闲核心1执行。
在所有8个工作组完成后，NDRangeKernel提交将被视为已完成，任何与enQueeNDRangeKernel关联的OpenCL事件都将其状态更新为CL_Complete。一旦事件状态更新为CL_Complete，事件的等待操作将被满足，并且线程将被允许进行。
以上数字可能是由类似于以下代码的代码造成的：

Event ev;
Q.enqueueNDRangeKernel(K, NullRange, NDRange(8), NDRange(1), NULL, &ev);
concurrent();
ev.wait()
post();

在这个示例代码中，已经存在一个名为Q的OpenCL队列和一个名为K的内核。第1行定义了OpenCL事件对象EV。第2行将内核K排队到队列Q，全局大小为8，局部大小为1，总共有8个工作组。第五个参数是在提交内核K之前必须完成的事件向量。在本例中，该依赖项集为空。第六个参数是与此内核提交相关联的事件的地址，该事件在第1行中定义为EV。
当执行第2行的enQueue命令时，它将把这个内核提交放在OpenCL命令-Queue Q中。然后，enQueue命令将返回，并且示例的第3行可以开始执行对并发的调用。
OpenCL运行时异步地监视内核提交的命令队列q，其中满足内核的所有依赖关系。当运行时标识一个时，它将为内核创建适当数量的工作组，并将它们放在与Q关联的设备的工作组队列中。
在本例中，内核K没有依赖项，因此运行时将立即执行此任务，前提是设备可用。此时，主机CPU正在执行函数并发，DSP核同时执行内核K的工作组。让我们假设主机函数并发首先完成，从而允许线程继续到第4行，在与提交内核K相关联的OpenCL事件上执行等待操作。等待操作将阻塞线程，直到事件的状态为CL_Complete为止。
OpenCL运行时将在完成内核K的所有8个工作组之后将EV的状态更新为CL_Complete。当发生这种情况时，第4行上的等待操作将被满足，线程将继续到第5行，其中主机执行函数POST。
13.4.3 NDRangeKernel Work-Item within a Work-Group Execution
在OpenCL应用程序中，内核函数的主体表示要为单个工作项完成的计算。在作为全局大小参数的内核的Enqueuendrangekernel命令中指定要计算的工作项的数量。“本地大小”参数定义了多个工作项在单个工作组中分组的方式。上一节描述了工作组的执行情况。本节将描述工作组内的工作项如何执行。OpenCL实现方式可以在工作小组中如何执行工作项的细节中显著地变化。可变性将基于工作项正在执行的设备的硬件体系结构。对于具有宽SIMD(单指令多数据)架构的GPU，一个工作组内的一些工作项将同时执行，每个SIMD通道的一个工作项将执行。对于本质上是迭代的TexasInstruments “DSP”，工作组内的工作项将顺序执行。为了以高效的方式顺序执行工作项，
“OpenCL实现”中的OpenCL C编译器将在内核函数的主体周围创建循环。例如，在下面的示例代码中，函数向量add是表示单个工作项的计算的内核的示例。它使用函数get_global_id查询全局工作项ID。然后，它使用ID索引到A和B数组中，并将B元素添加到A元素。

kernel vectorAdd(global int* A, global const int * B)
{
    int gid = get_global_id(0);
    A[gid] += B[gid];
}

在TI OpenCL C编译器编译期间，上面的内核表达式被转换成代码来表示整个工作组。转换会使上面的代码看起来像下面的代码。

extern uint32_t _local_size[3];
extern uint32_t _first_gid_in_wg[3];

kernel vectorAdd(global int* A, global const int * B)
{
    int _local_id_0;

    for (_local_id_0 = 0; _local_id_0 < _local_size[0]; _local_id_0++)
    {
        int gid = _local_id_0 + _first_gid_in_wg[0];
        A[gid] += B[gid];
    }
}

外部变量_local_size和_first_gid_in_wg是由三个元素组成的数组，其中一个是NDRange Kernel的每个潜在维度。OpenCL运行时将根据主机应用程序中的enQueue NDRange Kernel命令的参数填充这些变量。 上面的示例内核只表示一个访问维度，因此转换后的内核插入了一个级别的循环。对于具有二维或三维访问的内核，编译器将将内核转换为分别包含两个或三个级别的循环。
13.4.4 NDRangeKernels: Putting it All Together
从前面的部分中，我们已经表明：内核表示为源，表示对一个工作项的计算，在enQueeNDRangeKernel命令中指定的本地大小决定有多少工作项被分组到一个工作组中，在TI DSP设备上，工作组中的工作项通过编译器插入的循环顺序执行，在enQueeNDRangeKernel命令中指定的全局大小决定有多少总工作项，以及在TI DSP设备上，通过将全局大小除以局部大小来确定有多少工作组，工作组以工作池的方式跨DSP核并行执行，直到一个队列命令的所有工作组都完成为止，可以并发执行的工作组的数量由设备中的DSP核数决定(在OpenCL术语中，DSP核心是一个计算单元)。下面的图直观地总结了上面的内容，显示了两个在语义上将执行相同计算的队列NDRangeKernel命令，但是只要简单地更改本地大小参数，内核执行方式的平衡就会发生变化。在这两种情况下，全局大小为1024。在案例1中，本地大小为128，这将导致创建8个工作组的执行分区，每个工作组将遍历128个工作项。在案例2中，本地大小被更改为256，这将导致4个工作组，每个工作组都有256个工作项。

14.Extensions
TI的OpenCL实现已经通过一组超出OpenCL1.1规范的特性进行了扩展。添加这些特性是为了更好地支持在C66xDSP上执行代码，启用现有的DSP库，并更好地映射到TI的设备。
15.Calling Standard C Code From OpenCL C Code
此OpenCL实现支持从OpenCL C代码调用标准C代码的能力。这包括调用现有C66 DSP库中的函数，例如dsplib、mathlib或imglib。有关此功能的示例，请参阅调用您定义的C函数的CCODE示例，或调用库中函数的dsplib_FFT示例。
标准C代码还可能包含控制并行执行的OpenMP实用程序。有关此用例的描述，请参阅从OpenCL C代码调用具有OpenMP的标准C代码。
15.1 Global variables in C code
本节描述从OpenCL C内核调用的C代码中全局变量的行为。全局变量是在文件作用域中声明的变量。必须对全局变量进行注释，否则C代码必须使用–mem_model：data=ar-mem_model：const=data选项编译。有关详细信息，请参阅C 6000编译器用户指南7.5.5.1节。
int far global_var = 0;
在没有任何数据部分转换的情况下，C代码中的全局变量被放置在在OpenCL设备中的计算单元(C66xDSP)之间共享的外部内存(DDR)中。
要为每个计算单元提供自己的全局变量副本，有两个选项：
使用一个数组和使用核心id寄存器DNUM的索引。

#include <c6x.h> // for DNUM
int far global_var[MAX_COMPUTE_UNITS]; // 2 for AM572, 8 for K2H
...

void foo()
{
   ... = global_var[DNUM];
}

使用数据部分杂注将全局变量放在本地内存中，用于计算单元。但是，这些变量的任何显式初始化都会被忽略，并且不执行C标准所要求的隐式零初始化。

#pragma DATA_SECTION(global_var, ".mem_l2")
int global_var;   // Initialization has no effect

16. Calling Standard C code with OpenMP from OpenCL C code
    从OpenCL C代码调用的标准C代码可以包含OpenMP实用程序。在使用此特性时，包含对启用OpenMP的C函数的调用的OpenCL C内核必须作为任务(而不是NDRangeKernel)提交，并且必须将其提交到内部OpenCL命令队列(即，未使用CL队列无序exec模式启用标志定义)。在这个场景中，OpenCL将把内核分派给DSP加速器的一个计算单元，OpenMP运行时将管理跨计算单元的任务分发。请参阅vecadopenmp示例或vecadopenmp t示例。
    17.OpenMP dispatch from OpenCL
    OpenCL程序由在主机上执行的主机程序和在OpenCL设备上执行的内核组成。TI扩展允许OpenCL内核作为调用包含OpenMP区域的C函数的封装。这三个组件-主机程序、OpenCL内核和OpenMP区域-形成能够调度OpenMP区域的OpenCL程序的主要部分。下图提供了OpenMP区域如何使用OpenCL API发送的概述。主机程序在ARMCortex-A15上运行的SMP Linux上执行，而OpenCL内核和OpenMP区域在由DSP组成的OpenCL设备上执行。OpenCL内核调用DSP内核0上的OpenMP主线程。当在执行主线程时遇到OpenMP区域时，DSP内核0和其他DSP一起工作以执行OpenMP区域。
    

17.1 Structure of an OpenCL + OpenMP Application
本节使用OpenCL包中的vecadd_openmp示例来描述OpenCL + OpenMP应用程序的结构。
17.1.1 Host Code
主机代码负责在OpenCL设备上设置和管理内核的执行。
典型OpenCL+OpenMP应用程序的主机代码从创建OpenCL上下文开始。上下文有助于创建程序和内核，创建命令队列，管理主机和DSP之间的内存，以及提交内核以便在设备上执行。
下面示例vadd_openmp应用程序的主机代码说明了主机代码的责任。

...
float srcA  [NumElements];
float srcB  [NumElements];
float dst   [NumElements];
float Golden[NumElements];
...

// Create a context with the DSP device
Context context(CL_DEVICE_TYPE_ACCELERATOR);

// Get information about the device associated with the context
std::vector<Device> devices = context.getInfo<CL_CONTEXT_DEVICES>();
std::string str;
devices[d].getInfo(CL_DEVICE_NAME, &str);
cout << "DEVICE: " << str << endl << endl;

// Create input and output buffers
Buffer bufA(context, CL_MEM_READ_ONLY,  bufsize);
Buffer bufB(context, CL_MEM_READ_ONLY,  bufsize);
Buffer bufDst(context, CL_MEM_WRITE_ONLY, bufsize);

ifstream t("vadd_wrapper.cl");
std::string kSrc((istreambuf_iterator<char>(t)), istreambuf_iterator<char>());

// Create a program from the kernel source code
Program::Sources source(1, make_pair(kSrc.c_str(), kSrc.length()));
Program program = Program(context, source);

// Compile the kernel source and link it with the specified object file
program.build(devices, "vadd_openmp.obj");

// Specify which kernel from the program to execute
Kernel kernel(program, "vadd_wrapper");

// Set the argument list for the kernel command
kernel.setArg(0, bufA);
kernel.setArg(1, bufB);
kernel.setArg(2, bufDst);
kernel.setArg(3, NumElements);

Event ev1,ev2,ev3,ev4,ev5,ev6,ev7,ev8;

// Create command queue using the context and device
CommandQueue InO_Q(context, devices[d], CL_QUEUE_PROFILING_ENABLE);

// Enqueue commands to copy data into the input buffers
InO_Q.enqueueWriteBuffer(bufA, CL_FALSE, 0, bufsize, srcA, NULL, &ev1);
InO_Q.enqueueWriteBuffer(bufB, CL_FALSE, 0, bufsize, srcB, NULL, &ev2);

std::vector<Event> vec_ev5(1);

// Enqueue the kernel for execution as an OpenCL task
InO_Q.enqueueTask(kernel, NULL, &vec_ev5[0]);

// Enqueue command to copy results from out of the output buffer
InO_Q.enqueueReadBuffer(bufDst, CL_TRUE, 0, bufsize, dst, &vec_ev5, &ev6);

调用OpenMP应用程序的内核必须排队到有序命令队列中。顺序执行序列化命令队列中命令的执行顺序，从而确保设备一次只执行一个“OpenMP”内核。
注意到主机排队命令在内核排队之前写入Bufa和bufB(从驻留在主机上的srcA和srcB)。这确保了在执行OpenCL内核之前完成数据传输。当内核完成执行时，从bufDst读取到DST(驻留在主机上)的命令将运行。
对于OpenCL+ OpenMP应用程序，内核被分派为OpenCL任务。只有一个内核实例由其中一个DSP执行。
17.1.2 Kernel
OpenCL+ OpenMP应用程序中的内核本质上是调用包含OpenMP区域的函数的包装器。在vadd_openmp示例中，vadd_Wrapper内核在初始化c[]后调用vadd_openmp(其中包含一个OpenMP区域)。注意，内核最初由一个DSP执行，直到遇到OpenMP区域。

__kernel void vadd_wrapper(__global const float *a,
                           __global const float *b,
                           __global       float *c,
                                          int    size)
{
  vadd_openmp(a, b, c, size);
}

具有OpenMP区域的C函数
OpenMP指令用于表达和指导并行度。在vadd_OpenMP函数中，for循环之前的指令表示循环的chunks可以由多个核分布和并发执行。

void vadd_openmp(float *a, float *b, float *c, int size)
{
  int i;
  #pragma omp parallel for
  for (i = 0; i < size; i++)
    c[i] = a[i] + b[i];
}

17.1.3 Makefile
Makefile定义了生成和清理可执行文件的规则。它至少指定主机和设备编译器、编译器标志和链接器选项。在vecadd_openmp示例中，主机代码和OpenMP DSP代码分别由GCC和C 6000编译器分别编译。指定了–OMPc 6000编译器标志，以便为DSP代码启用OpenMP编译。

```bash
# Host compiler and compile flags/options. Used to compile host code
CXX   = g++ $(HOST_INCLUDE)
CXX_FLAGS  = -O3
HOST_INCLUDE = -I$(TI_OCL_INSTALL)/include

# Device OpenMP compiler and compile flags/options. Used to compile OpenMP code for device
CL6X  = cl6x -mv6600 --abi=eabi $(DSP_INCLUDE)
CL6X_FLAGS = -O1 --omp
DSP_INCLUDE  = -I$(TI_OCL_CGT_INSTALL)/include

# linker options and libraries for final executable
LIBS  = -L$(TI_OCL_INSTALL)/lib -lOpenCL -locl_util

# rule for building final executable
EXE        = vecadd_openmp
$(EXE): main.o vadd_openmp.obj
        @$(CXX) $(CXX_FLAGS) main.o $(LIBS) -o $@

# rules for compiling host C/C++ code
%.o: %.cpp
        @$(CXX) -c $(CXX_FLAGS) $<
        @echo Compiling $<
%.o: %.c
        @$(CXX) -c $(CXX_FLAGS) $<
        @echo Compiling $<

# rule for compiling OpenMP C code for device
%.obj: %.c
        @$(CL6X) -c $(CL6X_FLAGS) $<
        @echo Compiling $<

17.2 Guidelines for writing OpenCL + OpenMP applications
使用OpenCL内核来调度OpenMP区域是OpenCL的一个特定于TI的扩展。下面是在这种模式下编写应用程序的指南。
C 6000编译器目前支持C的OpenMP3.0规范，不支持C中的OpenMP指令。有关OpenMP应用程序接口版本3.0的完整C语言规范，请参阅OpenMP网站。
不支持OpenMP线程私有变量。OpenMP DSP运行时中只有一个级别的并行性可用。因此，嵌套的并行区域将由只包含一个线程的团队执行。OpenMP环境变量不受支持，并且对OpenMP运行时行为没有影响。OpenMP定时例程(OMP获得时间，OMP获得时间)假设DSP核心运行在1 GHz。OpenCL中OpenMP运行时的配置为每个核心的堆栈保留了10 KB的内存。由于堆栈大小很小，请避免在堆栈上放置大型数组/结构。此外，保持调用堆栈简短，以避免堆栈溢出。OpenMP运行时的当前配置为共享堆预留了8MB内存。由于这个堆的大小很小，所以避免动态内存分配(使用mallocs)。或者，在主机上创建OpenCL缓冲区，并将它们作为参数传递给内核。默认情况下，全局变量放在DDR中。
18. C66x standard C compiler intrinsic functions
用于C66x DSP的OpenCL C编译器支持C66x标准C编译器集的内部函数，但接受或导致40位值的内部函数除外。请参阅第7.5.5分节。使用Intrinsics访问C 6000编译器用户指南中的汇编语言语句，以了解这些内部函数。
19. OpenCL C code using printf
在DSP上运行的内核增加了printf功能。OpenCLC内核或从OpenCLC内核调用的标准C代码可以调用printf。由printf产生的字符串将被传送到主机ARM，并使用ARM侧的printf显示。此特性可用于帮助调试OpenCL内核。请注意，使用来自DSP的printf是有性能损失的，因此在评估DSP性能时不应该使用它。此特性不是OpenCL1.2 printf，它包含用于打印向量类型的附加格式代码。
20. DMA Control Using EdmaMgr Functions
该OpenCL实现还支持从DSP和OpenCLC内核直接访问EDMA系统。支持多种EDMA构建体。这些包括1D到1D、1D到2D、2D到1D以及链式传输。这些附加EDMAOpenCLC内置模块的原型见下文。也请参见EdmMgr示例、SGEMM示例和Dgem示例，例如用法。如果使用，请将#Include添加到您的源。标准OpenCL C内置函数ASYNC_WORKS_GROUP_COPY和ASYNC_WORKS_GROUP_STRDED_COPY还将使用片内DMA，并且可以与EDMAMGR功能组合使用。EDMAMGR功能在标准OpenCLC内置功能上提供了额外的功能。
EdmaMgr
 

```c
typedef void *EdmaMgr_Handle;
An opaque type used to hold EdmaMgr Handle data structures.

用于保存EdmaMgr句柄数据结构的不透明类型

EdmaMgr_Handle EdmaMgr_alloc(int32_t max_linked_transfers)

分配与所有其他EDMAMGRAPI一起使用的EDMAMGR句柄。对于与单个传输EdmMgrAPI一起使用的句柄，max_linked_transfer参数应为1。它可以大于用于多个传输API的句柄之一。

int32_t EdmaMgr_free(EdmaMgr_Handle h)
Free the resources associated with the EdmaMgr Handle.
void EdmaMgr_wait(EdmaMgr_Handle h)

等待所有与句柄H相关联的EDMA传送在继续执行该线程之前完成。
Single Transfer EdmaMgr APIs
这组EdmaMgr函数将在一个源地址和目标地址上启动EDMA操作。
Multiple Transfer EdmaMgr APIs
Multiple Transfer EdmaMgr APIs
Fast Global buffers in on-chip MSMC memory
已添加用于从MSMC内存中分配缓冲区的TI扩展。除了在缓冲区驻留的DSP上的位置之外，此MSMC定义的缓冲区将以所有其他方式充当标准全局缓冲区。示例：

Buffer bufMsmc(context, CL_MEM_READ_ONLY|CL_MEM_USE_MSMC_TI, size);
Buffer bufDdr (context, CL_MEM_READ_ONLY, size); }

Matmpy示例说明了MSMC缓冲区的使用。平台示例将查询DSP设备，并报告可供OpenCL使用的MSMC内存量。
OpenCL C Builtin Function Extensions
以下内置功能可从DSP设备上的OpenCL C代码中获得。它们也可从OpenCL C代码调用的标准C代码中获得。如果在标准C代码中使用这些功能，请将#Include <dsp_c.h>添加到您的源。此标题定义了函数原型。

uint32_t __core_num(void)

返回发出分配的DSP的核心ID。该值将在范围[0。n-1]其中n是设备中C66DSP内核的数量。

uint64_t __clock64(void)

返回一个64位无符号时间戳，表示DSP中的周期计数计数器。调用此函数两次并减去结果是确定DSPOpenCLC或标准C代码中两点之间经过的循环数的有效机制。此功能相当于C66DSP上的TSCH：TSCL寄存器对。

uint32_t __clock(void)

返回32位无符号时间戳，表示DSP中的循环计数计数器.两次调用此函数并减去结果是一种有效的机制，用于确定DSP、OpenCL C或标准C代码中两个点之间的运行周期数。这个函数相当于C66 DSP上的TSCL寄存器对。

void __cycle_delay(uint64_t cyclesToDelay)

给定要延迟的周期数，此函数在返回之前将繁忙循环大约有那么多个周期。

void __mfence(void)

为C66x DSP创建内存围栏。这个函数相当于OpenCLC内置函数mem_fence()，但是这个版本也可以从从OpenCLC代码调用的标准C代码中调用。

uint dot(uchar4 a, uchar4 b)

计算两个uchar4向量的点积，并将结果作为uint返回。C66x DSP可以作为一条指令支持这一点。

int dot(short2 a, short2 b)
Compute the dot product of two short2 vectors and return the result as an int. The C66x DSP can support this as a single instruction.
uint32_t __dsp_frequency(void)
Returns the clock frequency (Mhz) the DSP cores are running at.

21. Calling TI BIOS APIs from OpenCL C kernels
    将TI扩展到OpenCL，允许从OpenCL C内核调用通用的标准C函数，可以使用它在DSP上分配代码，这些代码使用TI的BIOS API。
    Use cases
    能够为从OpenCL发送的计算设置一个超时，注册一个任务或时钟函数，使其在正常OpenCL内核调度之外的DSP上持续存在，
    能够启动和停止具有独立OpenCL内核的DSP计算
    能够在DSP上创建额外的BIOS任务(线程)并与信号量进行协调
    能够通过OpenCL发送in-flight DSP代码与主机线程通信，而不需要一个新的OpenCL C内核来启动或完成。

允许BIOSAPI调用的扩展目前只支持TI的AM57x系列设备。来自DSP代码的Printf功能只支持OpenCL C内核或内核直接调用树中的代码。相反，在任何用户创建的BIOS任务或时钟处理程序中，DSP代码都不支持printf。OpenCL运行时为OpenCL运行时和任何OpenCL C内核代码的内部保留BIOS任务优先级级别6-10。任务优先级级别0-5可以用于比OpenCL优先级低的用户创建的任务，任务优先级级别11-15可以用于用户创建的优先级高于OpenCL的任务。所有分配给DSP的OpenCL C内核都作为BIOS任务的一部分运行。BIOS不是一个分时的实时操作系统。如果一个更高优先级的任务已经准备好运行，那么它将一直运行，直到它被阻塞，此时，下一个准备运行的最高级别任务将被调度。如果用户创建的任务被赋予比OpenCL运行时使用的调度代码更高的优先级，那么它必须完成或定期阻塞。如果没有，那么OpenCL运行时任务将不再被调度，系统可能会死锁。所有以时间单元作为参数或返回时间单元作为值的BIOS API都是BIOS时钟滴答，这些时钟被配置为每1毫秒发生一次，因此对Task_sleep (1000)的调用实际上是1秒的睡眠。计算DSP时钟在Task_sleep ()调用周围的滴答数可能不会给出预期值。如果DSP没有要执行的其他任务，则会发生这种情况，在这种情况下，它将处于空闲状态，CPU时钟刻度寄存器也不会提前。例如，在DSP运行频率为750 MHz和以下代码的情况下，可以预期经过的变量包含大约750，000，000。然而，在现实中，经过的价值可能要小得多。

unsigned t0      = __clock();  // value of DSP's TSCL register
Task_sleep(1000);              // equivalent of 1 second
unsigned t1      = __clock();  // value of DSP's TSCL register
unsigned elapsed = t1-t0;      // cycles elapsed between __clock() calls

建议您不要直接从OpenCLC代码调用BIOSAPI，而是从OpenCLC代码调用的标准C代码调用BIOSAPI。由于long和ulong数据类型的不同解释，因此建议使用此选项。OpenCLC将这些类型定义为64位，但DSP的标准C编译器将这些类型解释为32位。BIOSAPI使用由BIOS定义的数据类型，这些类型对应于标准C编译器解释，而不是OpenCLC解释。支持来自以下BIOS和IPC包的API：
o ti.sdo.ipc.MessageQ
o ti.sysbios.family.c66.Cache
o ti.sysbios.heaps.HeapBuf
o ti.sysbios.heaps.HeapMem
o ti.sysbios.knl.Clock
o ti.sysbios.knl.Idle
o ti.sysbios.knl.Intrinsics
o ti.sysbios.knl.Queue
o ti.sysbios.knl.Semaphore
o ti.sysbios.knl.Swi
o ti.sysbios.knl.Task
o
OpenCL软件包中包含的以下示例使用来自OpenCL C内核的C代码中的BIOS API来说明：
• persistent_clock_concurrent
• persistent_clock_spanning
• persistent_common
• persistent_kernel_timeout
• persistent_messageq_concurrent
• persistent_task_concurrent
• persistent_task_spanning
示例位于/usr/share/ti/examples/OpenCL中。建立这些示例需要从相应的处理器SDKRTOS包中安装XDC/BIOS/IPC软件包。以下是makefile正在寻找的、覆盖它们以设置每个软件包或整个处理器SDKRTOS安装的选项。

PRSDK_INSTALL_PATH     ?= /cgnas/ti-processor-sdk-rtos-am57xx-evm-04.00.00.04
XDC_DIR                ?= $(wildcard $(PRSDK_INSTALL_PATH)/xdc*)/packages
BIOS_DIR               ?= $(wildcard $(PRSDK_INSTALL_PATH)/bios*)/packages
IPC_DIR                ?= $(wildcard $(PRSDK_INSTALL_PATH)/ipc*)/packages

22. Setting Timeout Limit on OpenCL Kernels

此TI-ExtendedOpenCL实现支持在OpenCL内核上设置超时限制的能力。当在设备上的内核执行完成之前已经达到超时限制时，内核执行将在设备上终止，并且将返回到内核事件的负状态。之后，设备将准备好运行下一个内核。

支持超时的语义
我们目前仅在OpenCL计算单元上本地支持超时。也就是说，每个计算单元具有其自己的超时时钟，并且独立地对所指定的超时限制进行比较。关于多个工作组和多个计算单元，请执行以下规则：
如果内核有多个发送到同一计算单元的工作组，则当第一个此类工作组开始执行时，时钟将开始，并在所有工作组中保持运行。如果在执行任何工作组期间达到超时限制，则将终止内核。如果在任何计算单元上发生超时，则内核事件状态将设置为负值，CL_ERROR_Kernel_TIMEOUT_TI。
如果在内核上设置了超时限制，并且在执行期间发生超时，则用户可以查询相应的内核事件以查询超时错误状态。在其等待列表中具有超时内核事件的所有后续内核事件将根据OpenCL规范将其执行状态设置为CL_EXEC_status_error_for_events_in_wait_list以处理内核错误。
22.1 OpenCL extensions and APIs

OpenCL platform extension: cl_ti_kernel_timeout
OpenCL DSP device extension: cl_ti_kernel_timeout_compute_unit
OpenCL DSP device queue property: CL_QUEUE_KERNEL_TIMEOUT_COMPUTE_UNIT_TI
OpenCL command queue property: CL_QUEUE_KERNEL_TIMEOUT_COMPUTE_UNIT_TI
OpenCL kernel event status: CL_ERROR_KERNEL_TIMEOUT_TI
OpenCL host API: cl_int __ti_set_kernel_timeout_ms(cl_kernel d_kernel, cl_uint timeout_in_ms)

22.2 Example of querying, setting and checking timeout
timeout example说明了超时扩展如何工作。所涉及的步骤是：
1.检查设备队列属性，以确定是否支持超时扩展。

devices[0].getInfo(CL_DEVICE_QUEUE_PROPERTIES, &devq_prop);
if ((devq_prop & CL_QUEUE_KERNEL_TIMEOUT_COMPUTE_UNIT_TI) != 0)

2. 创建具有超时属性的CommandQueue

new CommandQueue(context, devices[0],CL_QUEUE_KERNEL_TIMEOUT_COMPUTE_UNIT_TI);

3. 在内核上设置以毫秒为单位的超时限制
   __ti_set_kernel_timeout_ms(K(), 100);
4. 在对内核进行排队时创建一个内核事件

ev = kernel_functor(...);
//or: enqueueNDRangeKernel(kernel, Range, Range, Range, wait_evs, &ev);
//or: enqueueTask(kernel, wait_evs, &ev);

5. 检查内核事件状态，以确定是否发生超时。

ev.getInfo(CL_EVENT_COMMAND_EXECUTION_STATUS, &status);
if (status == CL_ERROR_KERNEL_TIMEOUT_TI)

Environment Variables
这些环境变量可用于控制OpenCL行为并提供调试的可见性。
TI_OCL_KEEP_FILES
当为DSP编译OpenCL C内核时，结果是/tmp子目录中的二进制.out文件。它们随后可下载到DSP以供运行。编译过程为每个源文件生成多个中间文件。OpenCL通常会删除这些临时文件。但是，有时检查这些文件是有用的。可以将此环境变量设置为指示运行时将临时文件留在/tmp中。检查与OUT文件关联的程序集文件，可以帮助您查看代码优化得有多好。
TI_OCL_DEBUG
设置此环境变量将修改OpenCL应用程序的执行，以启用OpenCLC内核的调试。如果应用程序使用在线编译(即从字符串而不是二进制编译)，那么在线编译将调试标志断言给编译器。(如果应用程序使用脱机编译，即从二进制文件创建程序，则需要将-g作为选项传递给脱机编译器clocl。OpenCL运行时在分派所有内核之前暂停应用程序。当暂停运行时，运行时会向用户指示内核调度正在挂起。
目前，通过将TI_OCL_DEBUG设置为“gdb”或“ccs”，支持两种调试方法。如果设置为“gdb”，运行时将提供gdbc6x命令来连接到DSP并设置适当的断点。如果设置为“CCS”，运行时将提供CodeComposerStudio(CCS)指令，以连接到DSP并设置适当的断点。在这两种方法中，运行时强制内核中的所有内核和工作组只在dsp核心0上执行。详细信息可以在调试中找到。
TI_OCL_CACHE_KERNELS
内核的在线编译是便携式OpenCL程序的一个有用特性。为设备编译内核所需的所有细节都封装在OpenCL API调用中。然而，DSP的在线编译可能很费时.设置此环境变量会导致OpenCL运行时对内核执行一次在线编译，并将结果缓存在/tmp中的数据库中。在不修改内核源代码或编译它的选项的情况下再次运行应用程序，就会绕过编译步骤，并使用缓存的内核二进制文件。
TI_OCL_COMPUTE_UNIT_LIST
将OpenCL运行时可用的计算单元指定为一个逗号分隔的计算单元索引列表，从0开始。指定的计算单元列表必须是连续的，例如K2H上的“1，2，3，4”。如果未指定环境变量，则运行时默认使用设备上可用的所有计算单元(或从OpenCL产品版本1.1.13开始，所有可用的计算单元在/etc/ti-MCTD/ti_MCTD_config.json中指定)。
Example usage on AM572x:

Listing 9 runs the vecadd kernel only on DSP1¶
-> TI_OCL_COMPUTE_UNIT_LIST="0" ./vecadd
Listing 10 runs the vecadd kernel only on DSP2¶
-> TI_OCL_COMPUTE_UNIT_LIST="1" ./vecadd
Listing 11 runs the vecadd kernel on both DSP1 and DSP2 (default behavior)¶
-> TI_OCL_COMPUTE_UNIT_LIST="0, 1" ./vecadd

TI_OCL_LOAD_KERNELS_ONCHIP
默认情况下，OpenCL内核相关代码和全局数据将从DDR内存中分配。如果设置了此环境变量，则从MSMC内存中分配与内核相关的代码和全局数据。
TI_OCL_CPU_DEVICE_ENABLE
目前，OpenCL ARM CPU设备只支持本机内核(有关本机内核的描述，请参阅OpenCL1.1规范)。因此，默认情况下，在执行OpenCL平台查询时，ARM CPU不被视为计算设备。如果应用程序只对本机内核使用ARM CPU，那么可以使用这个环境变量作为OpenCL的计算设备。即使设置了这个环境变量，也不支持将NDRangeKernels或任务排队到CPU。
TI_OCL_WORKER_SLEEP
OpenCL运行时启动应用程序中定义的每个OpenCL命令队列的新CPU线程。这些线程管理OpenCL命令队列以及CPU与命令队列关联的设备之间的通信。如果在设备上有主动运行的OpenCL内核，则为监视与设备代表这些内核的通信而分配的线程消耗CPU资源，检查这些内核的状态。该环境变量可用于提供对消耗多少CPU资源的控制级别。当不设置 TI_OCL_WORKER_SLEEP时，OpenCL运行时使用更多的CPU容量，以确保内核执行上最快的周转延迟。当将ti_OCL_worker_sleep环境变量设置为几微秒时，它会降低内核执行的周转延迟，以降低监视内核所需的CPU容量。如果应用程序不是由CPU周期限制的性能，或者应用程序对许多细粒度内核进行排队，那么具有 TI_OCL_WORKER_SLEEP环境变量UNSET是合适的。在相反的情况下，当CPU周期限制了应用程序的性能或如果更少，但是更长的运行内核会排队，则将 TI_OCL_WORKER_SLEEP置为某个数量的微秒是合适的。要使用的微秒的正确数量取决于执行平台和特定应用。然而，使用从80到150的范围内的微秒值是合理的起点。
TI_OCL_ENABLE_FP64
C66xDSP是双精度浮点，在OpenCL实现中支持双精度浮点OpenCL规范中的所有可选功能，但双FP支持包括子正常行为或优雅下溢的要求除外。C66xDSP上的64位浮点硬件不支持子正常行为。它支持与零行为的刷新。为了支持双打的子正常行为，需要软件仿真，这将会对C66xDSP的硬件性能造成重大的性能损失。因此，默认情况下，在TIOpenCL实现中支持的平台和设备不支持双浮点的支持。也就是说，如果查询了平台或设备进行扩展，则默认情况下不列出CL_KHR_FP64。此外，OpenCLC预定宏CL_KHR_FP64不会默认定义。当设置TI_OCL_ENABLE_FP64环境变量时，TIOpenCL实现报告支持双浮点，即CL_KHR_FP64被列为平台的扩展，并在编译OpenCL内核时定义DSP设备和CL_KHR_FP64。此环境变量控制OpenCLImplementation是否支持Double。然而，双、所有双矢量类型和使用双打的所有内置函数都被支持和可用，而不考虑该环境变量的设置。
TI_OCL_VERBOSE_ERROR
OpenCL规范提供了一个明确定义的机制，用于从API函数返回错误代码。然而，通常情况是出于不同原因返回通用错误代码。当设置此环境变量时，除了定义的返回代码错误机制之外，OpenCL运行时可能会打印更多的描述错误消息。
TI_OCL_WG_SIZE_LIMIT
OpenCL将查询提供给一个设备，以查找工作组中允许的最大工作项数量。TI的实施中的DSP设备允许每个工作组大量的工作项。其他OpenCL实现具有小得多的最大工作组大小限制。当运行代码为其他OpenCL实现而设计和优化时，此环境变量可用于限制报告的最大工作组大小。
TI_OCL_PROFILING_EVENT_TYPE
指定要配置的硬件事件类型。两个基本的划分，失速周期事件和内存事件，描述在剖析。如果指定了1，OpenCL运行时将分析一个失速循环事件。如果指定了2，OpenCL运行时将分析一个或两个内存事件。否则，将禁用分析。
TI_OCL_PROFILING_EVENT_NUMBER1
指定要配置文件的事件号。此变量的确切值表示来自AET_GEM_START_EVT_START或AET_GEM_MEM_EVT_START的偏移量，这取决于事件类型。有关完整的事件列表，请参见分析。
TI_OCL_PROFILING_EVENT_NUMBER 2
指定要配置文件的第二个内存事件号。可以跳过。
TI_OCL_PROFILE_STACK_BLOCK
阈值指定了要计数的失速循环的阈值。计数器中只捕获失速周期高于此阈值的失速事件。OpenCL运行时中的默认值为1，即捕获所有暂停事件。
23. Dispatch from multiple Linux processes
在OpenCL版本1.1.11及更高版本上，与OpenCL应用程序对应的Linux进程在遇到第一个OpenCL API调用时获得OpenCL DSP设备的所有权，并在其结束时放弃所有权。互斥访问是使用文件锁/var/lock/OpenCL实现的。如果进程A和B包含OpenCL API调用并且A开始运行，A将在执行第一个OpenCL API调用时获得文件锁。进程A将保持锁，直到它完成执行，此时进程B可以获得锁。另外，如果一个进程分叉了一个子进程，并且两个进程都调用了OpenCL API，那么这些进程就会死锁。从1.1.12版本开始，取消了进程级锁/var/lock/OpenCL。多个Linux进程可以将OpenCL内核分派给DSP。每个内核运行到完成。由多个进程提交的内核的执行是交错的。
ti-mctd
为了启用并发进程来调度OpenCL API，添加了一个守护进程（ti-mctd）。守护进程提供以下服务：
管理OpenCL连续内存（CMEM）堆。由于OpenCL API的多个进程共享堆，因此守护进程提供了一个集中的位置来管理堆。 开放式DSP监控寿命管理（仅在K2x SoCs上）。在K2x设备上，ti-mctd守护进程使用MPM在启动期间在DSP上重置、加载和运行OpenCL监视器。 作为systemd服务，在引导期间启动ti-mctd守护进程。
关联的systemd单元文件位于/lib/systemd/system/ti-mct-daemon.service。
使用场景2，可以使用守护进程配置文件/etc/ti-mctd/ti_mcd_config.json中指定的更大的Linux共享内存区域来停止和重新启动ti-MCTD守护进程。例如：

# pkill ti-mctd
# edit /etc/ti-mctd/ti_mctd_config.json, increase linux-shmem-size-KB from 128 to 256
# rm /dev/shm/HeapManager (if still exists)
# ti-mctd
# ls -lh /dev/shm/HeapManager
-rw-r--r--    1 root     root      256.0K May  2 

15:44 /dev/shm/HeapManager
ti-mct-heapcheck
ti-mct-heap-check打印OpenCL CMEM堆的当前状态。例如，下面的输出是在执行OpenCL程序期间运行ti-mct-heap-check时生成的。
可以使用-c选项释放任何分配的块，而分配这些块的进程没有释放这些块。这种情况通常发生在进程异常终止时。
ti-mctd config file
从OpenCL产品版本1.1.13开始，ti-mctd守护进程在启动时读取配置文件etc/ti-mctd/ti_mctd_config.json。以下是来自K2H EVM的示例配置：

root@k2hk-evm:~# cat /etc/ti-mctd/ti_mctd_config.json
{
        "cmem-block-offchip" : "0",
        "cmem-block-onchip" : "1",
        "compute-unit-list" : "0,1,2,3,4,5,6,7",
        "linux-shmem-size-KB" : "128",
        "eve-devices-disable" : "0",
}

cmem-block-offchip和cmem-block-onchip指定用于OpenCL的CMEM块ID。OpenCL需要片外CMEM块，而片上CMEM块是可选的。提供了这两个配置项，以便OpenCL CMEM块可以与专用于其他用途的CMEM块共存。
compute-unit-list指定OpenCL应用程序可以使用的DSP。这是一个系统范围的配置。它可以被环境变量TI_OCL_COMPUTE_UNT_LIST在每个shell或每个应用程序的基础上重写。计算单元列表必须是连续的。但是它不需要从第一个可用开始，以系统中最后一个可用的DSP结束。例如，1，2，3，4，5，6是K2H上的有效计算单元列表。OpenCL内核将只被分派到此列表中指定的DSP。提供此配置项是为了使OpenCL DSP可以与专用于其他用途的DSP共存。
linux-shmem-size-KB指定守护进程和OpenCL应用程序用于管理共享数据结构的Linux共享内存的大小。正如前面一节所讨论的，如果您看到bad_alloc异常，请增加它。
eve-devices-disable指定是否禁用EVE设备，如果在SoC上可用(例如AM57x9s)。默认情况下，值为0，这意味着在OpenCL运行时启用了EVE设备。如果用户选择不使用OpenCL运行时中可用的EVE设备，请将值更改为1。
Restrictions on multiple OpenCL processes
从OpenCL产品版本1.1.13开始，当不使用OpenMP扩展时，OpenCL应用程序可以与其他OpenCL应用程序并行运行，只要对于任何两个并发应用程序，它们都使用相同的计算单元列表或两个不重叠的计算单元列表。例如，sgemm示例不使用OpenMP扩展，可以在K2H上并行运行以下sgemm实例：

root@k2hk-evm:~# ./sgemm -M 1024 -K 2000 -N 1000 -r & \
                 ./sgemm -M 1000 -K 2000 -N 1024 -r
## or launch them simultaneously in two separate windows/shells
root@k2hk-evm:~/examples/sgemm# \
TI_OCL_COMPUTE_UNIT_LIST="0,1,2,3" ./sgemm -M 1024 -K 2000 -N 1000 -r & \
TI_OCL_COMPUTE_UNIT_LIST="0,1,2,3" ./sgemm -M 1000 -K 2000 -N 1024 -r & \
TI_OCL_COMPUTE_UNIT_LIST="4,5" ./sgemm -M 1000 -K 2048 -N 1024 -r & \
TI_OCL_COMPUTE_UNIT_LIST="6,7" ./sgemm -M 1024 -K 2048 -N 1000 -r
## or launch them simultaneously in four separate windows/shells