2020-10-19

TI OpenCL 用户指南3

24. Optimization Tips
    OpenCL应用程序由主机应用程序和一组设备内核组成。主机代码和设备代码都有优化技术。存在跨越主机和设备之间的边界的一些技术。本节提供了编写OpenCL应用程序的技巧，该应用程序执行得很好。它以DSP为加速器设备，以TI SoCs为目标。这些提示被组织成基于尖端适用的部分，即主机或设备。
25. Optimization Techniques for Host Code
    使用离线的嵌入式编译模型
    OpenCL允许在主机代码运行时动态编译设备代码。这允许应用程序的可移植性，但显然，当编译发生时，它会减慢宿主应用程序的运行速度。为了加快主机应用程序的速度，应该离线编译设备代码，即在主机应用程序运行之前。编译部分中有两个使用离线编译的编译模型。对于最快的操作，带嵌入式对象模型的离线编译将是最快的.有关为该模型构造代码的详细信息，请参阅从二进制创建带有嵌入式二进制的OpenCL程序。

避免共享内存SoC平台上的读/写缓冲模型
在共享存储器SoC平台上，主机和设备具有读取和写入相同存储器区域的能力。但是，Linux系统内存不能与设备共享。因此，快速OpenCL应用程序应避免在Linux系统内存和可共享内存区域之间复制数据。有关Linux/OpenCL内存分区的详细信息，请参见如何将DDR3分区为Linux系统和OpenCL。
读缓冲区和写缓冲区OpenCL操作执行复制，应避免。或者，FASTOpenCL应用程序将在共享内存中分配OpenCL缓冲区，并允许主机直接读取和写入底层内存。这可以通过两种方式实现。
通常创建缓冲区，并使用map缓冲区和Unmap缓冲区OpenCLAPI将基础缓冲区内存映射到主机地址空间。有关缓冲区创建信息，请参阅OpenCL缓冲区，并查看缓冲区读/写与map/unmap缓冲区信息的对比。
使用__malloc_ddr or __malloc_msmc，并使用生成的指针创建具有CL_MEM_USE_HOST_PTR属性的缓冲区。有关__malloc_ddr and __malloc_msmc的详细信息，请参见“Alternate Host malloc/free Extension for Zero Copy OpenCL Kernels for details on __malloc_ddr and __malloc_msmc，并查看CL_MEM_USE_HOST_PTR属性用法的OpenCL缓冲区。
25.1 Use MSMC Buffers Whenever Possible
TI SoC通常有一个片上共享内存区域，称为MSMC.MSMC的内存访问延迟比DDR小得多，因此对MSMC缓冲区的操作将比对DDR缓冲区的操作执行得更好。这将是特别真实的计算成本每字节加载是低的，即带宽有限的算法。TI OpenCL实现有一个扩展，允许在MSMC内存中创建全局缓冲区。有关扩展的详细信息，请参阅片内MSMC内存中的快速全局缓冲区。您还可以使用__malloc_msmc内存分配扩展，并将返回的指针传递到缓冲区创建操作，还可以断言CL_MEM_USE_HOST_PTR属性，如上一小节所示。
25.2 Dispatch Appropriate Compute Loads
将计算从主机调度到设备自然需要一定的开销。调度个人，小的计算不会导致提高性能。如果你有灵活性来控制计算的大小，那么一个很好的规则，一个拇指将开销保持在总调度往返的10%以下。当然，您需要知道开销，以便计算最小目标计算负载。
设备派遣的开销有两个方面：根据设备频率和使用SOC平台的原始OpenCL调度开销通常每次调度在60和180微秒之间运行。使用TIOpenCL产品附带的NULL示例可用于测量开销的此组件。在向/从设备通信共享缓冲区时，CPU上的显式缓存操作的成本。该计算具有一些可变性，但公式microseconds = 3 + bytes/8096每个缓冲区的每个调度都是一个合理的近似。

例如，如果内核K接受两个1MB缓冲区作为输入，那么开销的粗略计算将是：180 + (3+1024/8) + (3+1024/8) = 442us，这意味着建议K的最小计算为10x开销或大约4.5毫秒（Ms）。
更喜欢每个工作组有一个工作项的内核，以获得更好的性能，
使用单个工作项创建工作组，并在工作组内使用迭代。
26. Optimization Techniques for Device (DSP) Code
26.1Prefer Kernels with 1 work-item per work-group
为了获得更好的性能，使用单个工作项创建工作组，并在工作组内使用迭代。以这种方式构
造的内核利用了C66xDSP高效执行循环的能力。
本地缓冲区不能用于主机和设备之间的直接通信，但它们非常适合在设备代码中存储临时中间值。在TI SoC上，本地缓冲区位于L2 SRAM存储器中，而全局缓冲区位于DDR 3存储器中。对L2的访问时间比DDR快10倍以上。在编写值时，本地而不是全局的影响会进一步放大。对于算法，如果值被写入缓冲区，并且缓冲区随后被另一个内核或CPU主机使用，则通常最好将值写入本地缓冲区，然后使用OpenCL异步_Work_group_Copy函数将该本地缓冲区复制回全局缓冲区。
下面两个内核执行相同的简单向量加法操作。区别在于，第一个从DDR读取两个输入并将结果写入DDR，其中第二个从DDR读取两个输入并将结果写入本地L2，然后使用异步_Work_GROUP_COMPY将本地缓冲区大容量移动回全局缓冲区。第二个版本比第一个版本快3倍。向量加法的第一个版本：

