1

我正在阅读这篇关于使用 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 和 MIC 设备，但不确定。
这是否与 MIMD SIMD 指令或多核和矢量编程有关？
现在使用矢量类型还有真正的优势吗？
最后，我真的不知道这两个 for 循环做了什么以及它们的目的。
为什么不直接做f[get_global_id(0)]？
谢谢，
埃里克。

4

2 回答 2

3

标量和向量只是在 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 中清理内存。因为它可以与其他内核执行并行完成。

于 2013-10-10T10:21:20.477 回答
1

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

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

于 2013-10-10T04:20:56.167 回答