title: Find max of matrix in CUDA
body: <p>I just started in CUDA. Now I have a question.
  I have N*N matrix, and a window scale is 8x8. I want subdivided this matrix into multiple sub-matrix and find max value of this.
  For example if I have 64*64 matrix so I will have 8 small matrix with 8*8 scale and find out 8 max values. Finally I save all max values into new array, but its order always change. I want find solution to keep them in right order</p>
  
  <pre><code>__global__ void calculate_emax_kernel(float emap[],float emax[], int img_height, int img_width,int windows_size)
  {
      int x_index = blockIdx.x*blockDim.x+threadIdx.x;
      int y_index = blockIdx.y*blockDim.y+threadIdx.y;
  
      int num_row_block = img_height/windows_size;
      int num_col_block = img_width/windows_size;
      __shared__ float window_elements[256];
      __shared__ int counter;
      __shared__ int emax_count;
  
      if (threadIdx.x == 0) emax_count = 0;
      __syncthreads();
      int index;
      int emax_idx = 0;
  
  
      if(y_index &gt;= img_height|| x_index &gt;= img_width) return;
      for(int i = 0; i &lt; num_row_block; i++)
      {
          for(int j = 0; j &lt; num_col_block; j++)
          {
              counter = 0;
              if(y_index &gt;= i*windows_size &amp;&amp; y_index &lt; (i+1)*windows_size
                      &amp;&amp; x_index &gt;= j*windows_size &amp;&amp; x_index &lt; (j+1)*windows_size)
              {
                  int idx = y_index*img_height + x_index;
                  index = atomicAdd(&amp;counter, 1);
  
                  window_elements[index] = emap[idx];
                  __syncthreads();
  
  
                  // reduction
                  unsigned int k = (windows_size*windows_size)/2;
                  while(k != 0)
                  {
                      if(index &lt; k)
                      {
                          window_elements[index] = fmaxf(window_elements[index], window_elements[index+k]);
  
                      }
                      k /= 2;
                  }
                  if(index == 0)
                  {
                      emax[i*num_row_block+j] = window_elements[index];
                  }
              }
              __syncthreads();
          }
          __syncthreads();
      }
      __syncthreads();
  }
  </code></pre>
  
  <p>This is my configuration</p>
  
  <pre><code>void construct_emax(float *input,float *output, int img_height, int img_width)
  {
      int windows_size = 4;
      float * d_input, * d_output;
      cudaMalloc(&amp;d_input, img_width*img_height*sizeof(float));
      cudaMalloc(&amp;d_output, img_width*img_height*sizeof(float));
  
      cudaMemcpy(d_input, input, img_width*img_height*sizeof(float), cudaMemcpyHostToDevice);
      dim3 blocksize(16,16);
      dim3 gridsize;
  
      gridsize.x=(img_width+blocksize.x-1)/blocksize.x;
      gridsize.y=(img_height+blocksize.y-1)/blocksize.y;
  
      calculate_emax_kernel&lt;&lt;&lt;gridsize,blocksize&gt;&gt;&gt;(d_input,d_output,img_height,img_width,windows_size);
  
  }
  </code></pre>