kernel void VectorAdd(global const short4* a,
                      global const short4* b,
                      global short4* c)
{
    int id = get_global_id(0);
    c[id] = a[id] + b[id];
}
The second version of vector addition, using local buffers
kernel void VectorAdd(global const short4* a,
                      global const short4* b,
                      global short4* c,
                      local  short4* temp)
{
    int id  = get_global_id(0);
    int lid = get_local_id(0);
    int lsz = get_local_size(0);

    temp[lid]  = a[id] + b[id];

    event_t ev = async_work_group_copy(&c[lsz*get_group_id(0)], temp, lsz, 0);
    wait_group_events(1,&ev);
}

26.2 Use async_work_group_copy and async_work_group_strided_copy

上一节演示了async_work_group_copy调用的用法。OpenCL内置函数async_work_group_copy和async_work_group_copy都使用系统DMA操作来执行数据从一个位置到另一个位置的移动。这可能是有益的原因有几个：
顾名思义，异步…函数是异步的，这意味着调用启动数据传输，但在返回之前不等待完成。随后的wait_group_events 调用阻塞，直到数据传输完成。这允许在数据传输的同时执行额外的工作。DDR通过系统DMA写入发生在最佳突发大小，而DSP写入DDR内存没有，因为缓存设置为写入模式上的DSP，以避免错误的共享问题，可能导致不正确的结果。
Avoid DSP writes directly to DDR
See the previous two subsections.
Use the reqd_work_group_size attribute on kernels
如果您按照主机优化技巧“更喜欢每个工作组有一个工作项的内核”，那么您应该用reqd_work_group_size属性对内核进行注释，以通知OpenCL C编译器内核只有一个工作项。这会将信息传递给OpenCL C编译器，否则它将不知道这些信息，并且有许多基于这些知识启用的优化。使用此属性的示例如下所示：

kernel __attribute__((reqd_work_group_size(1, 1, 1)))
void k1wi(global int *p)
{
   ...
}

即使内核在每个工作组具有>1个工作项，此属性对于OpenCLC编译器也是有用的。当然，要使用它，您将断言主机代码将以与在属性中指定的数字相同的本地大小对该内核进行排队。如果内核以与属性中指定的本地大小不同的本地大小进行排队，则运行时将给出明确定义的错误。以下内核使用属性来断言维度1的本地大小为640，维度2的本地大小为480，维度3未使用：

kernel __attribute__((reqd_work_group_size(640, 480, 1)))
void img_alg(global int *p)
{
   ...
}

Use the TI OpenCL extension than allows Standard C code to be called from OpenCL C code
Call existing, optimized, std C code library functions. Or write your own standard C code.
Avoid OpenCL C Barriers
Avoid OpenCL C barriers if possible. Particularly prevent private data from being live across barriers. barrier(), async…(), wait…()
Use the most efficient data type on the DSP
为应用程序选择最有效的数据类型。例如,如果足够,优选“char”类型来表示使用“float”类型的8位数据。这可能会产生重大影响，因为：
它更有效地利用可用的数据带宽它提高了C66xDSP的计算效率，单指令SIMD指令操作的SIMD元素的数量通常倾向于与元素宽度成反比。观察到，如果8位存储对于给定的应用程序是足够的，则更有效的使用使用Char和float计算资源和数据带宽。
不要使用大型向量类型
不要使用向量类型，其中向量类型的大小为>64位。对于宽向量类型，C66xDSP的指令支持有限，因此它们的使用对性能没有好处。
Vector types with total size <= 64 bits may be beneficial, but the benefit is not guaranteed.
Consecutive memory accesses
数据访问模式在生成有效代码中起着关键作用。连续的内存访问是最快的方法。此外，数据流可以发生在不同的数据大小，如。

1. Single Byte ld/st
2. Half Word ld/st
3. Single Word ld/st
4. Double Word ld/st
   在上述列表中，数据流量为存储器操作的上升速率。
   使用双字ld/st是最有利的，因为它具有最高的数据流速率。
   这可用于在不同包装粒度中传输数据。说双字id可以在不同的包装粒度中带来数据，例如：
   • Single 64-bit data
   • Two 32-bit data
   • Four 16-bit data
   • Eight 8-bit data
   根据应用程序的性质，可以选择不同大小的加载。这里的重点是设法实现更高的数据流速率。例如：
   MXN图像表示为‘char’类型的一维数组。该图像由高斯滤波核组成。为了像前面讨论的那样利用SIMD运算，选择了一个向量长度为4的向量。
   为了有效地输入数据，

char* image;
char4 r1_7654, r1_3210;

r1_7654 = vload4(0, image);
r1_3210 = vload4(4, image);
Prefer the CPU style of writing OpenCL code over the GPU style

有大量现有的OpenCL代码可用，而且大多数代码都针对GPU或CPU进行了优化。通常，应用程序将为每个应用程序优化不同的内核。通常，当在TI SoC上执行并使用DSP作为设备时，针对CPU的版本将比针对GPU的版本执行得更好。
27. Typical Steps to Optimize Device Code
下面的流程图描述了典型的顺序，其中应用各种优化步骤来改善C66xDSP上OpenCL内核的性能：

28. Example: Optimizing 1D convolution kernel

在此conv1d示例附带OpenCL产品中，我们展示了如何按步骤优化OpenCL1D卷积内核步骤。
通常，我们需要优化三个区域：指令管线：在低II（启动时间间隔）时，回路是否为软件管线？SIMD效率：是否充分利用了可用的SIMD指令（如C66X）？内存层次结构性能：输入和输出数据是否可以通过双重缓冲扩展到更快的内存，以重叠计算和数据移动？

