如何进一步优化这个 OpenCL 内核

Question

我正在尝试在 OpenCL 中实现 hausdorff 距离，以下内核构成了它的基础，或者我认为它确实如此，因为我仍然必须完全实现它。也就是说，我可以得到一些建议还是有办法优化这个内核？基本上我怎样才能删除调用辅助函数的内核函数中的for循环.....

OpenCL Kernel 及其辅助函数：

void helper( int a_1, __global int* b_1, __global int* c_1 ){
        int i = get_global_id(0);
        c_1[i] = a_1 - b_1[i];
}

__kernel void test_call( __global int* a,             //input buffer of size [100000, 1]
                         __global int* b,             //input buffer of size [100000, 1]
                         __global int* c ){           //output buffer of size [100000, 1]
        for ( int iter = 0 ; iter < 100000 ; iter++ ){
                helper ( a[iter], b, c );
                // once array c is obtained by calling the above function,
                // it will be used in further processing that will take place inside 
                // this for loop itself
}

基本上我在这里尝试做的是用输入缓冲区“b”中的每个元素减去输入缓冲区“a”中的每个元素。复杂度为 O(n²)。

顺便说一句，这个简单的实现本身会在 2.5 秒内产生结果。对此的串行实现将需要几分钟才能完成执行。

Answer 1

我想，您的代码仍然可以通过使用平铺方案来改进。 在您当前的实现中，所有工作项都加载缓冲区“a”的所有值。目前，他们以不同步的方式这样做。使用平铺方案，您可能可以更好地利用您的缓存架构，方法是让设备只从片外内存中加载“a”的每个值。

这个方案最好在这里解释：http://software.intel.com/file/24571（PDF 文件）。

在您的情况下，Tiled Parallel 方法可能类似于以下伪代码。

forall values of b in parallel {
    foreach tile q {
        forall work_items p in work_group in parallel {
            local[p] = a[q*tile_size + p]
        }
        synchronize work items in work group
        foreach body j in tile q {
            c[i] = local[j] - b[i];
        }
        synchronize work items in work group
    }
}

关键思想是每个工作项都可以使用缓冲区“a”的值，这些值已经被同一组中的其他工作项加载。 'a' 的每个条目将（理想情况下）仅从内存中获取一次，从缓存中获取 tile_size-1 次。