精华内容
下载资源
问答
  • c64_tools TI C64+ DSP 加载程序和 IPC 实用程序。 关于 该软件包提供了一些工具和库... CodeComposerStudio / DSPBIOS 5.xx. Licensed under the terms of the GPL. "c64.ko" -- A Linux kernel module that handle
  • 展示-源码

    2021-02-25 13:26:02
    打开CodeComposerStudio。 生成程序。 连接设备。 打开UnitFlash程序并找到设备。 在UnitFlash中(查找程序按钮)-> Flash Image->浏览-> display_CC1352R1_LAUNCHXL_tirtos_ccs(dir)->调试(dir)->选择...
  • TMS320C62xx系列DSP的集成开发环境CodeComposerStudio(简称CCS)支持标准C语言和汇编混合编程的方式编程,为了提高编解码算法的效率,本文对ITU_T的标准G.729ab的C语言原码进行汇编指令优化设计。同时,对于上层编...
  • 2020-10-19

    2020-10-19 16:17:00
    TI OpenCL 用户指南3 Optimization Tips OpenCL应用程序由主机应用程序和一组设备内核组成。主机代码和设备代码都有优化技术。存在跨越主机和设备之间的边界的一些技术。本节提供了编写OpenCL应用程序的技巧,该...

    TI OpenCL 用户指南3

    1. Optimization Tips
      OpenCL应用程序由主机应用程序和一组设备内核组成。主机代码和设备代码都有优化技术。存在跨越主机和设备之间的边界的一些技术。本节提供了编写OpenCL应用程序的技巧,该应用程序执行得很好。它以DSP为加速器设备,以TI SoCs为目标。这些提示被组织成基于尖端适用的部分,即主机或设备。
    2. 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内核的性能:

    1. 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的异步特性,我们使用了双缓冲(乒乓)来有效地重叠数据移动和计算。
    对于这个特定的内核,每一行都需要COLS
    sizeof(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
    
    1. 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
    
    1. 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核心号码。

    1. 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应用程序的过程。

    展开全文
  • CCS提供了配置、建立、调试、跟踪和分析程序的工具,它便于实时、嵌入式信号处理程序的编制和测试,它能够加速开发进程,提高工作效率。本文件为ccs的中文教程
  • CCS(Code Composer Studio)简介

    千次阅读 2018-04-25 20:47:03
    CCS,英文全称是Code Composer Studio,它是美国德州仪器公司(Texas Instrument,TI)出品的代码开发和调试套件。 英文全称:Code Composer Studio 中文译名:代码调试器,代码设计套件 解释: ...

    CCS,英文全称是Code Composer Studio,它是美国德州仪器公司(Texas Instrument,TI)出品的代码开发和调试套件。

    英文全称:Code Composer Studio

    中文译名:代码调试器,代码设计套件

    解释:

    CCS的全称是Code Composer Studio,它是美国德州仪器公司(Texas Instrument,TI)出品的代码开发和调试套件。TI公司的产品线中有一大块业务是数字信号处理器(DSP)和微处理器(MCU),CCS便是供用户开发和调试DSP和MCU程序的集成开发软件。

    经典的版本号是CCS 3.1、CCS 3.3,最新版本号已经更新到了v6,兼容64位系统。

    Code Composer Studio IDE 提供强健、成熟的核心功能与简便易用的配置和图形可视化工具,使系统设计更快。

    开发周期中的功能:应用设计 - 包括 DSP BIOS、参考框架和更新顾问

    编码与编译 - 包括 C/C++ 和汇编语言以及 CodeWright 集成编辑器

    调试 - RTDX 快速模拟和连接/断开连接

    分析与调优 - 包括实时分析、编译器分析和回卷

    展开全文
  •   在对MSP430进行项目开发时使用了CCS这款集成开发软件,但是它安装成功后没有默认的Debug快捷键,感觉很不方便,下面直接上图对CCS的快捷键进行设置 1、Window --> Prefrences 2、General -->...

      在对MSP430进行项目开发时使用了CCS这款集成开发软件,但是它安装成功后没有默认的Debug快捷键,感觉很不方便,下面直接上图对CCS的Build快捷键进行设置

    1、Window --> Prefrences
    在这里插入图片描述
    2、General --> keys
    在这里插入图片描述
    3、在上一步中我们进入到快捷键的设置页面,下面对Build的快捷键进行重置
    在这里插入图片描述
    4、配置完成
     说明:若有其他需求,请读者根据自己的需求进行个性化配置
     
    在这里插入图片描述

    展开全文
  • 我们在Windows10系统中安装了DSP编译软件Code Composer Studio后,在打开此软件后,经常会发生闪退现象。 这是因为CCS软件为英文软件,对中文的支持度非常低。我在当初在windows10上安装了各个版本的CCS,结果全部...

    我们在Windows10系统中安装了DSP编译软件Code Composer Studio后,在打开此软件后,经常会发生闪退现象。

    这是因为CCS软件为英文软件,对中文的支持度非常低。我在当初在windows10上安装了各个版本的CCS,结果全部都是闪退,没有一个版本能够使用。当时我的安装路径是英文路径,百思不得其解。

    后来机缘巧合下发现,这是因为我的windows的账户名是中文,然后我在windows设置——账户管理——新建本地账户——把新账户命名为英文账户——重启电脑——选择打开新账户,然后奇迹发生了,打开软件后CCS软件不再闪退,能够正常使用啦。

    所以,遇到CCS闪退的正确解决方案是:

    (1)查看自己的安装路径下是否有中文,有的话,要把路径改成英文

    (2)查看自己的window账户是不是英文,如果是中文账户,那么就再添加一个英文账户,(注意:把已有的中文Windows账户名改成英文账户并不起作用,本人亲自试过,所以建议重新创建一个英文账户),创建步骤如下:windows设置——账户管理——新建本地账户——把新账户命名为英文账户——重启电脑——选择打开新账户;

    (3)在打开软件后,软件会提示工作空间的保存位置,这里工作空间的路径一定要设置成英文!注意一定要设置成英文!

    展开全文
  • 很多人都说Code Composer Studio无法用gcc编译,各种报错,balabala。之前参加比赛用msp430(虽然最后没用,比赛也重在参与了,2333),下了CCS10.1,发现确实无法用gcc编译,具体是怎么回事呢,通过报错信息,找到...
  • DSP实验报告—实验1

    千次阅读 2019-05-04 13:56:00
    软件集成开发环境(CodeComposerStudio5.x):完成系统的软件开发,进行软件和硬件仿真调试。它也是硬件调试的辅助手段。 开发系统(ICETEK-XDS100V2+):实现硬件仿真调试时与硬件系统的通信,控制和读取硬件系统的...
  • Code Composer Studio ccs使用方法

    千次阅读 2019-09-01 19:01:18
    Code Composer Studio基本使用方法现已制作完成,如有需要的同学请加qq群:176242843
  • 1.打开CCS5.2,新建一个工程。单击左上角工具栏Project,选择New ccs project.建立一个新的工程。在弹出的任务框project name 内填入工程名称2.在左方框内右击自己的工程名文件,选择最后的properties选项。...
  • 2. 安装CCS3.3, 我用的是:codecomposerstudio3.3ccstudiov3.3白金版ccs3.3.ISO,默认安装,有一个警告,点x无视之。 3. 打补丁:参考了北京交通大学电气工程学院夏明超的课件:《DSP原理与应用2012-第二章_DSP...
  • Code Composer Studio 安装教程与入门

    千次阅读 多人点赞 2020-05-01 18:25:07
    实验一 Code Composer Studio 安装教程与入门 一、实验目的 1.掌握 Code Composer Studio v5 的安装和配置步骤过程。 2.了解 DSP 开发系统和计算机与目标系统的连接方法。 3.了解 Code Composer Studio v5 软件的...
  • 转 [转]JLink重刷固件(win7/win8/win10亲测可用) 2017年09月12日 14:31:55 醒着的树懒 阅读数:6329 ... ...
  • 【美】Texas Instruments Incorprated 著 牛金海 等编译 本书主要介绍CCS 开发环境的使用,非常适合于DSP开发的工程技术人员以及高校的学生参考.
  • MSP-EXP430G2 LaunchPad是德州仪器(TI)提供的一款开发工具,用于学习和练习如何使用他们的微控制器。该开发板属于MSP430产品线系列,我们可以对所有MSP430系列微控制器进行编程。学习如何使用TI微控制器肯定会成为...
  • 启动CodeComposerStudio5 3. 创建工程 选择菜单“View”的“ProjectExplorer”项,打开工程查看界面在ProjectExplorer窗口的空白地方右键选择New->CCSProject: 在弹出的窗口中设置工程名称,以及工程的...
  • Code Composer Studio Licensing Information 1. Help -> Code Composer Studio Licensing Information 2. Upgrade 3. Specify a license file C:\ti\ccsv6\ccs_base\DebugServer\license
  • 本书主要介绍 CCS 开发环境的使用。 适合于从事 TI DSP开发的工程技术人 员以及高校的学生参考。 本书根据 Texas Instrument(TI)的datasheet SPRU509F.pdf 编译而成。 为开始使用Code Composer Studio IDETM,...
  • 2.启动CodeComposerStudio53.导入工程文件 工程文件目录为:C:\ICETEK\F28335A-V5.0\Lab303-Timer 在项目浏览器中,双击timer.c,激活timer.c文件,浏览该文件的内容,理解各语句作用。 3.点击按钮,CCS...
  • --silicon_version=mspx -g -O2 --define=__MSP430F5529__ --define=__MSP430F5529 --include_path="D:/Driver/CodeComposerStudio/ccsv6/ccs_base/msp430/include" --include_path="D:/Driver/CodeComposerStudio/...

空空如也

空空如也

1 2
收藏数 23
精华内容 9
关键字:

codecomposerstudio