The example 1D卷积核应用于2D数据的每一行，这些数据可以表示图像、独立通道的集合等。一维卷积核/滤波器尺寸为5x1。我们为非对称过滤器编写了一个通用内核。如果您的滤波器是对称的，欢迎您优化两个乘法。
我们使用一个简单的直进内核的时间作为基线，并报告优化版本相对于该基线的加速比。执行时间由主机端的1920x1080输入映像测量，并报告为微秒。在AM572x(2个DSP核)和K2H(8个DSP核)EVMS上进行了同样的实验。通过改变代码中的NUMCOMPUNITS，我们获得了将内核分配到1、2、4和8个DSP核的性能。每个实验重复5次，每个类别中的最小值在这里报告。
在下面的内容中，我们将通过优化来实现性能的提高。
Driver code setup
将一维卷积核应用于二维图像。我们编写了一个驱动代码，它用随机数据初始化2D图像，并相应地调用OpenCL内核。我们选择高清图像大小1920x1080作为输入和输出。在主机端进行了相同的内核计算，并根据DSP的内核结果对其结果进行了验证，保证了算法的正确性。性能是根据主机端在内核排队之前和内核完成后经过的时间来衡量的。
最初，我们将OpenCL全局大小(1920，1080)划分为NUMCOMPUNITS工作组，每个工作组具有本地大小(1920，1080/NUMCOMPUNITS)，这样每个DSP核心将得到一个工作组。
k_baseline: Ensure correct measurements
TI的OpenCL运行时将延迟地将设备程序加载到程序的内核的第一个队列上，因此从第一个队列到第一个队列的运行时间将更长，以考虑程序的加载。为了从内核性能中删除程序加载开销，我们可以在运行其他内核之前对一个空内核进行排队。
k_baseline: Check software pipelining

我们可以查看程序集输出，以查看编译器是否成功地对内核进行了软件流水线操作。对于最初的二维内核，编译器将在内核周围添加两个隐式循环，以创建OpenCL工作组，并尝试将最内部的循环用于软件管道。使用选项-k运行编译器，以保留程序集代码：

clocl -k ti_kernels.cl
Look at the k_conv1d_5x1 function in ti_kernels.asm, search for SOFTWARE PIPELINE and we see these two lines:
;*      Searching for software pipeline schedule at ...
;*         ii = 7  Schedule found with 5 iterations in parallel

因此原始内核已经是软件流水线化的。每7个周期开始在最内侧维度上的循环迭代。
k_loop: Improve software pipelining

仔细看一下基线源代码，我们就会发现循环处理的是一些边界条件。如果我们可以将这些边界迭代从主循环中剥离出来，那么主循环可能被安排在较低的II上。为此，我们还需要将OpenCL内核的工作空间从2D减少到1D，以便内核代码中最内部的循环变得显式化。核k_loop是这种转换的结果。从组装文件中，我们可以看到主循环是调度在ii=3，这意味着每3个周期启动一次迭代：

;*      Searching for software pipeline schedule at ...
;*         ii = 3  Schedule found with 10 iterations in parallel

总结：在内核中显式地定义圆环，将OpenCL内核的工作空间从2D减少到1D，去掉边界条件，或者删除边界检查，这样就可以填充输入数据或减小输出大小，从而使边界条件消失。与基线版本相比，使用精简的II，我们并没有看到执行带来的性能改善。一个可能的原因是，由于缓存错误而导致的软件管道中断已经占据了执行的主导地位。现在是为内存层次结构进行优化的时候了。在此之前，让我们看看是否可以优化C66 DSP上可用的SIMD功能。
k_loop_simd: Improve software pipelining with SIMDization

有时，编译器可能无法auto-SIMDize循环.我们可以查看所涉及的内存访问和计算，并执行SIMD手动。由于OpenCL C向量语义，我们必须假设每一行在8字节的边界上正确地对齐，以便使用向量类型Float 2。首先我们对内存访问和计算进行SIMDISE，然后我们寻求在寄存器中流水线加载值的机会。K_loop_SIMD是SIMD化的结果。从程序集中可以看到，每5个周期启动一次展开迭代(对应于两个基线迭代)：

;*      Searching for software pipeline schedule at ...
;*         ii = 5  Schedule found with 5 iterations in parallel

总结：
Unroll col-loop by a factor of 2 by hand
Data layout requirement: each row is aligned on 8-byte double word boundary
SIMDize loads and stores
SIMDize computation
Pipeline loaded values in registers if possible
用手将圆环展开2倍
数据布局要求：每一行按8字节双对齐
k_loop_db: EDMA and double buffer k_loop
TI的OpenCL实现为OpenCL本地内存在每个核心上指定了L2 SRAM的一部分。我们可以使用EDMA将数据从全局缓冲区(DDR)移动到本地缓冲区(L2)，在本地缓冲区上执行计算，然后将本地缓冲区(L2)的结果存储回全局缓冲区(DDR)。OpenCL C内核语言内置了我们映射到TI的EDMA例程的异步_Work_group_()函数。为了更好地利用EDMA的异步特性，我们使用了双缓冲(乒乓)来有效地重叠数据移动和计算。
对于这个特定的内核，每一行都需要COLSsizeof(float) + COLSsizeof(float)字节进行输入和输出。使用双缓冲，每行输入和输出都需要16cols字节。给定我们选择的Cols=1920，我们可以在128 KB的本地内存中容纳最多4行，或者在768 KB的本地内存中容纳最多25行：

4  * (2 * (1920*4 + 1920*4)) <= 128 * 1024
25 * (2 * (1920*4 + 1920*4)) <= 768 * 1024

