diff --git a/Data/Array/Accelerate/CUDA/CodeGen/Reduction.hs b/Data/Array/Accelerate/CUDA/CodeGen/Reduction.hs
--- a/Data/Array/Accelerate/CUDA/CodeGen/Reduction.hs
+++ b/Data/Array/Accelerate/CUDA/CodeGen/Reduction.hs
@@ -167,7 +167,6 @@
          * cooperatively reduce the shared array to a single value.
          */
         $items:(sdata "threadIdx.x" .=. y)
-        __syncthreads();
 
         ix = min(shapeSize - blockIdx.x * blockDim.x, blockDim.x);
         $stms:(reduceBlock dev fun x y sdata (cvar "ix"))
@@ -317,8 +316,6 @@
              * cooperatively reduces this to a single value.
              */
             $items:(sdata "threadIdx.x" .=. y)
-            __syncthreads();
-
             $stms:(reduceBlock dev fun x y sdata (cvar "n"))
 
             /*
@@ -651,17 +648,28 @@
 
     reduce :: Int -> [C.Stm] -> [C.Stm]
     reduce i rest
+      -- Ensure that threads synchronise before reading from or writing to
+      -- shared memory. Synchronising after each reduction step is not enough,
+      -- because one warp could update the partial results before a different
+      -- warp has read in their data for this step.
+      --
       | i > warpSize dev
       = [cstm| if ( threadIdx.x + $int:i < $exp:n ) {
+                   __syncthreads();
                    $items:(x0 .=. sdata ("threadIdx.x + " ++ show i))
                    $items:(x1 .=. f x1 x0)
+                   __syncthreads();
                    $items:(sdata "threadIdx.x" .=. x1)
                } |]
-      : [cstm| __syncthreads(); |]
       : rest
 
+      -- The threads of a warp execute in lockstep, so it is only necessary to
+      -- synchronise at the top, to ensure all threads have written their
+      -- results into shared memory.
+      --
       | otherwise
-      = [cstm| if ( threadIdx.x < $int:(warpSize dev) ) {
+      = [cstm| __syncthreads(); |]
+      : [cstm| if ( threadIdx.x < $int:(warpSize dev) ) {
                    $stms:(reduceWarpTree dev fun x0 x1 sdata n (cvar "threadIdx.x"))
                } |]
       : rest
diff --git a/Data/Array/Accelerate/CUDA/Compile.hs b/Data/Array/Accelerate/CUDA/Compile.hs
--- a/Data/Array/Accelerate/CUDA/Compile.hs
+++ b/Data/Array/Accelerate/CUDA/Compile.hs
@@ -453,7 +453,6 @@
     [ "-I", ddir </> "cubits"
     , "-arch=sm_" ++ show m ++ show n
     , "-cubin"
-    , "--restrict"
     , "-o", cufile `replaceExtension` "cubin"
     , if D.mode D.dump_cc  then ""   else "--disable-warnings"
     , if D.mode D.debug_cc then "-G" else "-O3"
diff --git a/accelerate-cuda.cabal b/accelerate-cuda.cabal
--- a/accelerate-cuda.cabal
+++ b/accelerate-cuda.cabal
@@ -1,5 +1,5 @@
 Name:                   accelerate-cuda
-Version:                0.13.0.3
+Version:                0.13.0.4
 Cabal-version:          >= 1.6
 Tested-with:            GHC >= 7.4
 Build-type:             Custom
diff --git a/configure b/configure
--- a/configure
+++ b/configure
@@ -1,6 +1,6 @@
 #! /bin/sh
 # Guess values for system-dependent variables and create Makefiles.
-# Generated by GNU Autoconf 2.69 for accelerate-cuda 0.13.0.0.
+# Generated by GNU Autoconf 2.69 for accelerate-cuda 0.14.0.0.
 #
 # Report bugs to <accelerate-haskell@googlegroups.com>.
 #
@@ -579,8 +579,8 @@
 # Identity of this package.
 PACKAGE_NAME='accelerate-cuda'
 PACKAGE_TARNAME='accelerate-cuda'
-PACKAGE_VERSION='0.13.0.0'
-PACKAGE_STRING='accelerate-cuda 0.13.0.0'
+PACKAGE_VERSION='0.14.0.0'
+PACKAGE_STRING='accelerate-cuda 0.14.0.0'
 PACKAGE_BUGREPORT='accelerate-haskell@googlegroups.com'
 PACKAGE_URL=''
 
@@ -1192,7 +1192,7 @@
   # Omit some internal or obsolete options to make the list less imposing.
   # This message is too long to be a string in the A/UX 3.1 sh.
   cat <<_ACEOF
-\`configure' configures accelerate-cuda 0.13.0.0 to adapt to many kinds of systems.
+\`configure' configures accelerate-cuda 0.14.0.0 to adapt to many kinds of systems.
 
 Usage: $0 [OPTION]... [VAR=VALUE]...
 
@@ -1253,7 +1253,7 @@
 
 if test -n "$ac_init_help"; then
   case $ac_init_help in
-     short | recursive ) echo "Configuration of accelerate-cuda 0.13.0.0:";;
+     short | recursive ) echo "Configuration of accelerate-cuda 0.14.0.0:";;
    esac
   cat <<\_ACEOF
 
@@ -1339,7 +1339,7 @@
 test -n "$ac_init_help" && exit $ac_status
 if $ac_init_version; then
   cat <<\_ACEOF
-accelerate-cuda configure 0.13.0.0
+accelerate-cuda configure 0.14.0.0
 generated by GNU Autoconf 2.69
 
 Copyright (C) 2012 Free Software Foundation, Inc.
@@ -1394,7 +1394,7 @@
 This file contains any messages produced by compilers while
 running configure, to aid debugging if configure makes a mistake.
 
-It was created by accelerate-cuda $as_me 0.13.0.0, which was
+It was created by accelerate-cuda $as_me 0.14.0.0, which was
 generated by GNU Autoconf 2.69.  Invocation command line was
 
   $ $0 $@
@@ -2954,7 +2954,7 @@
 # report actual input values of CONFIG_FILES etc. instead of their
 # values after options handling.
 ac_log="
-This file was extended by accelerate-cuda $as_me 0.13.0.0, which was
+This file was extended by accelerate-cuda $as_me 0.14.0.0, which was
 generated by GNU Autoconf 2.69.  Invocation command line was
 
   CONFIG_FILES    = $CONFIG_FILES
@@ -3007,7 +3007,7 @@
 cat >>$CONFIG_STATUS <<_ACEOF || ac_write_fail=1
 ac_cs_config="`$as_echo "$ac_configure_args" | sed 's/^ //; s/[\\""\`\$]/\\\\&/g'`"
 ac_cs_version="\\
-accelerate-cuda config.status 0.13.0.0
+accelerate-cuda config.status 0.14.0.0
 configured by $0, generated by GNU Autoconf 2.69,
   with options \\"\$ac_cs_config\\"