[arrayfire] 17/84: Fix bug in identity cuda plaguing compute 5.2

[Ghislain Vaillant]

This is an automated email from the git hooks/post-receive script.

ghisvail-guest pushed a commit to branch master
in repository arrayfire.

commit 9459c62b577f5a102e3f53ec38a85d8c712f634c
Author: Shehzan Mohammed <shehzan at arrayfire.com>
Date:   Wed Dec 9 16:12:28 2015 -0500

    Fix bug in identity cuda plaguing compute 5.2
    
    * This is similar to the bug in triangle fixed in 144a2db
---
 src/backend/cuda/kernel/identity.hpp | 17 ++++++++++-------
 1 file changed, 10 insertions(+), 7 deletions(-)

diff --git a/src/backend/cuda/kernel/identity.hpp b/src/backend/cuda/kernel/identity.hpp
index ac97670..056838a 100644
--- a/src/backend/cuda/kernel/identity.hpp
+++ b/src/backend/cuda/kernel/identity.hpp
@@ -23,14 +23,14 @@ 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] ||
@@ -38,8 +38,11 @@ namespace kernel
            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;
     }
 

-- 
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/debian-science/packages/arrayfire.git