为了确保双缓冲管道至少执行几次，比如说8，我们可以将BLOCK_HEIGHT to ROWS / NUMCOMPUNITS / 8 + 1。在内核中，在计算本地存储器中的图像行的当前块之前，我们将下一块行插入到具有EDMA的本地存储器中。另一个转换是内核现在显式地迭代通过行维度，因为需要双缓冲。因此，我们需要将所需的内核工作组大小设置为(1，1，1)。在主机代码中，我们只需要指定工作组的数量，在将ND范围内核排队时，我们使用计算单位的数量。我们在内核中添加了三个附加的参数：块高度、用于输入的本地缓冲区和用于输出的本地缓冲区。OpenCL运行时自动分配本地缓冲区，OpenCL应用程序代码只需要指定大小。在所有这些变换过程中，我们看到，非正弦化的K_LOOP_DB不仅优于基线K_LOOP，而且与K_LOOP_SIMD进行了比较。

With all these transformation, we see that non-SIMDized k_loop_db outperforms not only baseline k_loop, but also SIMDized k_loop_simd.
Summary
1.	Require 8-byte alignment for each row
2.	Determine the block height for double buffering
3.	Set required work group size to (1,1,1) for kernel
4.	Set OpenCL workspace to (NUMCOMPUNITS, 1, 1), each work group will figure out which rows to work on

5. Double buffer with EDMA on input and

output, computation only loads from and stores to local buffers
k_loop_simd_db: EDMA and double buffer k_loop_simd
We apply the same EDMA and double buffering transformation on k_loop_simd as above. Now we see similar performance improvements upon k_loop_simd.
k_loop_simd_db_extc: Use external C function for k_loop_simd_db

虽然我们可以在OpenCL C语言中完全处理这个例子，但有时OpenCL C对我们C66DSP的表现力有限制。例如，C66DSP可以执行比异步工作组*()OpenCLC内置函数更多的EDMA传输模式，C66DSP支持非对齐SIMD负载和存储。当这些限制确实影响用户应用程序时，我们可以在标准C函数中使用它们，并从OpenCL C代码中调用它们。
我们将此版本用作示例如何将标准C函数纳入OpenCL。我们将K环SIMDDB的主体移动到外部C函数中，并将OpenCL函数视为一个简单的包装函数。外部C函数使用C66C编译器编译，您可以使用C66C内部函数。同样，您可以重新利用自己或TI开发的现有优化的C实现和库。当然，这是TI的扩展，不适用于其他供应商的OpenCL平台。
c_loop_simd_db_extc() in k_ext.c中是重写C函数。注意EDMAMGR函数和C66SIMD内部函数的显式使用。使用此版本，我们获得了稍微更好的性能。

Summary
1.	Move kernel body to an external standard C function
2.	Use EdmaMgr_*() functions directly, cover non-consecutive transfers
3.	Use C66 C SIMD intrinsic built-in functions, cover non-aligned SIMD loads and stores
4.	Link separately compiled C object back to kernel executable

29. Example: Optimizing 3x3 Gaussian smoothing filter
    本节描述了一种逐步优化C66xDSP的3x3高斯平滑滤波核的方法。
    高斯滤波器用作平滑滤波器。通过将NxN图像窗口与NxN高斯核卷积并获得加权和来应用滤波器。此处可提供更多关于过滤器的：http://homepages.inf.ed.ac.uk/rbf/HIPR2/gsmooth.htm我们使用的内核大小是3x3内核。设3x3图像窗口，b为3x3高斯核。通过对A和B进行卷积来施加滤波器，并且A以滑动窗口的方式获得。
    Natural C Code

The first listing is a snippet of C code for convolution:
 
	float image[image_size];
float gaussian_kernel[9];
float weight;
float filtered_image[(image_width - 2) * (image_height - 2)];

for (i = 1; i < img_height - 1; i++)
{
    for (j = 1; j < img_width - 1; j++)
    {
	sum = 0;
	for (p = 0; p < 3; p++)
        {
	    for (q = 0; q < 3; q++)
            {
		sum += image[(i + p) * img_width + j + q] * 
                       gaussian_kernel[p * 3 + q];
            }
	}
	sum /= weight;
	filtered_image[(i - 1) * img_width + j] = sum;
    }
}

Optimizing for DSP
An OpenCL C kernel for convolution. Note that the types are float.

	// Serves as bounds check
bool OutsideImage(int2 pos, int width, int height)
{
 if (pos.x < 1 || pos.y < 1)
    return true;

 if (pos.x >= width - 1 || pos.y >= height - 1)
    return true;

 return false;
}

kernel void gaussian_filter (global float* image,
                             global float* filtered_image,
                             global float* gaussian_kernel,
                             global int*   image_dims
                                    float  weight)
{
   const int image_height = image_dims[0];
   const int image_width = image_dims[1];
    
   const int global_x = get_global_id(0);
   const int global_y = get_global_id(1);
   const int2 pixel_pos = { global_x, global_y };

   if (OutsideImage(pixel_pos, image_width, image_height))
      return;

   float sum = 0;
   int index = 0;
   int2 pos;

   /* 3x3 Convolution */

   for(int y= -1; y<=1; y++)
      for(int x=-1; x<=1; x++)
      {
         pos = pixel_pos + (int2)(x,y);
         sum += gaussian_kernel[index++] * image[pos.y * image_width + pos.x];
      }

   sum /= weight;

   filtered_image[global_y * img_width + global_x] = sum; 
}

Step 1: Initial optimization for DSP:

•	Convert the float type to uchar
 
	// Serves as bounds check
bool OutsideImage(int2 pos, int width, int height)
{
    if (pos.x < 1 || pos.y < 1)
        return true;

    if (pos.x >= width - 1 || pos.y >= height - 1)
        return true;

    return false;
}

