accelerate-cuda 0.13.0.3 → 0.13.0.4
raw patch · 4 files changed
+23/−16 lines, 4 files
Files
- Data/Array/Accelerate/CUDA/CodeGen/Reduction.hs +13/−5
- Data/Array/Accelerate/CUDA/Compile.hs +0/−1
- accelerate-cuda.cabal +1/−1
- configure +9/−9
Data/Array/Accelerate/CUDA/CodeGen/Reduction.hs view
@@ -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
Data/Array/Accelerate/CUDA/Compile.hs view
@@ -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"
accelerate-cuda.cabal view
@@ -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
configure view
@@ -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\\"