Viết vào bộ nhớ toàn cầu và địa phương tăng thời gian thực hiện kernel bằng 10000%

Question

Tôi có hạt nhân OpenCL sau:

kernel void ndft(
    global float *re, global float *im, int num_values, 
    global float *spectrum_re, global float *spectrum_im, 
    global float *spectrum_abs, 
    global float *sin_array, global float *cos_array, 
    float sqrt_num_values_reciprocal) 
{ 
    // MATH MAGIC - DISREGARD FROM HERE ----------- 

    float x; 
    float y; 
    float sum_re = 0; 
    float sum_im = 0; 

    size_t thread_id = get_global_id(0); 
    //size_t local_id = get_local_id(0); 

    // num_values = 24 (live environment), 48 (test) 
    for (int i = 0; i < num_values; i++) 
    { 
     x = cos_array[thread_id * num_values + i] * sqrt_num_values_reciprocal; 
     y = sin_array[thread_id * num_values + i] * sqrt_num_values_reciprocal; 
     sum_re = sum_re + re[i] * x + im[i] * y; 
     sum_im = sum_im - re[i] * y + x * im[i]; 
    } 

    // MATH MAGIC DONE ---------------------------- 

    //spectrum_re[thread_id] = sum_re; 
    //spectrum_im[thread_id] = sum_im; 
    //spectrum_abs[thread_id] = hypot(sum_re, sum_im); 
    float asdf = hypot(sum_re, sum_im); // this is just a dummy calculation 
}

Như thế này, thời gian thực hiện là khoảng 15 chúng tôi (kích thước nhóm làm việc = 567, 14 nhóm làm việc , cho tổng số 7938 chủ đề).

Tuy nhiên, bằng cách nào đó, tôi bằng cách nào đó cần truy xuất kết quả của hoạt động, đó là những gì mà một vài dòng cuối cùng dành cho, (nhận xét). Ngay sau khi tôi thực hiện một hoạt động bộ nhớ nào đó (và không quan trọng nếu spectrum_X là global, như trong ví dụ hoặc local), thời gian exeuction của hạt nhân tăng lên ~ 1,4 đến 1,5 ms.

Tôi nghĩ thời gian thực hiện tăng lên là một số chi phí cố định, vì vậy tôi sẽ tích luỹ thêm dữ liệu để số lượng thời gian tương đối bị mất do hiệu ứng đó giảm thiểu. Nhưng khi tôi tăng gấp đôi số lượng chủ đề của mình (gấp đôi số lượng dữ liệu), thời gian thực hiện cũng tăng gấp đôi (lên 2.8 ~ 3.0 ms).

Tôi phát hiện ra rằng ngay cả khi tôi chỉ bỏ ghi chú một của các dòng đó, tôi có cùng thời gian thực hiện như thể tôi đã bỏ ghi chú cả ba. Ngay cả khi tôi thêm if (thread_id == 0) và chạy nó, tôi cũng có cùng thời gian thực hiện. Tuy nhiên, nó chỉ là cách quá chậm theo cách này (giới hạn trên cho ứng dụng của tôi là khoảng 30 chúng tôi). Nó thậm chí còn nhanh hơn gấp 5 lần khi tôi chạy nó trong mã C thông thường trên CPU của tôi.

Bây giờ tôi rõ ràng đang làm điều gì đó sai, nhưng tôi không chắc chắn nơi để bắt đầu tìm kiếm một giải pháp.

---

Như tôi đã nhận xét về câu trả lời talonmies', tôi cũng đã làm như sau:

Từ đoạn code trên, tôi đã 4 dòng cuối cùng trông giống như

//spectrum_re[thread_id] = sum_re; 
//spectrum_im[thread_id] = sum_im; 
spectrum_abs[thread_id] = hypot(sum_re, sum_im); 
//float asdf = hypot(sum_re, sum_im);

Theo dự kiến, thời gian thực hiện ~ 1,8 ms. Mã lắp ráp tạo ra cho hệ thống của tôi là:

// 
// Generated by NVIDIA NVVM Compiler 
// Compiler built on Tue Apr 03 12:42:39 2012 (1333449759) 
// Driver 
// 

.version 3.0 
.target sm_21, texmode_independent 
.address_size 32 