kernel void gaussian_filter (global uchar* image,
                             global uchar* filtered_image,
                             global char*  gaussian_kernel,
                             global int*   image_dims,
                                    short   weight)
{
   const int image_height = image_dims[0];
   const int image_width = image_dims[1];
    
   const int global_x = get_global_id(0);
   const int global_y = get_global_id(1);
   const int2 pixel_pos = { global_x, global_y };

   if (OutsideImage(pixel_pos, image_width, image_height))
      return;

   short sum = 0;
   int index = 0;
   int2 pos;

   /* 3x3 Convolution */
   for(int y= -1; y<=1; y++)
      for(int x=-1; x<=1; x++)
      {
         pos = pixel_pos + (int2)(x,y);
         sum += gaussian_kernel[index++] * image[pos.y * image_width + pos.x];
      }

   sum /= weight;

   filtered_image[global_y * img_width + global_x] = (uchar) sum; 
}

Step 2:

•	Switch to using vector types to take advantage of vector instructions available on the DSP
•	Annotate the kernel with a work-group size attribute

	inline int
dot_product (uchar4 mask, uchar4 data)
{
    int sum = 0;
    sum = (int) (mask.s0 * data.s0 +
		 mask.s1 * data.s1 + mask.s2 * data.s2 + mask.s3 * data.s3);
    return sum;
}

kernel __attribute__ ((reqd_work_group_size (1, 1, 1)))
void gaussian_filter (global const uchar4* restrict imgin_ptr,
		      global       uchar4* restrict imgout_ptr,
		      short                         width,
		      short                         pitch,
		      global const uchar*           kernel_coefficient,
		      short                         shift)
{

    int i;
    int sum0, sum1, sum2;
    int sum3;

    uchar4 mask1_0, mask2_0, mask3_0;
    uchar4 mask1_1, mask2_1, mask3_1;

    uchar4 r1_3210;
    uchar4 r2_3210;
    uchar4 r3_3210;
    uchar4 r1_5432;
    uchar4 r2_5432;
    uchar4 r3_5432;

    uchar8 r1_76543210, r2_76543210, r3_76543210;

    mask1_0 =
	(uchar4) (kernel_coefficient[0], kernel_coefficient[1],
		  kernel_coefficient[2], 0);
    mask2_0 =
	(uchar4) (kernel_coefficient[3], kernel_coefficient[4],
		  kernel_coefficient[5], 0);
    mask3_0 =
	(uchar4) (kernel_coefficient[6], kernel_coefficient[7],
		  kernel_coefficient[8], 0);

    mask1_1 =
	(uchar4) (0, kernel_coefficient[0], kernel_coefficient[1],
		  kernel_coefficient[2]);
    mask2_1 =
	(uchar4) (0, kernel_coefficient[3], kernel_coefficient[4],
		  kernel_coefficient[5]);
    mask3_1 =
	(uchar4) (0, kernel_coefficient[6], kernel_coefficient[7],
		  kernel_coefficient[8]);

    for (i = 0; i < width; i += 1)
    {
	  r1_76543210 = vload8 (i, imgin_ptr);
	  r1_76543210 = vload8 (pitch + i, imgin_ptr);
	  r1_76543210 = vload8 (2 * pitch + i, imgin_ptr);

	  r1_3210 = (uchar4) (r1_76543210.s0123);
	  r2_3210 = (uchar4) (r2_76543210.s0123);
	  r3_3210 = (uchar4) (r3_76543210.s0123);

	  sum0 = (dot_product (mask1_0, r1_3210) +
		  dot_product (mask2_0, r2_3210) +
		  dot_product (mask3_0, r3_3210)) >> shift;

	  sum1 = (dot_product (mask1_1, r1_3210) +
		  dot_product (mask2_1, r2_3210) +
		  dot_product (mask3_1, r3_3210)) >> shift;

	  r1_5432 = (uchar4) (r1_76543210.s2345);
	  r2_5432 = (uchar4) (r2_76543210.s2345);
	  r3_5432 = (uchar4) (r3_76543210.s2345);

	  sum2 = (dot_product (mask1_0, r1_5432) +
		  dot_product (mask2_0, r2_5432) +
		  dot_product (mask3_0, r3_5432)) >> shift;

	  sum3 = (dot_product (mask1_1, r1_5432) +
		  dot_product (mask2_1, r2_5432) +
		  dot_product (mask3_1, r3_5432)) >> shift;

	  imgout_ptr[i] = (uchar4) (sum0, sum1, sum2, sum3);
    }
}

Step 3: Use double buffering to overlap data movement with computation
Pseudo-code for a double-buffered version of the OpenCL C kernel:

#define ARRAY_SIZE n
#define NUM_BATCHES n_b
#define BATCH_SIZE (ARRAY_SIZE / NUM_BATCHES)	// This is the size of a single buffer
#define LOCAL_SIZE (BATCH_SIZE * 2)	// This is the size of the double buffer

kernel __attribute__ ((reqd_work_group_size (1, 1, 1)))
void gaussian_filter (global const uchar4 * restrict imgin_ptr,
	              global uchar4 * restrict imgout_ptr,
	              short width,
	              short pitch, 
                      global const uchar * kernel_coefficient, 
                      short shift)
{
    //Initialize the required variables

    //Copy content in the double buffer

    //Compute for the buffer in batch 1

    for (batch = 0; batch < NUM_BATCHES - 2; batch++)
    {
      if (batch % 2 == 0)
      {
	  Copy content in buffer batch 1 || Compute for buffer in batch 2
      }
      else
      {
	Copy content in buffer batch 2 || Compute for buffer in batch 1
      }
    }

    if (batch % 2 == 0)
    {
        Copy content in buffer batch 1 || Compute for buffer in batch 2
    }
    else
    {
        Copy content in buffer batch 2 || Compute for buffer in batch 1
    }
}

