title: CUDA: Best number of pixel computed per thread (grayscale)
body: <p>I'm working on a program to convert an image in grayscale. I'm using the CImg library. I have to read for each pixel, the 3 values R-G-B, calculate the corresponding gray value and store the gray pixel on the output image. I'm working with an <strong>NVIDIA GTX 480</strong>. Some details about the card:</p>
  
  <ul>
  <li>Microarchitecture: <strong>Fermi</strong></li>
  <li>Compute capability (version): <strong>2.0</strong></li>
  <li>Cores per SM (warp size): <strong>32</strong></li>
  <li>Streaming Multiprocessors: <strong>15</strong></li>
  <li>Maximum number of resident warps per multiprocessor: <strong>48</strong></li>
  <li>Maximum amount of shared memory per multiprocessor: <strong>48KB</strong></li>
  <li>Maximum number of resident threads per multiprocessor: <strong>1536</strong></li>
  <li>Number of 32-bit registers per multiprocessor: <strong>32K</strong></li>
  </ul>
  
  <p>I'm using a square grid with blocks of 256 threads.
  This program can have as input images of different sizes (e.g. 512x512 px, 10000x10000 px). I observed that incrementing the number of the pixels assigned to each thread increments the performance, so it's better than compute one pixel per thread. The problem is, how can I determine the number of pixels to assign to each thread statically? Computing tests with every possible number? I know that on the GTX 480, 1536 is the maximum number of resident threads per multiprocessor. Have I to consider this number? The following, is the code executed by the kernel.</p>
  
  <pre><code>for(i = ((gridDim.x + blockIdx.x) * blockDim.x) + threadIdx.x; i &lt; width * height; i += (gridDim.x * blockDim.x)) {
      float grayPix = 0.0f;
      float r = static_cast&lt; float &gt;(inputImage[i]);
      float g = static_cast&lt; float &gt;(inputImage[(width * height) + i]);
      float b = static_cast&lt; float &gt;(inputImage[(2 * width * height) + i]);
  
      grayPix = ((0.3f * r) + (0.59f * g) + (0.11f * b));
      grayPix = (grayPix * 0.6f) + 0.5f;
      darkGrayImage[i] = static_cast&lt; unsigned char &gt;(grayPix);
  }
  </code></pre>