2

基本上，我正在编写一个 OpenCL 内核，它以随机/不可预测的方式访问全局内存（光线跟踪器的当前未优化的路径跟踪组件），这几乎完全否定了 GPU 相对于 CPU 的并行化性能优势（供参考，我在 i7-2630QM CPU、GTX 560m GPU 上运行——性能数据如下）。为了便于调整/测试，我编写了一个“测试”内核来模拟这种内存访问模式；它本质上为 GPU 提供了大量的三角形坐标和要处理的索引列表 - 对于每个索引，它将在该三角形和之后的 63 上运行射线三角形交集，模拟八叉树中对象的迭代。

我尝试了一系列优化，包括：合并内存访问、使用只读纹理内存代替“全局”、循环展开、调整工作组大小和线程分布、本地内存和屏障以及手动内联函数。这些都提供了，充其量，增量性能改进。在运行内核之前对索引进行排序确实可以显着加快速度，但是对于八叉树遍历，这将需要在 GPU 上重新排序每次迭代，并结合其他因素，让我怀疑它是否会有很大帮助。

我试图弄清楚是否有一些可以修复的主要漏洞——滥用数据类型、看不见的优化、驱动程序太旧（使用 OpenCL 1.0，它不允许 1d 纹理）等等——或者如果我期待考虑到我正在使用的硬件，性能提升太多了（关于光线追踪方面的各种优化仍有待完成，但我想在深入研究之前解决这个更普遍的问题）。非常感谢您提前提出任何见解或建议。

409,600 个 64 个三角形的块（以 409,600 个线程运行）的性能数据（秒）：

CPU (Single Thread):

Unsorted:   2.21
Sorted:     1.48 

GPU:

        Sorted  Unsorted
Texture     0.07    0.15
Global      0.02    0.25

代码：

#define IMG_WIDTH_MINUS_ONE 32767
#define IMG_HEIGHT_LOG_2 15
#define SUB(dest,v1,v2) \
      dest[0]=v1[0]-v2[0]; \
      dest[1]=v1[1]-v2[1]; \
      dest[2]=v1[2]-v2[2];

#define EPSILON 0.00001

#define CROSS(dest,v1,v2) \
      dest[0]=v1[1]*v2[2]-v1[2]*v2[1]; \
      dest[1]=v1[2]*v2[0]-v1[0]*v2[2]; \
      dest[2]=v1[0]*v2[1]-v1[1]*v2[0];

#define DOT(v1,v2) (v1[0]*v2[0]+v1[1]*v2[1]+v1[2]*v2[2])

__kernel void square(
   __global int4 *inputIndeces,
   __read_only image2d_t image,
   __global float* output,
   const unsigned int count)
{

    int global_id = get_global_id(0);
    float r_orig[3];
    float r_dir[3];
    float4 trianglePoints[3];
    int cpuStartIndex = inputIndeces[global_id].x;
    int outputIndex = inputIndeces[global_id].w;
    output[outputIndex] = 0.0;
    r_orig[0] = 0.0;
    r_orig[1] = 0.0;
    r_orig[2] = 500.0;
    r_dir[0] = 0.0;
    r_dir[1] = 0.0;
    local int counter;
    counter = 0;
    r_dir[2]= -1.0;
    float tvec[3], pvec[3], qvec[3], edgeA[3], edgeB[3];
    float det, inv_det, t, u, v;
    #pragma unroll 64
    for (int ind=cpuStartIndex;ind<cpuStartIndex+64;++ind) {



        int tIndex = ind<<2;

        int2 coords[3];


        coords[0] = (int2)(tIndex & IMG_WIDTH_MINUS_ONE,tIndex >> IMG_HEIGHT_LOG_2);
        coords[1] = (int2)((tIndex + 1) & IMG_WIDTH_MINUS_ONE,(tIndex + 1) >> IMG_HEIGHT_LOG_2);
        coords[2] = (int2)((tIndex + 2) & IMG_WIDTH_MINUS_ONE,(tIndex + 2) >> IMG_HEIGHT_LOG_2);

        trianglePoints[0] = read_imagef(image, sampler, coords[0]);
        trianglePoints[1] = read_imagef(image, sampler, coords[1]);
        trianglePoints[2] = read_imagef(image, sampler, coords[2]);

        edgeA[0] = (trianglePoints[0].w - trianglePoints[0].x);
        edgeA[1] = (trianglePoints[1].x - trianglePoints[0].y);
        edgeA[2] = (trianglePoints[1].y - trianglePoints[0].z);

        edgeB[0] = (trianglePoints[1].z - trianglePoints[0].x);
        edgeB[1] = (trianglePoints[1].w - trianglePoints[0].y);
        edgeB[2] = (trianglePoints[2].x - trianglePoints[0].z);

        CROSS(pvec,r_dir,edgeB);
        det = DOT(edgeA, pvec);

        if (det > -EPSILON && det < EPSILON) {
             continue;
        }

        inv_det = 1.0 / det;

        tvec[0] = r_orig[0] - trianglePoints[0].x;
        tvec[1] = r_orig[1] - trianglePoints[0].y;
        tvec[2] = r_orig[2] - trianglePoints[0].z;

        u = DOT(tvec, pvec) * inv_det;
        if (u < 0.0 || u > 1.0) {
                continue;
            }

            CROSS(qvec,tvec,edgeA);
            v = DOT(r_dir, qvec) * inv_det;
            if (v < 0.0 || u + v > 1.0) {
                continue;
            }
            t = DOT(edgeB, qvec) * inv_det;
            if (t > 0.001) {
                ++counter;
            }
        else {
            continue;
        }

    }
    output[outputIndex] = (float)counter;
}
4

0 回答 0