Now, we have an optimized OpenCL C kernel for the DSP. Note that the kernel is a generic OpenCL C kernel and can be compiled/run on any OpenCL device.
Performance Improvement
Description	Performance in cycles per pixel
Generic OpenCL C kernel	12.0
OpenCL C kernel optimized for DSP	5.0
30. performance Data
The table below show improvements obtained from optimizing a set of image processing kernels for the C66x DSP device using the techniques described in this chapter.
Kernel	Generic OpenCL C (cycles/pixel)	Optimized OpenCL C (cycles/pixel)	Improvement (times faster)
Convolution	12	5	2.40
Histogram	56	1.75	32.00
X_Gradient	12.4	1.25	9.92
Edge Relaxation	530	48	11.04

31.Debug with printf
虽然它必须手动插入到代码中，但是一个简单的printf函数可以帮助您调试程序进度、感兴趣的数据的值等等。您也可以使用printf调试OpenCL应用程序。
Host side OpenCL application code
显然，只要您的主机编译器支持它（GCC），您可以将printf置于主机侧OpenCL应用程序代码中。
DSP side OpenCL kernel code
尽管TI的OpenCL实现目前处于我们支持的SoC上最OpenCL版本1.1上，但我们确实支持OpenCL1.2版功能，printf，如OpenCLv1.2规范第6.12.13节所述。来自DSP侧的printf的输出被重定向到主机侧，打印在stdout中，例如，启动OpenCL应用程序的Linux窗口/终端。
您不仅可以将printf放入OpenCLC内核代码中，如OpenCL1.2规范中所述，还可以将printf放入链接到OpenCLC内核的标准C代码中，如TI扩展中所示（从OpenCLC代码调用标准C代码）。它们将全部打印在主机侧。在OpenCLC内核中使用printf时，不需要包括任何头文件，在链接到的标准C代码中使用时，需要像通常一样包括stdio.h。
注意，在printf的格式字符串中，TI的实现现在支持OpenCL1.2规范中描述的所有特性。例如，现在支持表示OpenCL矢量类型的%v。一个已知的问题是printf的使用(“%s\n”、“String”)；导致clocl/clang断言失败。您可以通过简单地使用printf(“string\n”)来避免这一点；

注意，DSP内核的printf的输出在默认情况下仍然与DSP核心数字相加。OpenCL1.2规范不需要它。因此，现在可以使用一个新的环境变量TI_OCL_PRINTF_COREID在显示DSP核号和不显示DSP核号之间切换。如果使用TI_OCL_PRINTF_COREID=0，则不会显示DSP核心号码。

32. Debug with gdb
    Host side gdb
    您可以以与调试其他主机侧应用程序相同的方式调试您的主机侧OpenCL应用程序。在编译过程中，您需要使用标志“-g”。有关详细信息，请参阅GDB文档。
    如果gdb未预装在您的主机文件系统上，您需要下载软件包并安装它，或者您需要下载gdb发行版、构建和安装在您的文件系统上。
    DSP side debug with host side client gdbc6x
    DSP侧内核代码可以通过托管调试器gdbc6x进行调试。在OpenCL应用程序中调试内核代码的过程如下。您需要两个Windows/控制台，一个用于运行OpenCL应用程序的窗口，另一个用于调试DSP侧内核。
    在窗口1中，在运行应用程序之前设置环境变量TI_OCL_DEBUG，例如，TI_OCL_DEBUG=GDB。/您的_OCL_APP如果使用Bash应用程序运行后，在启动内核到DSP之前，OpenCL运行时将在窗口1中打印GDBC6x命令，例如，GDBC6x-Q-IEX“目标远程/dev/gdbtyp0”-IEX“设置确认”-IEX“符号-文件/usr/share/ti/OpenCL/dsp.out”-IEX“添加-符号-文件/TMP/opencl7mNBld.out0x86000000”-IEX“B出口”-IEX“B矢量添加”将GDBC6X命令复制并粘贴到窗口2中，运行它点击窗口1中的任何键1开始调试窗口2中的内核

1.	In window 1, set environment variable TI_OCL_DEBUG before running application, for example, TI_OCL_DEBUG=gdb ./your_ocl_app if you use bash
2.	Once the application is running, before launching your kernel to DSP, OpenCL runtime will print out a gdbc6x command in window 1, for example, gdbc6x -q -iex “target remote /dev/gdbtty0” -iex “set confirm off” -iex “symbol-file /usr/share/ti/opencl/dsp.out” -iex “add-symbol-file /tmp/opencl7mNBld.out 0x86000000” -iex “b exit” -iex “b VectorAdd”
3.	Copy and paste the gdbc6x command into window 2, run it
4.	Hit any key in window 1
5.	Start debugging the kernel in window 2
The following are the sample output of window 1:
root@am57xx-evm:~/oclexamples/vecadd# TI_OCL_DEBUG=gdb ./vecadd
DEVICE: TI Multicore C66 DSP

Offloading vector addition of 8192K elements...

gdbc6x -q -iex "target remote /dev/gdbtty0" -iex "set confirm off" -iex "symbol-file /usr/share/ti/opencl/dsp.out" -iex "add-symbol-file /tmp/openclXmObdu.out 0x86000000" -iex "b exit" -iex "b VectorAdd"
Press any key, then enter to continue
c
Kernel Exec : Queue  to Submit: 4 us
Kernel Exec : Submit to Start : 45 us
Kernel Exec : Start  to End   : 83717229 us

Success!
and window 2:
root@am57xx-evm:~# gdbc6x -q -iex "target remote /dev/gdbtty0" -iex "set confirm off" -iex "symbol-file /usr/share/ti/opencl/dsp.out" -iex "add-symbol-file /tmp/openclXmObdu.out 0x86000000" -iex "b exit" -iex "b VectorAdd"
Remote debugging using /dev/gdbtty0
0xfeabec64 in ?? ()
Reading symbols from /usr/share/ti/opencl/dsp.out...done.
add symbol table from file "/tmp/openclXmObdu.out" at
    .text_addr = 0x86000000