.entry ndft(
    .param .u32 .ptr .global .align 4 ndft_param_0, 
    .param .u32 .ptr .global .align 4 ndft_param_1, 
    .param .u32 ndft_param_2, 
    .param .u32 .ptr .global .align 4 ndft_param_3, 
    .param .u32 .ptr .global .align 4 ndft_param_4, 
    .param .u32 .ptr .global .align 4 ndft_param_5, 
    .param .u32 .ptr .global .align 4 ndft_param_6, 
    .param .u32 .ptr .global .align 4 ndft_param_7, 
    .param .f32 ndft_param_8 
) 
{ 
    .reg .f32 %f; 
    .reg .pred %p; 
    .reg .s32 %r; 


    ld.param.u32 %r3, [ndft_param_2]; 
    // inline asm 
    mov.u32  %r18, %envreg3; 
    // inline asm 
    // inline asm 
    mov.u32  %r19, %ntid.x; 
    // inline asm 
    // inline asm 
    mov.u32  %r20, %ctaid.x; 
    // inline asm 
    // inline asm 
    mov.u32  %r21, %tid.x; 
    // inline asm 
    add.s32  %r22, %r21, %r18; 
    mad.lo.s32 %r11, %r20, %r19, %r22; 
    setp.gt.s32  %p1, %r3, 0; 
    @%p1 bra BB0_2; 

    mov.f32  %f46, 0f00000000; 
    mov.f32  %f45, %f46; 
    bra.uni  BB0_4; 

BB0_2: 
    ld.param.u32 %r38, [ndft_param_2]; 
    mul.lo.s32 %r27, %r38, %r11; 
    shl.b32  %r28, %r27, 2; 
    ld.param.u32 %r40, [ndft_param_6]; 
    add.s32  %r12, %r40, %r28; 
    ld.param.u32 %r41, [ndft_param_7]; 
    add.s32  %r13, %r41, %r28; 
    mov.f32  %f46, 0f00000000; 
    mov.f32  %f45, %f46; 
    mov.u32  %r43, 0; 
    mov.u32  %r42, %r43; 

BB0_3: 
    add.s32  %r29, %r13, %r42; 
    ld.global.f32 %f18, [%r29]; 
    ld.param.f32 %f44, [ndft_param_8]; 
    mul.f32  %f19, %f18, %f44; 
    add.s32  %r30, %r12, %r42; 
    ld.global.f32 %f20, [%r30]; 
    mul.f32  %f21, %f20, %f44; 
    ld.param.u32 %r35, [ndft_param_0]; 
    add.s32  %r31, %r35, %r42; 
    ld.global.f32 %f22, [%r31]; 
    fma.rn.f32 %f23, %f22, %f19, %f46; 
    ld.param.u32 %r36, [ndft_param_1]; 
    add.s32  %r32, %r36, %r42; 
    ld.global.f32 %f24, [%r32]; 
    fma.rn.f32 %f46, %f24, %f21, %f23; 
    neg.f32  %f25, %f22; 
    fma.rn.f32 %f26, %f25, %f21, %f45; 
    fma.rn.f32 %f45, %f24, %f19, %f26; 
    add.s32  %r42, %r42, 4; 
    add.s32  %r43, %r43, 1; 
    ld.param.u32 %r37, [ndft_param_2]; 
    setp.lt.s32  %p2, %r43, %r37; 
    @%p2 bra BB0_3; 

BB0_4: 
    // inline asm 
    abs.f32  %f27, %f46; 
    // inline asm 
    // inline asm 
    abs.f32  %f29, %f45; 
    // inline asm 
    setp.gt.f32  %p3, %f27, %f29; 
    selp.f32 %f8, %f29, %f27, %p3; 
    selp.f32 %f32, %f27, %f29, %p3; 
    // inline asm 
    abs.f32  %f31, %f32; 
    // inline asm 
    setp.gt.f32  %p4, %f31, 0f7E800000; 
    mov.f32  %f47, %f32; 
    @%p4 bra BB0_6; 

    mov.f32  %f48, %f8; 
    bra.uni  BB0_7; 

BB0_6: 
    mov.f32  %f33, 0f3E800000; 
    mul.rn.f32 %f10, %f8, %f33; 
    mul.rn.f32 %f47, %f32, %f33; 
    mov.f32  %f48, %f10; 

BB0_7: 
    mov.f32  %f13, %f48; 
    // inline asm 
    div.approx.f32 %f34, %f13, %f47; 
    // inline asm 
    mul.rn.f32 %f39, %f34, %f34; 
    add.f32  %f38, %f39, 0f3F800000; 
    // inline asm 
    sqrt.approx.f32  %f37, %f38;  // <-- this is part of hypot() 
    // inline asm 
    mul.rn.f32 %f40, %f32, %f37; 
    add.f32  %f41, %f32, %f8; 
    setp.eq.f32  %p5, %f32, 0f00000000; 
    selp.f32 %f42, %f41, %f40, %p5; 
    setp.eq.f32  %p6, %f32, 0f7F800000; 
    setp.eq.f32  %p7, %f8, 0f7F800000; 
    or.pred  %p8, %p6, %p7; 
    selp.f32 %f43, 0f7F800000, %f42, %p8; 
    shl.b32  %r33, %r11, 2; 
    ld.param.u32 %r39, [ndft_param_5]; 
    add.s32  %r34, %r39, %r33; 
    st.global.f32 [%r34], %f43; // <-- stores the hypot's result in spectrum_abs 
    ret; 
} 

