वैश्विक या स्थानीय स्मृति के लिए लेखन 10000% से बढ़ जाती है कर्नेल निष्पादन समय

Question

मैं निम्नलिखित OpenCL गिरी है:

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 
}

इस तरह, निष्पादन समय के बारे में 15 हमें (कार्य समूह आकार = 567, 14 कार्य समूहों है , कुल 7 9 38 धागे के लिए)।

हालांकि, मुझे किसी भी तरह ऑपरेशन के परिणामों को पुनर्प्राप्त करने की आवश्यकता है, जो पिछले कुछ पंक्तियों के लिए है, (टिप्पणी की गई)। जैसे ही मैं उन स्मृति कार्यों में से एक करता हूं (और इससे कोई फर्क नहीं पड़ता कि spectrum_Xglobal है, उदाहरण के लिए, या local), कर्नेल का बहिष्करण समय ~ 1.4 से 1.5 एमएस तक बढ़ जाता है।

मैंने सोचा कि निष्पादन समय में वृद्धि कुछ प्रकार का निश्चित ओवरहेड था, इसलिए मैं केवल अधिक डेटा जमा करूंगा, ताकि उस प्रभाव के कारण खोए गए समय की सापेक्ष मात्रा कम हो। लेकिन जब मैं अपने धागे की संख्या दोगुना करता हूं (i. ई। डेटा की मात्रा से दोगुना), निष्पादन समय भी दोगुना हो जाता है (2.8 ~ 3.0 एमएस तक)।

मुझे पता चला कि अगर मैं केवल उन पंक्तियों के एक को असम्बद्ध करता हूं, तो मेरे पास एक ही निष्पादन समय है जैसे कि मैंने तीनों को अपूर्ण किया है। यहां तक ​​कि अगर मैं if (thread_id == 0) जोड़ता हूं और इसे चलाता हूं, तो मेरे पास एक ही निष्पादन समय होता है। हालांकि, यह इस तरह से बहुत धीमा तरीका है (मेरे आवेदन के लिए ऊपरी सीमा लगभग 30 है)। जब मैं इसे अपने सीपीयू पर सामान्य सी कोड में चलाता हूं तो यह लगभग 5 गुना तेज होता है।

अब मैं स्पष्ट रूप से कुछ गलत कर रहा हूं लेकिन मुझे यकीन नहीं है कि समाधान कहां से शुरू करना है।

---

जैसा कि मैंने talonmies 'उत्तर पर टिप्पणी की, मैं भी निम्नलिखित किया:

ऊपर कोड से, मैं पिछले 4 लाइनों

//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);

जैसी उम्मीद थी की तरह लग रहे बनाया है, निष्पादन समय ~ 1.8 एमएस अपने सिस्टम के लिए उत्पन्न कोडांतरक कोड है:

// 
// 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; 
} 

दरअसल मेरे सभी गणना संचालन कर रहे हैं वहाँ -/कहते हैं की बहुत सारी mults के साथ-साथ hypot समारोह के लिए एक sqrt।

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

जो लाइन है कि वास्तव में वैश्विक सरणी spectrum_abs में डेटा संग्रहीत करता है: उपरोक्त एएसएम कोड से, मैं दूसरा अंतिम पंक्ति हटा दिया। तब मैंने clCreateProgramWithBinary का उपयोग किया और इनपुट के रूप में संशोधित एएसएम कोड फ़ाइल का उपयोग किया। निष्पादन का समय 20 तक गिर गया।

Answer 1

मुझे लगता है कि आप कंपाइलर अनुकूलन के प्रभाव देख रहे हैं।

एनवीआईडीआईए संकलक "मृत कोड" को खत्म करने में बहुत आक्रामक है जो सीधे वैश्विक स्मृति को लिखने में भाग नहीं लेता है। तो अपने कर्नेल में, यदि आप sum_re या sum_im नहीं लिखते हैं, तो संकलक पूरे गणना लूप (और शायद बाकी सब कुछ) को अनुकूलित करेगा और एक खाली कर्नेल के साथ छोड़ देगा जिसमें नो-ऑप से अधिक कुछ नहीं है। 15 माइक्रोसॉन्ड निष्पादन समय जो आप देख रहे हैं वह ज्यादातर कर्नेल लॉन्च ओवरहेड है और बहुत कुछ नहीं। जब आप वैश्विक मेमोरी लिखते हैं, तो कंपाइलर सभी गणना कोड को जगह में छोड़ देता है, और आप अपने कोड का सही निष्पादन समय देखते हैं।