Reading symbols from /tmp/openclXmObdu.out...done.
Breakpoint 1 at 0xfea53254: file exit.c, line 64.
Breakpoint 2 at 0x8600000c: file /tmp/openclXmObdu.cl, line 4.
(gdb) continue
Continuing.

Breakpoint 2, VectorAdd () at /tmp/openclXmObdu.cl:4
4   {
(gdb) list
1   kernel void VectorAdd(global const short4* a,
2                         global const short4* b,
3                         global short4* c)
4   {
5       int id = get_global_id(0);
6       c[id] = a[id] + b[id];
7   }
(gdb) break 6
Breakpoint 3 at 0x8600008a: file /tmp/openclXmObdu.cl, line 6.
(gdb) cont
Continuing.

Breakpoint 3, $C$L6 () at /tmp/openclXmObdu.cl:6
6       c[id] = a[id] + b[id];
(gdb) print a[0]
$1 = {0, 4, 8, 12}
(gdb) print b[0]
$2 = {0, 4, 8, 12}
(gdb) print c[0]
$3 = {0, 0, 0, 0}
(gdb) next
7   }
(gdb) print c[0]
$4 = {0, 8, 16, 24}
(gdb) info locals
dim = 0
dim = 0
a = 0x80000000
b = 0x82000000
c = 0x84000000
id = 0
(gdb) delete 3
(gdb) delete 2
(gdb) cont
Continuing.
^C
Program received signal SIGTRAP, Trace/breakpoint trap.
0xfea7ec04 in $C$RL54 ()
    at /home/gtbldadm/processor-sdk-linux-daisy-build/build-CORTEX_1/arago-tmp-external-linaro-toolchain/sysroots/am57xx-evm/usr/share/ti/ti-sysbios-tree/packages/ti/sysbios/knl/Idle.c:72
72  /home/gtbldadm/processor-sdk-linux-daisy-build/build-CORTEX_1/arago-tmp-external-linaro-toolchain/sysroots/am57xx-evm/usr/share/ti/ti-sysbios-tree/packages/ti/sysbios/knl/Idle.c: No such file or directory.
(gdb) quit
Detaching from program: , Remote target
Ending remote debugging.
root@am57xx-evm:~#
33.Debug with CCS

您还可以使用CodeComposerStudio(CCS)中的调试功能调试OpenCL DSP侧代码。为此，您将需要一个额外的硬件，例如XDS560v2这样的仿真器来连接到EVM的JTAG端口。
只有当您的TI OpenCL安装版本为01.01.06.00或更高时，此功能才可用。
Connect emulator to EVM and CCS
首先，通过JTAG端口将仿真器连接到EVM。接下来，取决于您是否希望CCS通过以太网或USB与仿真器通信，您可以将以太网电缆连接到仿真器或将仿真器连接到主机计算机，在此您使用USB电缆运行CCS。将仿真器加电。接下来，启动CCS，“查看->目标配置”，创建一个“用户定义”“新目标配置”。对于“连接的连接”，从列表中选择仿真器模型，例如“频谱数字XDS560V2STMUSB仿真器”。对于“板或装置”，从列表中选择EVM模型，例如“66AK2H”、“TMS320C6678”或“GPPEVM_AM572X”。随着CCS版本的推移，设备支持通常随时间增加。如果您看不到正在列出的设备，且您的CCS安装已过时，请更新至最新的CCS版本。选择“连接的连接”和“板或装置”后，请保存配置和测试连接，以确保CCS可以通过仿真器与EVM进行对话。如果您选择通过以太网与仿真器进行CCS对话，则可以使用带有CCS安装的配置实用程序来查找仿真器的IP地址，例如启动“XDS560v2STM配置实用程序”，然后在““ETH””选项卡下单击““查找以太网设备””。或者您可以使用它的MAC地址查找仿真器的IP地址。一旦IP地址已知，请单击““先进的””选项卡，单击““仿真器””，输入““仿真器IP地址””、“保存”和“测试连接”。最后，右键单击刚才所做的目标配置和“启动选定的配置”。一旦启动，您应该在列表中看到DSP内核0（“C66xx_DSP1”），连接到DSP内核0并恢复运行。
Debug DSP side code with CCS
带有CCS的DSP侧代码调试采用与GDBC6X的调试类似的步骤。在运行应用程序之前将环境变量TI_OCL_DEBUG设置为“CCS”，例如，TI_OCL_DEBUG=CCS。/您的_OCL_APP如果使用Bash。
在应用程序正在运行之前，在将内核启动到DSP之前，OpenCL运行时将打印您在继续之前应执行的CCS指令列表，例如，CCS挂起DSP内核0CCS负载符号：/TMP/openclwsNYLl.out，代码偏移：0x86000000CCS添加符号：/usr/share/ti/OpenCL/dsp.out，无代码偏移CCS添加断点:矢量添加CCS恢复DSP核心0按任何键，然后输入继续您可能需要将内核可执行文件和dsp.out复制到您运行CCS的主机文件系统（也许是内核源代码），以便在调试时CCS可以显示它。
一旦已加载/添加了CCS符号并设置了CCS断点，则当您按下一个键继续在主机侧时，CCS应停止内核功能条目上的DSP内核0。在这一点上，您可以通过代码、检查内存内容、检查变量值，就像通常使用CCS调试一样。
34.Debug with dsptop
dsptop是一种与Linux实用程序top类似的TI实用程序，可用于调试哪些DSP核心参与计算、OpenCL缓冲区的内存使用以及带有时间戳的内核活动，如工作组启动、工作组完成和缓存操作。和gdbc6x一样，dsptop的使用也需要两个窗口/控制台：首先在窗口1中启动dsptop，然后在Window 2中启动OpenCL应用程序。有关dsptop使用的详细信息，可以通过运行dsptop-h和这个dsptop维基页面找到。
35. Profiling
您可以使用通用分析工具（如gprof）将OpenCL应用程序配置为任何其他应用程序。在此，我们解释如何在OpenCL命令队列中配置命令。
此外，还可以通过AET库配置诸如L2高速缓存未命中和流水线停顿之类的硬件事件。有关详细信息，请参见“分析硬件事件”部分。
35.1 Host Side Profiling
要在主机侧配置命令，需要在创建命令队列时指定“CL_QUEUE_PROFILING_ENABLE”属性。当命令被排队时，OpenCL运行时将在nano-seconds中记录主机端时间戳，并开始执行并完成执行。用户代码可以使用“clGetEventProfilingInfo”API在命令enqueue时间返回的相应事件上查询这些时间戳。
35.2 DSP Side Profiling
在DSP端，OpenCL运行时还记录预定义的OpenCL活动的时间戳，并使用轻量级ULM（使用和负载监视器）将这些数据传回主机。这些预定义的活动包括工作组执行的开始和完成，以及DSP缓存一致性操作。在主机端，用户需要使用dsptop实用程序来检索信息。例如，在运行dsptop之后，您可能需要两个windows/console并运行OpenCL应用程序。