Trên thực tế tất cả các hoạt động tính toán của tôi ở đó - rất nhiều thêm/mults cũng như một sqrt cho hypot chức năng. Từ mã asm trên, tôi loại bỏ dòng cuối cùng thứ hai:

st.global.f32 [%r34], %f43;

đó là dòng mà thực sự lưu trữ các dữ liệu trong mảng toàn cầu spectrum_abs. Sau đó, tôi đã sử dụng clCreateProgramWithBinary và đã sử dụng tệp mã asm đã sửa đổi làm đầu vào. Thời gian thực hiện đã giảm xuống còn 20 chúng tôi.

Answer 1

Tôi đoán bạn sẽ thấy các hiệu ứng của tối ưu hóa trình biên dịch.

Trình biên dịch NVIDIA rất tích cực trong việc loại bỏ "mã chết" không trực tiếp tham gia ghi vào bộ nhớ chung. Vì vậy, trong hạt nhân của bạn, nếu bạn không viết sum_re hoặc sum_im, trình biên dịch sẽ tối ưu hóa toàn bộ vòng lặp tính toán (và có thể là mọi thứ khác) và để lại của bạn với một hạt nhân trống không chứa gì hơn một no-op. Thời gian thực hiện 15 micro giây mà bạn đang thấy hầu hết chỉ là khởi chạy hạt nhân và không còn nhiều thứ khác. Khi bạn bỏ ghi nhớ toàn cục, thì trình biên dịch để lại tất cả mã tính toán tại chỗ và bạn thấy thời gian thực thi thực sự của mã của bạn.Vì vậy, câu hỏi thực sự bạn có lẽ nên hỏi là làm thế nào để tối ưu hóa hạt nhân đó để giảm thời gian thực hiện của nó từ 1,5 mili giây mà nó hiện đang hướng tới mục tiêu 30 micro giây (rất đầy tham vọng) của bạn.

---

Mặc dù hoài nghi bày tỏ với câu trả lời ban đầu, đây là một trường hợp repro hoàn chỉnh, hỗ trợ sự khẳng định rằng đây là một trình biên dịch ảnh hưởng liên quan:

#include <iostream> 
#include <OpenCL/opencl.h> 

size_t source_size; 
const char * source_str = 
"kernel void ndft(                 \n" \ 
" global float *re, global float *im, int num_values,        \n" \ 
" global float *spectrum_re, global float *spectrum_im,        \n" \ 
" global float *spectrum_abs,              \n" \ 
" global float *sin_array, global float *cos_array,         \n" \ 
" float sqrt_num_values_reciprocal)             \n" \ 
"{                      \n" \ 
" // MATH MAGIC - DISREGARD FROM HERE -----------         \n" \ 
"                      \n" \ 
" float x;                   \n" \ 
" float y;                   \n" \ 
" float sum_re = 0;                 \n" \ 
" float sum_im = 0;                 \n" \ 
"                      \n" \ 
" size_t thread_id = get_global_id(0);            \n" \ 
"                      \n" \ 
" for (int i = 0; i < num_values; i++)            \n" \ 
" {                     \n" \ 
"  x = cos_array[thread_id * num_values + i] * sqrt_num_values_reciprocal;  \n" \ 
"  y = sin_array[thread_id * num_values + i] * sqrt_num_values_reciprocal;  \n" \ 
"  sum_re += re[i] * x + im[i] * y;            \n" \ 
"  sum_im -= re[i] * y + x * im[i];            \n" \ 
" }                     \n" \ 
"                      \n" \ 
" // MATH MAGIC DONE ----------------------------         \n" \ 
"                      \n" \ 
" //spectrum_re[thread_id] = sum_re;            \n" \ 
" //spectrum_im[thread_id] = sum_im;            \n" \ 
" //spectrum_abs[thread_id] = hypot(sum_re, sum_im);        \n" \ 
"}                      \n"; 

