title: Writing to global or local memory increases kernel execution time by 10000 %
body: <p>I have the following OpenCL kernel:</p>
  
  <pre><code>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 &lt; 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
  }</code></pre>
  
  <p>Like this, the execution time is about 15 us (work group size = 567, 14 work groups, for a total of 7938 threads).</p>
  
  <p>However, of course, I somehow need to retrieve the results of the operation, which is what the last few lines are for, (commented out). As soon as I perform a single of those memory operations (and it does not matter if <code>spectrum_X</code> is <code>global</code>, as in the example, or <code>local</code>), the exeuction time of the kernel increases to ~1.4 to 1.5 ms.</p>
  
  <p>I thought the increase in execution time was some kind of fixed overhead, so I would just accumulate more data, so that the relative amount of time lost due to that effect minimizes. But when I double my number of threads (i. e. twice the amount of data), the execution time also doubles (to 2.8 ~ 3.0 ms).</p>
  
  <p>I found out, that even if I only uncomment <em>one</em> of the those lines, I have the same execution time as if I uncommented all three. Even if I add a <code>if (thread_id == 0)</code> and run it, I have the same execution time. However, it is just way too slow this way (the upper limit for my application is about 30 us). It even perfoms about 5 times faster when I run it in ordinary C code on my CPU.</p>
  
  <p>Now I am obviously doing something wrong but I am not sure where to start looking for a solution.</p>
  
  <hr>
  
  <p>As I commented on talonmies' answer, I also did the following:</p>
  
  <p>From the above code, I made the last 4 lines look like</p>
  
  <pre><code>//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);</code></pre>
  
  <p>As expected, execution time ~1.8 ms.
  The generated assembler code for my system is:</p>
  
  <pre><code>//
  // 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;      // &lt;-- 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;    // &lt;-- stores the hypot's result in spectrum_abs
      ret;
  }
  </code></pre>
  
  <p>Indeed all my calculation operations are there - lots of adds/mults as well as a <code>sqrt</code> for the <code>hypot</code> function. From the above asm code, I removed the second last line:</p>
  
  <p><code>st.global.f32  [%r34], %f43;</code></p>
  
  <p>which is the line that actually stores the data in the global array <code>spectrum_abs</code>. Then I used <code>clCreateProgramWithBinary</code> and used the modified asm code file as input. Execution time went down to 20 us.</p>