1.	In window 1, run dsptop -l last
2.	In window 2, launch your OpenCL application, wait for it to finish
3.	Back in window 1, type “q” to quit dsptop. dsptop should print out the information sent back from the DSP side.
Details about usage of dsptop can be found by running dsptop -h and by this dsptop wikipage .

35.3 Profiling Hardware Events
分析硬件事件是通过CTools AET库提供的一种有用的功能。从TI OpenCL产品v1.1.14开始，此功能将集成到OpenCL运行时中。请参阅下面AET分析事件部分中的所有失速和内存事件。不能描述未以AET_EVT_STACK_或AET_EVT_MEM_为前缀的事件。要分析硬件事件，有三种选择：

1. Profile All Possible Events
2. Profile a Select Few Events
3. Manually Profile 1 or 2 Events
   Profile All Events
   To profile all hardware events, run the profiling script as follows:
   /usr/share/ti/opencl/profiling/oclaet.sh [-g] oclapp

e.g oclapp is ./vecadd

当可执行文件正在运行时，原始分析数据将记录到相对于当前目录的分析/aetdata.txt中。在分析脚本完成后，它运行一个Python脚本，该脚本形成JSON表、HTML表，以及可选的每个内核的分析信息IF-G选项的绘图。请注意，生成图可能需要安装其他Python包。

Profile Select Events
To profile selected hardware events, run the profiling script as follows:
/usr/share/ti/opencl/profiling/oclaet.sh oclapp event_type event_number [event_number ...]
# event_type: 1 for stall events, 2 for meeory events
# event_number is the event offset from AET_GEM_STALL_EVT_START, if
#   profiling stall cycles, or from AET_GEM_MEM_EVT_START, if profiling
#   memory events.  At each script run, you can only profile one or more
#   events of the same type.
Manually Profile 1 or 2 Events
To manually profile 1 stall event, or 1 to 2 memory events, simply set the environment variables described in Environment Variables, as follows.
TI_OCL_EVENT_TYPE=1 TI_OCL_EVENT_NUMBER1=13 TI_OCL_STALL_CYCLE_THRESHOLD=0 ./vecadd
TI_OCL_EVENT_TYPE=2 TI_OCL_EVENT_NUMBER1=11 TI_OCL_EVENT_NUMBER1=12 ./vecadd
Note that profiling data is appended to profiling/aetdata.txt at each profiling run. If a fresh profiling is needed, remove profiling/aetdata.txt before profiling run.
Analyzing Profiling Data
If you manually profile, you will have to run the python script to obtain the JSON file or html table.
python /usr/share/ti/opencl/profiling/oclaet.py -t -g profiling/aetdata.txt
The -t flag and -g flags tell the script to produce an html table and matplot plot of profiling data, respectively. If neither of these flags are specified, then only the json file of raw counter data will be formed. This json is easier to read than the raw data dump in profiling/aetdata.txt. The current format for raw data is:
EVENT_TYPE            # the event type
EVENT_NUMBER1         # the first event number (offset from base AET event)
EVENT_NUMBER2         # the second event number (offset from base AET event)
STALL_CYCLE_THRESHOLD # the stall cycle threshold
Core number           # number of core
Counter0_Value        # hardware counter 0 value: memory event 1
Counter1_Value       # hardware counter 1 value: memory event 2 or stall event
~~~~End Core          # End of core data for Core number
...                   # MORE CORE DATA CAN FOLLOW THIS
EVENT_TYPE
EVENT_NUMBER1
EVENT_NUMBER2
STALL_CYCLE_THRESHOLD
Core number
Counter0_Value
Counter1_Value
~~~~End Core
VectorAdd             # Kernel Name
---End Kernel         # Ends Kernel Data
Profiling Data Plotting Requirements
Plotting profiling data with python script requires matplotlib, pandas, and seaborn python packages to be installed. If they are not already installed on your system, you can follow the instructions below.
which pip
# install pip if it is not already on your system
wget https://bootstrap.pypa.io/get-pip.py
python get-pip.py

pip install matplotlib
pip install pandas
pip install seaborn

35.4 OpenCL on TI-RTOS

从版本01.01.09.01开始，OpenCL运行时支持运行TI-RTOS作为主机的Cortex-A15S。本章介绍了OpenCLRTOS软件包，展示了如何运行实例，并讨论了OpenCLRTOS软件包开发OpenCL应用程序的过程。