int main(void) 
{ 
    int err; 

    cl_device_id device_id; 
    clGetDeviceIDs(NULL, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL); 
    cl_context context = clCreateContext(0, 1, &device_id, NULL, NULL, &err); 
    cl_program program = clCreateProgramWithSource(context, 1, (const char **)&source_str, (const size_t *)&source_size, &err); 

    err = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL); 

    cl_uint program_num_devices; 
    clGetProgramInfo(program, CL_PROGRAM_NUM_DEVICES, sizeof(cl_uint), &program_num_devices, NULL); 

    size_t * binaries_sizes = new size_t[program_num_devices]; 
    clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES, program_num_devices*sizeof(size_t), binaries_sizes, NULL); 

    char **binaries = new char*[program_num_devices]; 
    for (size_t i = 0; i < program_num_devices; i++) 
     binaries[i] = new char[binaries_sizes[i]+1]; 

    clGetProgramInfo(program, CL_PROGRAM_BINARIES, program_num_devices*sizeof(size_t), binaries, NULL); 
    for (size_t i = 0; i < program_num_devices; i++) 
    { 
     binaries[i][binaries_sizes[i]] = '\0'; 
     std::cout << "Program " << i << ":" << std::endl; 
     std::cout << binaries[i]; 
    } 
    return 0; 
} 

Khi biên soạn và chạy, nó phát ra tiếp theo Mã PTX từ thời gian chạy OpenCL:

Program 0: 
bplist00?^clBinaryDriver\clBinaryData_clBinaryVersionWCLH 1.0O!.version 1.5 
.target sm_12 
.target texmode_independent 

.reg .b32 r<126>; /* define r0..125 */ 
.reg .b64 x<126>; /* define r0..125 */ 
.reg .b32 f<128>; /* define f0..127 */ 
.reg .pred p<32>; /* define p0..31 */ 
.reg .u32 sp; 

.reg .b8 wb0,wb1,wb2,wb3; /* 8-bit write buffer */ 
.reg .b16 ws0,ws1,ws2,ws3; /* 16-bit write buffer */ 
.reg .b32 tb0,tb1,tb2,tb3; /* read tex buffer */ 
.reg .b64 vl0,vl1; /* 64-bit vector buffer */ 
.reg .b16 cvt16_0,cvt16_1; /* tmps for conversions */ 


.const .align 1 .b8 ndft_gid_base[52]; 
.local .align 16 .b8 ndft_stack[8]; 
.entry ndft(
    .param.b32 ndft_0 /* re */, 
    .param.b32 ndft_1 /* im */, 
    .param.b32 ndft_2 /* num_values */, 
    .param.b32 ndft_3 /* spectrum_re */, 
    .param.b32 ndft_4 /* spectrum_im */, 
    .param.b32 ndft_5 /* spectrum_abs */, 
    .param.b32 ndft_6 /* sin_array */, 
    .param.b32 ndft_7 /* cos_array */, 
    .param.f32 ndft_8 /* sqrt_num_values_reciprocal */ 
) { 
    mov.u32 sp, ndft_stack; 
    mov.u32 r0, 4294967295; 
    ld.param.u32 r1, [ndft_2 + 0]; 
LBB1_1: 
    add.u32 r0, r0, 1; 
    setp.lt.s32 p0, r0, r1; 
    @p0 bra LBB1_1; 
LBB1_2: 
    ret; 
} 

tức là. một nhánh hạt nhân không chứa vòng lặp tính toán. Khi ba bộ nhớ toàn cầu viết trong ba dòng cuối cùng của hạt nhân được uncomment, nó phát ra này:

Program 0: 
S.version 1.5inaryDriver\clBinaryData_clBinaryVersionWCLH 1.0O 
.target sm_12 
.target texmode_independent 

.reg .b32 r<126>; /* define r0..125 */ 
.reg .b64 x<126>; /* define r0..125 */ 
.reg .b32 f<128>; /* define f0..127 */ 
.reg .pred p<32>; /* define p0..31 */ 
.reg .u32 sp; 

