OpenCl中的标量核和向量核

Question

我正在读一篇关于使用OpenCl的迷你分子动力学应用的论文http://eprints.dcs.warwick.ac.uk/1694/1/miniMD_opencl.pdf。代码位于此处。

我被内核的实现方式卡住了。我不明白的是

#if defined(SCALAR_KERNELS)
__kernel void f_clear(
    __global float* f,
    __const int nall) {

    for (unsigned i = get_global_id(0)+1; i <= nall; i += get_global_size(0)) {
        const int i4 = i << 2;
        f[i4+0] = 0.0f;
        f[i4+1] = 0.0f;
        f[i4+2] = 0.0f;
        f[i4+3] = 0.0f;
    }

}
#elif defined(VECTOR_KERNELS)
__kernel __attribute__((vec_type_hint(float4)))
void f_clear(
    __global float4* f,
    __const int nall) {

    const float4 zeroes = (float4) (0.0f, 0.0f, 0.0f, 0.0f);
    for (unsigned i = get_global_id(0)+1; i <= nall; i += get_global_size(0)) {
        f[i] = zeroes;
    }

}
#endif

假设VECTOR_KERNELS和SCALAR_KERNELS对应于GPU和麦克风设备，但不确定。

这与MIMD SIMD指令或多核和向量编程有关吗？

另外，现在使用向量类型有真正的优势吗？

最后，我真的不明白这两个for循环是做什么的，以及它们的用途。

为什么不直接使用f[get_global_id(0)]呢？

谢谢,

埃里克。

Answer 1

标量和向量只是在OpenCL中执行相同操作的不同方式。但是矢量是的方式，因为它们应该被编译器( CPU或GPU/FPGA)优化得更好。这样，编译器就可以自然地开发出SIMD单元。因此，如果可能并且对你来说更容易，使用它们。

正如奥斯汀所说，这两个循环都清除了nall的全局内存大小。

然而，查看代码是非常低效的。同一工作组中的工作项正在访问完全不同的全局内存区域，打破了合并。就像你说的那样，这会好得多：

__kernel __attribute__((vec_type_hint(float4)))
void f_clear(
    __global float4* f) {
    f[get_global_id(0)] = (float4) (0.0f, 0.0f, 0.0f, 0.0f);
}

并使用适当的全局大小(global_size = nall)启动此内核，并让编译器决定本地工作组大小。

PS:如果必须这样做，我更喜欢调用clEnqueueWriteBuffer并从CPU中清除内存。因为它可以与其他内核执行并行完成。

Answer 2

一些设备，如AMD、ATI和Intel，在支持向量类型方面做得非常好。这些向量是SIMD，如果可能的话，使用它会更快。NVIDIA在支持OpenCL中的向量方面似乎不是很好(至少是我测试过的所有的)。

这两个循环似乎都清除了一大块大小为nall的全局内存。