私有内存索引 OpenCL

Question

我正在为 OpenCL specification 苦苦挣扎，因为我发现它有时模棱两可，有人可以尝试回答以下问题吗？

考虑以下代码：

__kernel void myKernel(...)
{    
    // Buffer 1
    __local float *buffer1[64];

    // Buffer 2
    __local float *buffer2;

    // Buffer 3
    __private float *buffer3[64];

    // Buffer 4
    float *buffer4[64];

    int var1 = 1, var2 = 2;
    nonKernelFunction(&var1, &var2);

    // ...

}

void nonKernelFunction(int *pvar1, int *pvar2)
{
    int *pvar;
    if (someRunTimeCondition)
        pvar = pvar1;
    else
        pvar = pvar2;
    *pvar += 1; 
}

1) buffer1 和 buffer2 之间是否有区别（静态或动态）？

2) buffer3 和 buffer4 的声明是否等效（它们是用于变量，但我不确定是否用于指针）？

3) 在 GPU 上（我认为私有内存只是寄存器），编译器将在哪里分配资源？如果它在全局内存中，是否可以从主机知道运行时将使用多少内存？

4) 假设 buffer3 和 buffer4 存储在寄存器中，如何允许像 buffer3[i] = buffer4[i] （其中 i 在运行时已知）这样的指令？

5) 如果buffer3 和buffer4 没有存入寄存器，那么如何允许nonKernelFunction 代码（var1 和var2 肯定不在内存中）？

谢谢

Answer 1

阿法伊克：

1) 内核代码中的静态规范与主机通过缓冲区的“动态”规范之间没有技术差异；

2) 默认情况下变量是 __private，所以这不应该有任何区别；

3) 如果私有内存很小，则可以在寄存器中分配私有内存，否则将使用全局内存； 您可以使用 clGetKernelWorkGroupInfo 查询内核的最低内存要求；

4) 为什么不允许它们，因为它可能导致越界错误？

5) var1 和 var2 在 GPU 的地址空间中，即使不在私有内存中；访问可能会更慢，仅此而已。

编辑1： var1 和 var2 在寄存器中的事实，比如 reg1 和 reg2，应该不是问题，因为代码可能会导致伪汇编，例如：

myKernel:
    ...
    push reg1
    push reg2
    call nonKernelFunction
    ...

nonKernelFunction:
    test someRunTimeCondition
    jz ko
        mov [SP+2] reg1
        jmp end:
    ko:
        mov [SP+1] reg1
    end:
    mov [reg1] reg2
    inc reg2
    mov reg2 [reg1]

我不知道 GPU 组件/核心架构是否有很大不同，但在标准 CPU 上没有问题，因为您使用堆栈来抽象有效位置。

请注意，这里有更新版本的规范 :) http://www.khronos.org/registry/cl/specs/opencl-1.2.pdf