关于我的OpenCL内核（DSP窗口函数）分析结果的困惑

我在OpenCL中完成了Window Function内核。基本上，窗口函数只是将一组系数应用于另一组数字（维基百科更好地解释它）。在大多数情况下，我能够将窗口系数浮点数组填充到常量缓存中。关于我的OpenCL内核（DSP窗口函数）分析结果的困惑

我预计Compute Prof的结果表明，主机到设备和设备到主机的内存传输将花费超过95％的处理时间。几乎所有的情况下，只有80％的处理时间。我正在写和读一个420万浮点阵列，并写出另一个浮点数组，通常保持在一百万以下。

在内核中是否有任何东西看起来很腥？任何意见，如果它是一个问题，应该在GPU上运行得比CPU更快（我仍然不是100％）。我有点震惊，为什么我的gld_efficiency和gst_efficiency在0.1和0.2之间徘徊。我在考虑G80全局内存的情况下制作了这个内核。我的全局内存整体吞吐量似乎是40gbs。内核非常简单，并在下面发布。

__kernel void window(__global float* inputArray, // first frame to ingest starts at 0. Sized to nFramesToIngest*framesize samples 
    __constant float* windowArray, // may already be partly filled 
    int windowSize, // size of window frame, in floats 
    int primitivesPerDataFrame, //amount of primitives in each frame of inputArray parameter 
    int nInFramesThisCall, //each thread solves a frame, so this integer represent how many threads this kernel launches 
    int isRealNumbers //0 for complex, non-zero for real 
) 
{ 
int gid = get_global_id(0) + get_global_size(0) * get_global_id(1); 

if(gid < nInFramesThisCall) //make sure we don't execute unnecessary threads 
{ 
    if(isRealNumbers) 
    { 
     for(int i = 0; i < primitivesPerDataFrame; i++) 
     { 
      int inputArrayIndex = (gid*primitivesPerDataFrame)+i; 
      inputArray[inputArrayIndex] = inputArray[inputArrayIndex] * windowArray[i%windowSize]; 
     } 
    } 
    else //complex 
    { 
     for(int i = 0; i < primitivesPerDataFrame; i++) 
     { 
      int inputArrayIndex = (gid*primitivesPerDataFrame)+i; 
      inputArray[inputArrayIndex] = inputArray[inputArrayIndex] * windowArray[i%windowSize/2]; 
     } 
    } 
}

}

来源

2010-11-09 smuggledPancakes

多少个线程（OpenCL的术语工作项，顺便说一句）您使用的？您至少需要几百个内容来高效加载大型GPU。

你说你想使用的聚结的内存访问，但有一个负载偏移像

int inputArrayIndex = (gid*primitivesPerDataFrame)+i;

不会在大多数情况下，使之成为可能。 NVidia的G80在合并时有非常严格的限制，请参阅“OpenCL最佳实践指南”以获取更多信息。基本上，来自一个warp的工作项必须以某种方式同时访问64或128个字节对齐块的元素，以使加载和存储发生合并。

或者举一个例子：如果primitivesPerDataFrame为16，那么经纱的加载和存储是在间隔16个元件的偏移处完成的，这使得不可能有效地进行合并。

来源

2010-11-11 16:58:05 dietr

我通常每块启动256个线程。我相信我的条件语句，如果isRealNumbers == true，那么我的warp在InputArray中访问一个128字节对齐的块。如果isRealNumbers == false，那么我的warp会在InputArray中访问两个128字节对齐的块？我觉得读取和写入请求的总数应该是1：1，包括读取和写入事务的总数。如果我没有看到什么东西，请分享，因为分析器不同意我:) – smuggledPancakes 2010-11-12 14:08:03

只需计算'primitivesPerDataFrame'> 1的一些示例的偏移量，您就会看到访问不是连续的，而是有一定的步伐迅速变得足够大，使不可能的联合。此外，您正在为真实和复杂数字执行完全相同的全局内存访问。 – dietr 2010-11-12 20:24:57

对不起，如果我是密集的，但我不明白我对inputArray的任何访问是不是连续的。你能解释一下吗？ – smuggledPancakes 2010-11-15 14:19:32

关于我的OpenCL内核（DSP窗口函数）分析结果的困惑

回答

相关问题