Fix bug in identity cuda plaguing compute 5.2

shehzan10 · shehzan10 · commit 9459c62b577f · 2015-12-10T15:04:28.000-05:00
* This is similar to the bug in triangle fixed in 144a2db
diff --git a/src/backend/cuda/kernel/identity.hpp b/src/backend/cuda/kernel/identity.hpp
@@ -23,23 +23,26 @@ namespace kernel
     __global__
     static void identity_kernel(Param<T> out, int blocks_x, int blocks_y)
     {
-        unsigned idz = blockIdx.x / blocks_x;
-        unsigned idw = blockIdx.y / blocks_y;
+        const dim_t idz = blockIdx.x / blocks_x;
+        const dim_t idw = blockIdx.y / blocks_y;
 
-        unsigned blockIdx_x = blockIdx.x - idz * blocks_x;
-        unsigned blockIdx_y = blockIdx.y - idw * blocks_y;
+        const dim_t blockIdx_x = blockIdx.x - idz * blocks_x;
+        const dim_t blockIdx_y = blockIdx.y - idw * blocks_y;
 
-        unsigned idx = threadIdx.x + blockIdx_x * blockDim.x;
-        unsigned idy = threadIdx.y + blockIdx_y * blockDim.y;
+        const dim_t idx = threadIdx.x + blockIdx_x * blockDim.x;
+        const dim_t idy = threadIdx.y + blockIdx_y * blockDim.y;
 
         if(idx >= out.dims[0] ||
            idy >= out.dims[1] ||
            idz >= out.dims[2] ||
            idw >= out.dims[3])
             return;
 
+        const T one  = scalar<T>(1);
+        const T zero = scalar<T>(0);
+
         T *ptr = out.ptr + idz * out.strides[2] + idw * out.strides[3];
-        T val = (idx == idy) ? scalar<T>(1) : scalar<T>(0);
+        T val = (idx == idy) ? one : zero;
         ptr[idx + idy * out.strides[1]] = val;
     }
 
