在OpenCL中写入全局内存

我试图优化一些最初在Fortran中编写的代码。在OpenCL中写入全局内存

算法涉及在多个迭代上对大阵列（~27百万个细胞）进行操作。每个细胞可以在一次迭代过程中独立评估。然而，由于在t + 1完成的计算取决于在t完成的计算结果，所以迭代不能并行化。

粗糙，简化不平行的伪代码示例：

for (t=0; t<tmax; t++) 
{ 
A = A + B; 

B = B + A /2; 
}

其中A和B是大阵列。

目前，我已经通过调用EnqueueNDRangeKernel在宿主C++代码中的一个循环中实现了这一点。因为我需要前一次迭代的结果，所以我每次都写入全局内存。

每次迭代执行2700万次全局内存写入会杀死我的性能。我有两个我正在玩的内核版本;与Fortran相比，版本1的速度快了2.5倍;版本2快4倍。

我已经尝试摆弄算法，以及做指针（版本2）的东西。

我的问题如下：有没有办法避免这个全局内存写chokepoint？

谢谢！

请求代码：

呼叫在C++：

NDRange global(nxp1*ny*nz); 
NDRange local(nz); 

    for (w=0; w<100; w++) 
    { 
     queue.enqueueNDRangeKernel(kernA, NullRange, global, local); 
     queue.enqueueBarrierWithWaitList(); 
     queue.enqueueNDRangeKernel(kernB, NullRange, global, local); 
    } 

queue.finish();

仁：

__kernel void kernA(__global double *A, __global double *B) 
    { 
    int i = get_global_id(0); 

    double A_l; 
    A_l = A[i]; 
    double B_l; 
    B_l = B[i]; 

    A_l = A_l + B_l; 

    A[i] = A_l; //if this line is removed, everything goes much faster. 

    } 

    __kernel void kernB(__global double *A, __global double *B) 
    { 
    int i = get_global_id(0); 

    double A_l; 
    A_l = A[i]; 
    double B_l; 
    B_l = B[i]; 

    B_l = B_l + A_l/2; 

    B[i] = B_l; //if this line is removed, everything goes much faster. 
    }

我已经简化内核代码澄清算法的缘故。但是这个想法是我基于B更新A;然后我基于A更新B，重复这个过程进行几次迭代。

来源

2014-09-26 Vlad

为了计算新的A或B的“（i，j）”条目，是否需要引用除了旧的A和B的“（i，j）”条目之外的任何东西？如果不是，那么除了第一次迭代的输入和最后一次迭代的输出之外，您不需要使用全局内存，并且可以在单个内核调用中计算多次迭代。 – user57368 2014-09-26 18:28:54

我只需要知道前一次迭代中的旧A（i）和B（i）。你是否打算在内核中运行迭代循环？这会保持迭代的顺序吗？ – Vlad 2014-09-26 18:39:37

细胞是由什么组成的？单一的价值？ int，float，double或其他东西？你可以在同一个内核调用中计算多于一次的迭代。 – mfa 2014-09-26 18:43:07

没有办法完全避免全局写入问题。你只写了一次数值，而且你的速度是硬件限制的。尽管您可以一次计算多个步骤，但您可以减少全局读取次数。这仍然节省了每一步。

__kernel void myKernel(__global double *A, __global double *B, __global uint outDataMultiple)    
{                      
    const uint gid = get_global_id(0); 
    const uint inDataSize = get_global_size(0); 

    double2 nextValue; 
    nextValue.x = A[gid]; 
    nextValue.y = B[gid]; 
    for(uint i=0; i<outDataMultiple; i++){ 
     nextValue.x = nextValue.x + nextValue.y; 
     nextValue.y = nextValue.y + nextValue.x /2; 
     A[gid+i+1] = nextValue.x; 
     B[gid+i+1] = nextValue.y; 
    } 
}

在上面的内核中，一个工作项目将处理单个单元的多次迭代。你需要分配outDataMultiple倍多的内存，而内核将填充剩下的内存。全局工作项目数决定了初始输入的大小。 outDataMultiple仅受全局内存分配的限制，并且可能会影响每次迭代的数学复杂性。

全球总存储器需求： 27M *的sizeof（double2）*（1 + outDataMultiple）

__kernel void myKernel(__global double2 *data, __global uint outDataMultiple)    
{                      
    const uint gid = get_global_id(0); 
    const uint inDataSize = get_global_size(0); 

    double2 nextValue = data[gid]; 
    for(uint i=0; i<outDataMultiple; i++){ 
     nextValue.x = nextValue.x + nextValue.y; 
     nextValue.y = nextValue.y + nextValue.x /2; 
     data[gid+i+1] = nextValue; 
    } 
}

相同的内核的double2版本可以是可能的，只要可以交错的A和B向量。这将结合读取和写入来保证8字节块，并且可能会提高性能。

来源

2014-09-26 19:20:30 mfa

严格地说，从概念上讲，步骤如何1.从全局内存中加载内容2.启动循环3.在private上运行A = f（B），B = f（A）4.结束循环中5.复制回全局？我认为这些失败是因为寄存器中没有足够的可用内存，并且因为opencl产生了一堆线程，所以循环会独立运行在每个单元格上（这是因为[i]依赖于例如b [i]但是b [i + 1]）。 – Vlad 2014-09-26 19:32:52

27M线程不会全部同时启动。驱动程序/硬件将为您管理所有这些。您仍然应该使用enqueueNDrange指定工作组大小。在计算示例代码中的A [i + 1]时，我看不到B [i + 1]在哪里。我列出的代码将重现原始循环，但是在不同的内存位置，因此每次outDataMultiple写入只能进行一次全局读取。 – mfa 2014-09-26 19:40:37

减少OpenCL设备从全局内存中获取花费时间的简单方法是将全局内存批量缓存到本地内存，在本地内存上运行，然后将批处理本地内存写入全局内存。

本地内存的延迟与线程内存基本相同，可以从全局内存中以块的形式读取。本地内存可以在主机上声明并传递给内核（请参见下面的示例）或在内核中分配并使用（请参阅下面列出的AMD优化指南中的示例）。例如：

__kernel void kernA(__global double *A, 
        __global double *B, 
        __local double *BufferA, 
        __local double *BufferB) 
    { 
    BufferA[get_local_id(0)] = A[get_global_id(0)]; 
    BufferB[get_local_id(0)] = B[get_global_id(0)]; 
    mem_fence(CLK_LOCAL_MEM_FENCE); 

    double tmp = BufferA[get_local_id(0)] + BufferB[get_local_id(0)]; 

    A[get_global_id(0)] = BufferA[get_local_id(0)]; 
    mem_fence(CLK_GLOBAL_MEM_FENCE); 
    }

有进一步的东西可以进行，包括：

存储器平铺 - 使用2D全局ID
优化本地存储器大小为设备
存取存储器如果可以的话使用浮点数
尽可能多地进入内核
Read this guide by AMD for optmizing OpenCL kernels

来源

2014-09-26 22:17:55 Austin

在OpenCL中写入全局内存

回答

相关问题