ex6: fix correct handling of remainder pixels

exercise6/mandel.cu

@@ -30,16 +30,23 @@ typedef struct {
 } my_complex_t;
 
 #define PIXEL(i, j) ((i) + (j) * XSIZE)
+#define CEILDIV(a, b) (((a) + (b) - 1) / (b))
 
 /********** SUBTASK1: Create kernel device_calculate *************************/
 
-__global__ void device_calculate(double step, double xleft, double yupper, int *gpu_buf) {
+__global__ void device_calculate(
+    double step, double xleft, double yupper, int *gpu_buf) {
   // assumes the threads are organized in a grid, such that each thread
   // calculates a single pixel. thus a block corresponds to a region of the
   // output image.
   int i = threadIdx.x + blockIdx.x * blockDim.x;
   int j = threadIdx.y + blockIdx.y * blockDim.y;
 
+  // combined with CEILDIVed dimensions, we handle the remainder pixels when the
+  // divisibility doesn't work out to an integer.
+  if (i >= XSIZE || j >= YSIZE)
+    return;
+
   // same as inner-most part of host_calculate
   my_complex_t c, z, temp;
   int iter = 0;
@@ -162,7 +169,7 @@ int main(int argc, char **argv) {
   start = walltime();
   /********** SUBTASK3: Execute the kernel on the device *******************/
 
-  dim3 grid(XSIZE / BLOCKX, YSIZE / BLOCKY);
+  dim3 grid(CEILDIV(XSIZE, BLOCKX), CEILDIV(YSIZE, BLOCKY));
   dim3 block(BLOCKX, BLOCKY);
   device_calculate<<<grid, block>>>(step, xleft, yupper, gpu_buf);
 
@@ -218,8 +225,8 @@ int main(int argc, char **argv) {
     /* create nice image from iteration counts. take care to create it upside
              down (bmp format) */
     unsigned char *buffer = (unsigned char *)calloc(XSIZE * YSIZE * 3, 1);
-    for (int i = 0; i < XSIZE; i++) {
-      for (int j = 0; j < YSIZE; j++) {
+    for (int j = 0; j < YSIZE; j++) {
+      for (int i = 0; i < XSIZE; i++) {
         int p = ((YSIZE - j - 1) * XSIZE + i) * 3;
         fancycolour(buffer + p, device_pixel[PIXEL(i, j)]);
       }