为什么这个 OpenCL 算法中的本地内存这么慢?
Why is local memory in this OpenCL algorithm so slow?
我正在写一些 OpenCL 代码。我的内核应该根据输入图像创建一个特殊的 "accumulator" 输出。我尝试了两个概念,但都同样慢,尽管第二个使用本地内存。你能帮我确定为什么本地内存版本这么慢吗?内核的目标 GPU 是 AMD Radeon Pro 450。
// version one
__kernel void find_points(__global const unsigned char* input, __global unsigned int* output) {
const unsigned int x = get_global_id(0);
const unsigned int y = get_global_id(1);
int ind;
for(k = SOME_BEGINNING; k <= SOME_END; k++) {
// some pretty wild calculation
// ind is not linear and accesses different areas of the output
ind = ...
if(input[y * WIDTH + x] == 255) {
atomic_inc(&output[ind]);
}
}
}
// variant two
__kernel void find_points(__global const unsigned char* input, __global unsigned int* output) {
const unsigned int x = get_global_id(0);
const unsigned int y = get_global_id(1);
__local int buf[7072];
if(y < 221 && x < 32) {
buf[y * 32 + x] = 0;
}
barrier(CLK_LOCAL_MEM_FENCE);
int ind;
int k;
for(k = SOME_BEGINNING; k <= SOME_END; k++) {
// some pretty wild calculation
// ind is not linear and access different areas of the output
ind = ...
if(input[y * WIDTH + x] == 255) {
atomic_inc(&buf[ind]);
}
}
barrier(CLK_LOCAL_MEM_FENCE);
if(get_local_id(0) == get_local_size(0) - 1)
for(k = 0; k < 7072; k++)
output[k] = buf[k];
}
}
我希望第二个变体比第一个变体更快,但事实并非如此。有时甚至更慢。
本地缓冲区大小__local int buf[7072]
(28288 字节)太大。我不知道 AMD Radeon Pro 450 的共享内存有多大,但可能是每个计算单元 32kB 或 64kB。
32768/28288 = 1
、65536/28288 = 2
意味着只有1个或最多2个波阵面(64个工作项)只能同时运行,所以计算单元的占用率非常非常低,因此性能很差。
您的目标应该是尽可能减少本地缓冲区,以便可以同时处理更多波前。
使用 CodeXL
分析您的内核 - 有一些工具可以向您展示所有这些。
或者,如果您不想 运行 探查器更好地了解它的内容,您可以查看 CUDA occupancy calculator
excel 电子表格。
我正在写一些 OpenCL 代码。我的内核应该根据输入图像创建一个特殊的 "accumulator" 输出。我尝试了两个概念,但都同样慢,尽管第二个使用本地内存。你能帮我确定为什么本地内存版本这么慢吗?内核的目标 GPU 是 AMD Radeon Pro 450。
// version one
__kernel void find_points(__global const unsigned char* input, __global unsigned int* output) {
const unsigned int x = get_global_id(0);
const unsigned int y = get_global_id(1);
int ind;
for(k = SOME_BEGINNING; k <= SOME_END; k++) {
// some pretty wild calculation
// ind is not linear and accesses different areas of the output
ind = ...
if(input[y * WIDTH + x] == 255) {
atomic_inc(&output[ind]);
}
}
}
// variant two
__kernel void find_points(__global const unsigned char* input, __global unsigned int* output) {
const unsigned int x = get_global_id(0);
const unsigned int y = get_global_id(1);
__local int buf[7072];
if(y < 221 && x < 32) {
buf[y * 32 + x] = 0;
}
barrier(CLK_LOCAL_MEM_FENCE);
int ind;
int k;
for(k = SOME_BEGINNING; k <= SOME_END; k++) {
// some pretty wild calculation
// ind is not linear and access different areas of the output
ind = ...
if(input[y * WIDTH + x] == 255) {
atomic_inc(&buf[ind]);
}
}
barrier(CLK_LOCAL_MEM_FENCE);
if(get_local_id(0) == get_local_size(0) - 1)
for(k = 0; k < 7072; k++)
output[k] = buf[k];
}
}
我希望第二个变体比第一个变体更快,但事实并非如此。有时甚至更慢。
本地缓冲区大小__local int buf[7072]
(28288 字节)太大。我不知道 AMD Radeon Pro 450 的共享内存有多大,但可能是每个计算单元 32kB 或 64kB。
32768/28288 = 1
、65536/28288 = 2
意味着只有1个或最多2个波阵面(64个工作项)只能同时运行,所以计算单元的占用率非常非常低,因此性能很差。
您的目标应该是尽可能减少本地缓冲区,以便可以同时处理更多波前。
使用 CodeXL
分析您的内核 - 有一些工具可以向您展示所有这些。
或者,如果您不想 运行 探查器更好地了解它的内容,您可以查看 CUDA occupancy calculator
excel 电子表格。