title: OpenCL output buffer
body: <p>I am trying to learn OpenCL basics.</p>
  
  <p>I thought this code in a kernel:</p>
  
  <pre><code>out[ 1 &amp; ((a+b)==(b+a)) ] = (char)1;
  </code></pre>
  
  <p>Would generate the same results as:</p>
  
  <pre><code>out[ 1 ] = (char)1;
  </code></pre>
  
  <p>Can anybody tell me why it generates different result?</p>
  
  <p>Here is the full source code for my kernel:</p>
  
  <pre><code>#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable
  __kernel void hello(__global char * out)
  {
      size_t tid = get_global_id(0);
      int a = tid &amp; 0xff;
      int b = (tid &gt;&gt; 8) &amp; 0xff;
      out[ 1 &amp; ((a+b)==(b+a)) ] = (char)1;
  }
  </code></pre>
  
  <p>If I replace that last statement with "out[1]=(char)1;" then "out[0]" will not be written. But the kernel above will write "1" to out[0].</p>
  
  <p>EDIT: This is my C++ code:</p>
  
  <pre><code>#include &lt;utility&gt;
  #define __NO_STD_VECTOR // Use cl::vector instead of STL version
  #include &lt;CL/cl.hpp&gt;
  
  #include &lt;cstdio&gt;
  #include &lt;cstdlib&gt;
  #include &lt;fstream&gt;
  #include &lt;iostream&gt;
  #include &lt;string&gt;
  #include &lt;iterator&gt;
  
  inline void checkErr(cl_int err, const char * name)
  {
      if (err != CL_SUCCESS) {
          std::cerr &lt;&lt; "ERROR: " &lt;&lt; name
                   &lt;&lt; " (" &lt;&lt; err &lt;&lt; ")" &lt;&lt; std::endl;
          std::exit(EXIT_FAILURE);
      }
  }
  
  int main()
  {
      cl_int err;
  
      cl::vector&lt; cl::Platform &gt; platformList;
      cl::Platform::get(&amp;platformList);
      checkErr(platformList.size()!=0 ? CL_SUCCESS : -1, "cl::Platform::get");
      std::cerr &lt;&lt; "Platform number is: " &lt;&lt; platformList.size() &lt;&lt; std::endl;
  
      std::string platformVendor;
      platformList[0].getInfo((cl_platform_info)CL_PLATFORM_VENDOR, &amp;platformVendor);
      std::cerr &lt;&lt; "Platform is by: " &lt;&lt; platformVendor &lt;&lt; "\n";
      cl_context_properties cprops[3] = 
          {CL_CONTEXT_PLATFORM, (cl_context_properties)(platformList[0])(), 0};
  
      cl::Context context(
         CL_DEVICE_TYPE_GPU, 
         cprops,
         NULL,
         NULL,
         &amp;err);
      checkErr(err, "Conext::Context()"); 
  
  
      unsigned char outH[2] = {0};
      cl::Buffer outCL(
          context,
          CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR,
          sizeof(outH),
          outH,
          &amp;err);
      checkErr(err, "Buffer::Buffer()");
  
  
      cl::vector&lt;cl::Device&gt; devices;
      devices = context.getInfo&lt;CL_CONTEXT_DEVICES&gt;();
      checkErr(
          devices.size() &gt; 0 ? CL_SUCCESS : -1, "devices.size() &gt; 0");
  
  
      std::ifstream file("condition1.cl");
      checkErr(file.is_open() ? CL_SUCCESS:-1, "condition1.cl");
      const std::string prog(std::istreambuf_iterator&lt;char&gt;(file), (std::istreambuf_iterator&lt;char&gt;()));
  
      cl::Program::Sources source(1, std::make_pair(prog.c_str(), prog.length()+1));
      cl::Program program(context, source);
      err = program.build(devices,""); 
      checkErr(err, "Program::build()"); 
  
  
      cl::Kernel kernel(program, "hello", &amp;err);
      checkErr(err, "Kernel::Kernel()");
  
      err = kernel.setArg(0, outCL);
      checkErr(err, "Kernel::setArg()");
  
  
      cl::CommandQueue queue(context, devices[0], 0, &amp;err);
      checkErr(err, "CommandQueue::CommandQueue()");
  
  
      cl::Event event;
      err = queue.enqueueNDRangeKernel(
          kernel, 
          cl::NullRange,
          cl::NDRange(65536),
           cl::NDRange(1, 1), 
          NULL, 
          &amp;event);
      checkErr(err, "ComamndQueue::enqueueNDRangeKernel()");
  
      event.wait();    
      err = queue.enqueueReadBuffer(
          outCL,
          CL_TRUE,
          0,
          sizeof(outH),
          outH);
      checkErr(err, "ComamndQueue::enqueueReadBuffer()");
  
      for (int i = 0; i &lt; sizeof(outH); i++)
          std::cout &lt;&lt; (int)outH[i] &lt;&lt; " ";
  
      std::string str;
      std::getline(std::cin, str);
  
      return EXIT_SUCCESS;
  }
  </code></pre>
  
  <p>EDIT2:</p>
  
  <p>It is not caused by "undefined behaviour" at least. I get weird results for this kernel code too:</p>
  
  <pre><code>char result = 0;
  for (int a = 0; a &lt; 2; a++) {
      for (int b = 0; b &lt; 2; b++) {
          if ((a+b) != (b+a))
              result |= (1 &lt;&lt; (a+2*b));
      }
  }
  </code></pre>
  
  <p>I would expect that "result" will get the value 0. But it gets the value 6. If I change the != to an == then result gets the value 9. So in this code, when "a" is not equal with "b" then (a+b) does not equal (b+a).</p>
  
  <p>If I change the code and set a known value for either "a" or "b" then the result will be 0 as I expect. For example:</p>
  
  <pre><code>char result = 0;
  int a = 1;
  /*for (int a = 0; a &lt; 2; a++)*/ {
      for (int b = 0; b &lt; 2; b++) {
          if ((a+b) != (b+a))
              result |= (1 &lt;&lt; (a+2*b));
      }
  }
  </code></pre>