तो असली सवाल यह है कि आपको शायद यह पूछना चाहिए कि उस निष्कर्ष को 1.5 मिलीसेकंड से अपने निष्पादन समय को कम करने के लिए कैसे अनुकूलित किया जाए, जो वर्तमान में आपके (बहुत महत्वाकांक्षी) 30 माइक्रोसॉन्ड लक्ष्य की ओर ले जाता है।

---

संदेह के बावजूद मूल जवाब देने के लिए व्यक्त की, यहाँ एक पूरा रेप्रो मामले में जो दावे का समर्थन करता है कि यह एक संकलक से संबंधित प्रभाव है:

#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; 
} 

जब संकलित और चलाने के लिए, यह पालन का उत्सर्जन करता है ओपनसीएल रनटाइम से पीटीएक्स कोड:

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; 
} 

यानी। एक कर्नेल स्टब जिसमें गणना गणना लूप नहीं है। जब तीन वैश्विक स्मृति गिरी के अंतिम तीन लाइनों में लिखते हैं uncommented कर रहे हैं, यह इस का उत्सर्जन करता है:

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; 
} 

मुझे लगता है कि यह बहुत अकाट्य सबूत है कि यह संकलक अनुकूलन जो क्रम में अंतर का कारण बनता है है, और निर्भर करता है केवल स्मृति संदेश लिखने के लिए कर्नेल कोड में शामिल हैं या नहीं।

---

मुझे लगता है कि अंतिम सवाल तो बन जाता है यही कारण है कि यह इतना धीमा है (के बारे में इस या नहीं संकलक अनुकूलन के कारण होता है कि क्या बहस के बावजूद)। 1.5 मिलीसेकंड रनटाइम जो आप देख रहे हैं वह कोड के प्रदर्शन का सही प्रतिबिंब है और असली सवाल यह है कि क्यों। आपके कर्नेल कोड के पढ़ने से, उत्तर मेमोरी एक्सेस पैटर्न में झूठ बोलता है जो GPU के लिए बहुत भयानक हैं। गणना पाश के अंदर आप एक दो वैश्विक स्मृति, बहुत बड़ी प्रगति के साथ पढ़ता इस तरह है: अपने कोड में टिप्पणी के अनुसार num_values या तो 24 या 48. मतलब यह है कि स्मृति पढ़ता

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

संभवतः नहीं कर सकते coalesce, और एक Fermi जीपीयू पर एल 1 कैश या तो बहुत मदद नहीं होने जा रहा है। मेमोरी बैंडविड्थ उपयोग पर इसका बहुत नकारात्मक प्रभाव होगा और कोड को बहुत धीमा कर देगा। यदि आप उस इनपुट डेटा ऑर्डरिंग से फंस गए हैं, तो एक आउटपुट की गणना करने के लिए एक तेज समाधान का उपयोग करना होगा (इसलिए अंतिम योग में एक व्यापक व्यापक कमी करें)। यह 24 या 48 से 1 तक पढ़ने के चरण को कम कर देगा और उन दो बड़े इनपुट सरणी से वैश्विक स्मृति को पढ़ता है।

पाश वहाँ भी दोहराया है अंदर re के दोनों 24 या 48 तत्वों और im के लिए वैश्विक स्मृति को हासिल करेगा:

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

यह अनावश्यक है, और अपशिष्ट वैश्विक स्मृति बैंडविड्थ या कैश क्षमता का एक बहुत (जीपीयू में पर्याप्त रजिस्ट्रार नहीं हैं ताकि संकलक को रजिस्टर में प्रत्येक सरणी को पूरा किया जा सके)। प्रत्येक कार्य समूह को उन दो सरणी को __local मेमोरी सरणी में एक बार पढ़ने और गणना लूप के अंदर स्थानीय मेमोरी प्रति का उपयोग करना बेहतर होगा। यदि आपके पास प्रत्येक कार्य समूह एक बार की बजाय कई बार गणना करता है, तो आप संभावित रूप से बहुत सारी वैश्विक मेमोरी बैंडविड्थ को सहेज सकते हैं और शुरुआती पढ़ने को लगभग मुक्त कर सकते हैं।