.reg .b8 wb0,wb1,wb2,wb3; /* 8-bit write buffer */ 
.reg .b16 ws0,ws1,ws2,ws3; /* 16-bit write buffer */ 
.reg .b32 tb0,tb1,tb2,tb3; /* read tex buffer */ 
.reg .b64 vl0,vl1; /* 64-bit vector buffer */ 
.reg .b16 cvt16_0,cvt16_1; /* tmps for conversions */ 


.const .align 1 .b8 ndft_gid_base[52]; 
.local .align 16 .b8 ndft_stack[8]; 
.entry ndft(
    .param.b32 ndft_0 /* re */, 
    .param.b32 ndft_1 /* im */, 
    .param.b32 ndft_2 /* num_values */, 
    .param.b32 ndft_3 /* spectrum_re */, 
    .param.b32 ndft_4 /* spectrum_im */, 
    .param.b32 ndft_5 /* spectrum_abs */, 
    .param.b32 ndft_6 /* sin_array */, 
    .param.b32 ndft_7 /* cos_array */, 
    .param.f32 ndft_8 /* sqrt_num_values_reciprocal */ 
) { 
    mov.u32 sp, ndft_stack; 
    cvt.u32.u16 r0, %tid.x; 
    cvt.u32.u16 r1, %ntid.x; 
    cvt.u32.u16 r2, %ctaid.x; 
    mad24.lo.u32 r0, r2, r1, r0; 
    mov.u32 r1, 0; 
    shl.b32 r2, r1, 2; 
    mov.u32 r3, ndft_gid_base; 
    add.u32 r2, r2, r3; 
    ld.const.u32 r2, [r2 + 40]; 
    add.u32 r0, r0, r2; 
    ld.param.u32 r2, [ndft_2 + 0]; 
    mul.lo.u32 r3, r0, r2; 
    shl.b32 r3, r3, 2; 
    mov.f32 f0, 0f00000000 /* 0.000000e+00 */; 
    ld.param.f32 f1, [ndft_8 + 0]; 
    ld.param.u32 r4, [ndft_7 + 0]; 
    ld.param.u32 r5, [ndft_6 + 0]; 
    ld.param.u32 r6, [ndft_5 + 0]; 
    ld.param.u32 r7, [ndft_4 + 0]; 
    ld.param.u32 r8, [ndft_3 + 0]; 
    ld.param.u32 r9, [ndft_1 + 0]; 
    ld.param.u32 r10, [ndft_0 + 0]; 
    mov.u32 r11, r1; 
    mov.f32 f2, f0; 
LBB1_1: 
    setp.ge.s32 p0, r11, r2; 
    @!p0 bra LBB1_7; 
LBB1_2: 
    shl.b32 r1, r0, 2; 
    add.u32 r2, r8, r1; 
    st.global.f32 [r2+0], f0; 
    add.u32 r1, r7, r1; 
    st.global.f32 [r1+0], f2; 
    abs.f32 f1, f2; 
    abs.f32 f0, f0; 
    setp.gt.f32 p0, f0, f1; 
    selp.f32 f2, f0, f1, p0; 
    abs.f32 f3, f2; 
    mov.f32 f4, 0f7E800000 /* 8.507059e+37 */; 
    setp.gt.f32 p1, f3, f4; 
    selp.f32 f0, f1, f0, p0; 
    shl.b32 r0, r0, 2; 
    add.u32 r0, r6, r0; 
    @!p1 bra LBB1_8; 
LBB1_3: 
    mul.rn.f32 f3, f2, 0f3E800000 /* 2.500000e-01 */; 
    mul.rn.f32 f1, f0, 0f3E800000 /* 2.500000e-01 */; 
LBB1_4: 
    mov.f32 f4, 0f00000000 /* 0.000000e+00 */; 
    setp.eq.f32 p0, f2, f4; 
    @!p0 bra LBB1_9; 
LBB1_5: 
    add.f32 f1, f2, f0; 
LBB1_6: 
    mov.f32 f3, 0f7F800000 /* inf */; 
    setp.eq.f32 p0, f0, f3; 
    setp.eq.f32 p1, f2, f3; 
    or.pred p0, p1, p0; 
    selp.f32 f0, f3, f1, p0; 
    st.global.f32 [r0+0], f0; 
    ret; 
LBB1_7: 
    add.u32 r12, r3, r1; 
    add.u32 r13, r4, r12; 
    ld.global.f32 f3, [r13+0]; 
    mul.rn.f32 f3, f3, f1; 
    add.u32 r13, r9, r1; 
    ld.global.f32 f4, [r13+0]; 
    mul.rn.f32 f5, f3, f4; 
    add.u32 r12, r5, r12; 
    ld.global.f32 f6, [r12+0]; 
    mul.rn.f32 f6, f6, f1; 
    add.u32 r12, r10, r1; 
    ld.global.f32 f7, [r12+0]; 
    mul.rn.f32 f8, f7, f6; 
    add.f32 f5, f8, f5; 
    sub.f32 f2, f2, f5; 
    mul.rn.f32 f4, f4, f6; 
    mul.rn.f32 f3, f7, f3; 
    add.f32 f3, f3, f4; 
    add.f32 f0, f0, f3; 
    add.u32 r11, r11, 1; 
    add.u32 r1, r1, 4; 
    bra LBB1_1; 
LBB1_8: 
    mov.f32 f1, f0; 
    mov.f32 f3, f2; 
    bra LBB1_4; 
LBB1_9: 
    div.approx.f32 f1, f1, f3; 
    mul.rn.f32 f1, f1, f1; 
    add.f32 f1, f1, 0f3F800000 /* 1.000000e+00 */; 
    sqrt.approx.ftz.f32 f1, f1; 
    mul.rn.f32 f1, f2, f1; 
    bra LBB1_6; 
} 

