title: How to determine max size of images I can safely pass to/from OpenCL kernel?
body: <p>I'm developing an <strong>OpenCL 1.2</strong> application that deals with large imagery. 
  At the moment, the image I'm testing with is 16507x21244 pixels. My <strong>kernel is run in a loop that operates on chunks of the image</strong>. The kernel takes 32bpp (rgba) chunks of the image in, and passes float4-pixel chunks out.</p>
  
  <p>Let's define one side of a (square) chunk in pixels to be the chunk size. That is to say, an <strong>8192x8192 pixel square has chunk size 8192</strong>. Of course, on the right and bottom sides we have smaller rectangular chunks if the image is not cleanly divisable by the chunk size. My code handles this, but for the rest of this post, let us ignore this for simplicity.</p>
  
  <p>I am <strong>trying to determine the maximum chunk size</strong> that I can operate on in each iteration of my loop, as well as the optimal chunk size (which may not be the maximum chunk size).</p>
  
  <p>For reference, <a href="https://gist.github.com/miotatsu/f57d3c9e7c8b0680c874">here is the information reported by the clinfo utility on my machine</a>. I am running my kernel on the <code>Geforce GTX 560 Ti</code> with the Nvidia platform using their proprietary linux drivers.</p>
  
  <p>My initial naïve assumption was that I could operate on the max 2d image size.
  However, this results in <code>clEnqueueNDRangeKernel</code> returning an error code of -4 (<code>CL_MEM_OBJECT_ALLOCATION_FAILURE</code>).</p>
  
  <p>Thinking about it, this makes sense to me. With 1 GiB of video memory, one would expect to be able to hold a single 16384x16384 pixel texture (32bpp), or an 8192x8192 pixel texture (float4). If both need to be cached on the card while the kernel runs, we could expect to use the following amount of memory:</p>
  
  <pre><code>   4 bytes-per-pixel * chunk size^2 (input image) 
  + 16 bytes-per-pixel * chunk size^2 (output image) 
  = 1 GiB total video memory
  </code></pre>
  
  <p>Solving for chunk size we get</p>
  
  <pre><code>chunk size = sqrt(1GiB/20)
  </code></pre>
  
  <p>Plugging in the amount of memory reported by OpenCL (which is slightly less than 1GiB - 1023 MiB) and flooring the result, we get:</p>
  
  <pre><code>floor(sqrt(1072889856/20)) = 7324
  </code></pre>
  
  <p>However, a chunk size of 7324 still results in <code>CL_MEM_OBJECT_ALLOCATION_FAILURE</code>.</p>
  
  <p>My next guess was that we can't pass an image larger than the <strong>max allocation size</strong>, which OpenCL reports as 268222464 bytes for my card. Because my output image has the larger pixel width, it would dictate my chunk size.</p>
  
  <pre><code>floor(sqrt(268222464/16)) = 4094
  </code></pre>
  
  <p>Hey, that actually works! <strong>Now what if we try to go larger? To my surprise, it doesn't fail</strong>. Through trial and error, I narrowed in on 6784 as the actual max chunk size. At 6785, it begins complaining with <code>CL_MEM_OBJECT_ALLOCATION_FAILURE</code>. I do not know why the max appears to be 6784, and I do not know if this is repeatable or if the value fluctuates (such as other state existing in the video memory affecting how much it can hold.) I also find that running with a chunk size of 6784 is a few seconds slower than running with the size based on max allocation. I wonder if this is because OpenCL is needing to perform multiple (expensive) allocations under-the-hood? <strike>I also noticed the "max size of kernel argument" that OpenCL is able to report (<code>CL_DEVICE_MAX_PARAMETER_SIZE</code>). However, that value seems bogus. If I could only pass 4096 bytes in, that would limit me to 16x16 pixels!</strike></p>
  
  <p>So I'm left with two fundamental questions:</p>
  
  <ol>
  <li><strong>How do I determine the absolute maximum chunk size?</strong></li>
  <li><strong>How do I determine the fastest chunk size? (Is there a method other than trial and error?)</strong></li>
  </ol>
  
  <p>As a bonus question, are there any <strong>good resources</strong> I could turn to for future questions of this nature regarding <strong>low-level OpenCL-hardware interactions?</strong></p>
  
  <p>And finally, I will provide some code snippets for peer-review; I would be extremely grateful for any constructive criticism!</p>
  
  <ul>
  <li><a href="https://gist.github.com/miotatsu/5dfc2a3104aeb2add3ac">The host code module (dlang)</a></li>
  <li><a href="https://gist.github.com/miotatsu/82ea8bfdef82d0fad02e">Wrapper that sets up the OpenCL program (dlang)</a></li>
  <li><a href="https://gist.github.com/miotatsu/040a2ec8c2dfdc101a57">The kernel code</a></li>
  </ul>
  
  <p>As always, <strong>thanks in advance for any help!</strong></p>