Tôi nghĩ rằng đây là bằng chứng khá không thể chối cãi rằng đó là trình biên dịch tối ưu hóa mà là làm cho sự khác biệt trong thời gian chạy, và phụ thuộc chỉ vào việc ghi nhớ được bao gồm trong mã hạt nhân hay không.

---

Tôi đoán câu hỏi cuối cùng sau đó trở thành lý do tại sao điều này quá chậm (không phân biệt cuộc tranh luận về việc điều này có gây ra bởi tối ưu hóa trình biên dịch hay không). Thời gian chạy 1,5 mili giây bạn đang thấy là một sự phản ánh thực sự về hiệu suất của mã và câu hỏi thực sự là lý do tại sao. Từ đọc mã hạt nhân của tôi, câu trả lời có vẻ nằm trong các mẫu truy cập bộ nhớ khá khủng khiếp đối với GPU. Bên trong vòng lặp tính toán bạn có một bộ nhớ toàn cầu hai lần đọc với những bước tiến rất lớn, như thế này một:

x = cos_array[thread_id * num_values + i] * sqrt_num_values_reciprocal; 

Theo nhận xét trong mã của bạn num_values hoặc là 24 hoặc 48. Điều đó có nghĩa là bộ nhớ đọc có thể không có khả năng kết hợp lại, và bộ đệm L1 trên GPU Fermi cũng sẽ không giúp được gì nhiều. Điều này sẽ có tác động tiêu cực rất lớn đến việc sử dụng băng thông bộ nhớ và làm cho mã rất chậm. Nếu bạn bị mắc kẹt với thứ tự dữ liệu đầu vào đó, thì một giải pháp nhanh hơn sẽ là sử dụng một sợi dọc để thực hiện tính toán của một đầu ra (do đó làm giảm toàn bộ sợi dọc thành tổng cuối cùng). Điều này sẽ làm giảm bước đọc từ 24 hoặc 48 xuống 1 và kết hợp bộ nhớ toàn cầu đọc từ hai mảng đầu vào lớn đó.

Bên trong vòng lặp cũng được lặp lại lấy về vào bộ nhớ toàn cầu cho cả hai 24 hoặc 48 yếu tố của re và im:

sum_re += re[i] * x + im[i] * y; 
    sum_im -= re[i] * y + x * im[i]; 

này là không cần thiết và lãng phí, rất nhiều băng thông bộ nhớ toàn cục hoặc hiệu quả bộ nhớ cache (các GPU không có đủ thanh ghi để cho phép trình biên dịch giữ toàn bộ mỗi mảng trong thanh ghi). Sẽ tốt hơn nếu mỗi nhóm làm việc đọc hai mảng đó thành __local mảng bộ nhớ một lần và sử dụng bản sao bộ nhớ cục bộ bên trong vòng lặp tính toán. Nếu bạn có từng nhóm công việc tính toán nhiều lần, thay vì chỉ một lần, thì bạn có khả năng có thể tiết kiệm rất nhiều băng thông bộ nhớ toàn cầu và phân bổ đọc ban đầu cho đến khi nó gần như là miễn phí.