Update samples list to include additional samples.

2026-10-11 23:38:25 +08:00 · 2018-03-09 18:05:01 -08:00
parent 8bb8c5fac0
commit d08d485c67
83 changed files with 8530 additions and 7 deletions
--- a/Samples/shfl_scan/Makefile
+++ b/Samples/shfl_scan/Makefile
@@ -0,0 +1,304 @@
+################################################################################
+#
+# Copyright 1993-2015 NVIDIA Corporation.  All rights reserved.
+#
+# NOTICE TO USER:
+#
+# This source code is subject to NVIDIA ownership rights under U.S. and
+# international Copyright laws.
+#
+# NVIDIA MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THIS SOURCE
+# CODE FOR ANY PURPOSE.  IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR
+# IMPLIED WARRANTY OF ANY KIND.  NVIDIA DISCLAIMS ALL WARRANTIES WITH
+# REGARD TO THIS SOURCE CODE, INCLUDING ALL IMPLIED WARRANTIES OF
+# MERCHANTABILITY, NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE.
+# IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL,
+# OR CONSEQUENTIAL DAMAGES, OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS
+# OF USE, DATA OR PROFITS, WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE
+# OR OTHER TORTIOUS ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE
+# OR PERFORMANCE OF THIS SOURCE CODE.
+#
+# U.S. Government End Users.  This source code is a "commercial item" as
+# that term is defined at 48 C.F.R. 2.101 (OCT 1995), consisting  of
+# "commercial computer software" and "commercial computer software
+# documentation" as such terms are used in 48 C.F.R. 12.212 (SEPT 1995)
+# and is provided to the U.S. Government only as a commercial end item.
+# Consistent with 48 C.F.R.12.212 and 48 C.F.R. 227.7202-1 through
+# 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the
+# source code with only those rights set forth herein.
+#
+################################################################################
+#
+# Makefile project only supported on Mac OS X and Linux Platforms)
+#
+################################################################################
+
+# Location of the CUDA Toolkit
+CUDA_PATH ?= /usr/local/cuda
+
+##############################
+# start deprecated interface #
+##############################
+ifeq ($(x86_64),1)
+    $(info WARNING - x86_64 variable has been deprecated)
+    $(info WARNING - please use TARGET_ARCH=x86_64 instead)
+    TARGET_ARCH ?= x86_64
+endif
+ifeq ($(ARMv7),1)
+    $(info WARNING - ARMv7 variable has been deprecated)
+    $(info WARNING - please use TARGET_ARCH=armv7l instead)
+    TARGET_ARCH ?= armv7l
+endif
+ifeq ($(aarch64),1)
+    $(info WARNING - aarch64 variable has been deprecated)
+    $(info WARNING - please use TARGET_ARCH=aarch64 instead)
+    TARGET_ARCH ?= aarch64
+endif
+ifeq ($(ppc64le),1)
+    $(info WARNING - ppc64le variable has been deprecated)
+    $(info WARNING - please use TARGET_ARCH=ppc64le instead)
+    TARGET_ARCH ?= ppc64le
+endif
+ifneq ($(GCC),)
+    $(info WARNING - GCC variable has been deprecated)
+    $(info WARNING - please use HOST_COMPILER=$(GCC) instead)
+    HOST_COMPILER ?= $(GCC)
+endif
+ifneq ($(abi),)
+    $(error ERROR - abi variable has been removed)
+endif
+############################
+# end deprecated interface #
+############################
+
+# architecture
+HOST_ARCH   := $(shell uname -m)
+TARGET_ARCH ?= $(HOST_ARCH)
+ifneq (,$(filter $(TARGET_ARCH),x86_64 aarch64 ppc64le armv7l))
+    ifneq ($(TARGET_ARCH),$(HOST_ARCH))
+        ifneq (,$(filter $(TARGET_ARCH),x86_64 aarch64 ppc64le))
+            TARGET_SIZE := 64
+        else ifneq (,$(filter $(TARGET_ARCH),armv7l))
+            TARGET_SIZE := 32
+        endif
+    else
+        TARGET_SIZE := $(shell getconf LONG_BIT)
+    endif
+else
+    $(error ERROR - unsupported value $(TARGET_ARCH) for TARGET_ARCH!)
+endif
+ifneq ($(TARGET_ARCH),$(HOST_ARCH))
+    ifeq (,$(filter $(HOST_ARCH)-$(TARGET_ARCH),aarch64-armv7l x86_64-armv7l x86_64-aarch64 x86_64-ppc64le))
+        $(error ERROR - cross compiling from $(HOST_ARCH) to $(TARGET_ARCH) is not supported!)
+    endif
+endif
+
+# When on native aarch64 system with userspace of 32-bit, change TARGET_ARCH to armv7l
+ifeq ($(HOST_ARCH)-$(TARGET_ARCH)-$(TARGET_SIZE),aarch64-aarch64-32)
+    TARGET_ARCH = armv7l
+endif
+
+# operating system
+HOST_OS   := $(shell uname -s 2>/dev/null | tr "[:upper:]" "[:lower:]")
+TARGET_OS ?= $(HOST_OS)
+ifeq (,$(filter $(TARGET_OS),linux darwin qnx android))
+    $(error ERROR - unsupported value $(TARGET_OS) for TARGET_OS!)
+endif
+
+# host compiler
+ifeq ($(TARGET_OS),darwin)
+    ifeq ($(shell expr `xcodebuild -version | grep -i xcode | awk '{print $$2}' | cut -d'.' -f1` \>= 5),1)
+        HOST_COMPILER ?= clang++
+    endif
+else ifneq ($(TARGET_ARCH),$(HOST_ARCH))
+    ifeq ($(HOST_ARCH)-$(TARGET_ARCH),x86_64-armv7l)
+        ifeq ($(TARGET_OS),linux)
+            HOST_COMPILER ?= arm-linux-gnueabihf-g++
+        else ifeq ($(TARGET_OS),qnx)
+            ifeq ($(QNX_HOST),)
+                $(error ERROR - QNX_HOST must be passed to the QNX host toolchain)
+            endif
+            ifeq ($(QNX_TARGET),)
+                $(error ERROR - QNX_TARGET must be passed to the QNX target toolchain)
+            endif
+            export QNX_HOST
+            export QNX_TARGET
+            HOST_COMPILER ?= $(QNX_HOST)/usr/bin/arm-unknown-nto-qnx6.6.0eabi-g++
+        else ifeq ($(TARGET_OS),android)
+            HOST_COMPILER ?= arm-linux-androideabi-g++
+        endif
+    else ifeq ($(TARGET_ARCH),aarch64)
+        ifeq ($(TARGET_OS), linux)
+            HOST_COMPILER ?= aarch64-linux-gnu-g++
+        else ifeq ($(TARGET_OS),qnx)
+            ifeq ($(QNX_HOST),)
+                $(error ERROR - QNX_HOST must be passed to the QNX host toolchain)
+            endif
+            ifeq ($(QNX_TARGET),)
+                $(error ERROR - QNX_TARGET must be passed to the QNX target toolchain)
+            endif
+            export QNX_HOST
+            export QNX_TARGET
+            HOST_COMPILER ?= $(QNX_HOST)/usr/bin/aarch64-unknown-nto-qnx7.0.0-g++
+        else ifeq ($(TARGET_OS), android)
+            HOST_COMPILER ?= aarch64-linux-android-g++
+        endif
+    else ifeq ($(TARGET_ARCH),ppc64le)
+        HOST_COMPILER ?= powerpc64le-linux-gnu-g++
+    endif
+endif
+HOST_COMPILER ?= g++
+NVCC          := $(CUDA_PATH)/bin/nvcc -ccbin $(HOST_COMPILER)
+
+# internal flags
+NVCCFLAGS   := -m${TARGET_SIZE}
+CCFLAGS     :=
+LDFLAGS     :=
+
+# build flags
+ifeq ($(TARGET_OS),darwin)
+    LDFLAGS += -rpath $(CUDA_PATH)/lib
+    CCFLAGS += -arch $(HOST_ARCH)
+else ifeq ($(HOST_ARCH)-$(TARGET_ARCH)-$(TARGET_OS),x86_64-armv7l-linux)
+    LDFLAGS += --dynamic-linker=/lib/ld-linux-armhf.so.3
+    CCFLAGS += -mfloat-abi=hard
+else ifeq ($(TARGET_OS),android)
+    LDFLAGS += -pie
+    CCFLAGS += -fpie -fpic -fexceptions
+endif
+
+ifneq ($(TARGET_ARCH),$(HOST_ARCH))
+    ifeq ($(TARGET_ARCH)-$(TARGET_OS),armv7l-linux)
+        ifneq ($(TARGET_FS),)
+            GCCVERSIONLTEQ46 := $(shell expr `$(HOST_COMPILER) -dumpversion` \<= 4.6)
+            ifeq ($(GCCVERSIONLTEQ46),1)
+                CCFLAGS += --sysroot=$(TARGET_FS)
+            endif
+            LDFLAGS += --sysroot=$(TARGET_FS)
+            LDFLAGS += -rpath-link=$(TARGET_FS)/lib
+            LDFLAGS += -rpath-link=$(TARGET_FS)/usr/lib
+            LDFLAGS += -rpath-link=$(TARGET_FS)/usr/lib/arm-linux-gnueabihf
+        endif
+    endif
+    ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-linux)
+        ifneq ($(TARGET_FS),)
+            GCCVERSIONLTEQ46 := $(shell expr `$(HOST_COMPILER) -dumpversion` \<= 4.6)
+            ifeq ($(GCCVERSIONLTEQ46),1)
+                CCFLAGS += --sysroot=$(TARGET_FS)
+            endif
+            LDFLAGS += --sysroot=$(TARGET_FS)
+            LDFLAGS += -rpath-link=$(TARGET_FS)/lib -L $(TARGET_FS)/lib
+            LDFLAGS += -rpath-link=$(TARGET_FS)/usr/lib -L $(TARGET_FS)/usr/lib
+            LDFLAGS += -rpath-link=$(TARGET_FS)/usr/lib/aarch64-linux-gnu -L $(TARGET_FS)/usr/lib/aarch64-linux-gnu
+            LDFLAGS += --unresolved-symbols=ignore-in-shared-libs
+            CCFLAGS += -isystem=$(TARGET_FS)/usr/include
+            CCFLAGS += -isystem=$(TARGET_FS)/usr/include/aarch64-linux-gnu
+        endif
+    endif
+endif
+
+ifeq ($(TARGET_OS),qnx)
+    CCFLAGS += -DWIN_INTERFACE_CUSTOM
+    LDFLAGS += -lsocket
+endif
+
+# Install directory of different arch
+CUDA_INSTALL_TARGET_DIR :=
+ifeq ($(TARGET_ARCH)-$(TARGET_OS),armv7l-linux)
+    CUDA_INSTALL_TARGET_DIR = targets/armv7-linux-gnueabihf/
+else ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-linux)
+    CUDA_INSTALL_TARGET_DIR = targets/aarch64-linux/
+else ifeq ($(TARGET_ARCH)-$(TARGET_OS),armv7l-android)
+    CUDA_INSTALL_TARGET_DIR = targets/armv7-linux-androideabi/
+else ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-android)
+    CUDA_INSTALL_TARGET_DIR = targets/aarch64-linux-androideabi/
+else ifeq ($(TARGET_ARCH)-$(TARGET_OS),armv7l-qnx)
+    CUDA_INSTALL_TARGET_DIR = targets/ARMv7-linux-QNX/
+else ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-qnx)
+    CUDA_INSTALL_TARGET_DIR = targets/aarch64-qnx/
+else ifeq ($(TARGET_ARCH),ppc64le)
+    CUDA_INSTALL_TARGET_DIR = targets/ppc64le-linux/
+endif
+
+# Debug build flags
+ifeq ($(dbg),1)
+      NVCCFLAGS += -g -G
+      BUILD_TYPE := debug
+else
+      BUILD_TYPE := release
+endif
+
+ALL_CCFLAGS :=
+ALL_CCFLAGS += $(NVCCFLAGS)
+ALL_CCFLAGS += $(EXTRA_NVCCFLAGS)
+ALL_CCFLAGS += $(addprefix -Xcompiler ,$(CCFLAGS))
+ALL_CCFLAGS += $(addprefix -Xcompiler ,$(EXTRA_CCFLAGS))
+
+SAMPLE_ENABLED := 1
+
+ALL_LDFLAGS :=
+ALL_LDFLAGS += $(ALL_CCFLAGS)
+ALL_LDFLAGS += $(addprefix -Xlinker ,$(LDFLAGS))
+ALL_LDFLAGS += $(addprefix -Xlinker ,$(EXTRA_LDFLAGS))
+
+# Common includes and paths for CUDA
+INCLUDES  := -I../../Common
+LIBRARIES :=
+
+################################################################################
+
+# Gencode arguments
+SMS ?= 30 35 37 50 52 60 61 70
+
+ifeq ($(SMS),)
+$(info >>> WARNING - no SM architectures have been specified - waiving sample <<<)
+SAMPLE_ENABLED := 0
+endif
+
+ifeq ($(GENCODE_FLAGS),)
+# Generate SASS code for each SM architecture listed in $(SMS)
+$(foreach sm,$(SMS),$(eval GENCODE_FLAGS += -gencode arch=compute_$(sm),code=sm_$(sm)))
+
+# Generate PTX code from the highest SM architecture in $(SMS) to guarantee forward-compatibility
+HIGHEST_SM := $(lastword $(sort $(SMS)))
+ifneq ($(HIGHEST_SM),)
+GENCODE_FLAGS += -gencode arch=compute_$(HIGHEST_SM),code=compute_$(HIGHEST_SM)
+endif
+endif
+
+ALL_CCFLAGS += -O3
+
+ifeq ($(SAMPLE_ENABLED),0)
+EXEC ?= @echo "[@]"
+endif
+
+################################################################################
+
+# Target rules
+all: build
+
+build: shfl_scan
+
+check.deps:
+ifeq ($(SAMPLE_ENABLED),0)
+	@echo "Sample will be waived due to the above missing dependencies"
+else
+	@echo "Sample is ready - all dependencies have been met"
+endif
+
+shfl_scan.o:shfl_scan.cu
+	$(EXEC) $(NVCC) $(INCLUDES) $(ALL_CCFLAGS) $(GENCODE_FLAGS) -o $@ -c $<
+
+shfl_scan: shfl_scan.o
+	$(EXEC) $(NVCC) $(ALL_LDFLAGS) $(GENCODE_FLAGS) -o $@ $+ $(LIBRARIES)
+	$(EXEC) mkdir -p ../../bin/$(TARGET_ARCH)/$(TARGET_OS)/$(BUILD_TYPE)
+	$(EXEC) cp $@ ../../bin/$(TARGET_ARCH)/$(TARGET_OS)/$(BUILD_TYPE)
+
+run: build
+	$(EXEC) ./shfl_scan
+
+clean:
+	rm -f shfl_scan shfl_scan.o
+	rm -rf ../../bin/$(TARGET_ARCH)/$(TARGET_OS)/$(BUILD_TYPE)/shfl_scan
+
+clobber: clean
--- a/Samples/shfl_scan/NsightEclipse.xml
+++ b/Samples/shfl_scan/NsightEclipse.xml
@@ -0,0 +1,73 @@
+<?xml version="1.0" encoding="UTF-8"?> 
+<!DOCTYPE entry SYSTEM "SamplesInfo.dtd">
+<entry>
+  <name>shfl_scan</name>
+  <cflags>
+    <flag>-O3</flag>
+  </cflags>
+  <description><![CDATA[This example demonstrates how to use the shuffle intrinsic __shfl_up_sync to perform a scan operation across a thread block. ]]></description>
+  <devicecompilation>whole</devicecompilation>
+  <includepaths>
+    <path>./</path>
+    <path>../</path>
+    <path>../../common/inc</path>
+  </includepaths>
+  <keyconcepts>
+    <concept level="advanced">Data-Parallel Algorithms</concept>
+    <concept level="advanced">Performance Strategies</concept>
+  </keyconcepts>
+  <keywords>
+    <keyword>GPGPU</keyword>
+    <keyword>CUDA</keyword>
+    <keyword>scan</keyword>
+    <keyword>parallel prefix sum</keyword>
+    <keyword>Data-Parallel Algorithms</keyword>
+  </keywords>
+  <libraries>
+  </libraries>
+  <librarypaths>
+  </librarypaths>
+  <nsight_eclipse>true</nsight_eclipse>
+  <primary_file>shfl_scan.cu</primary_file>
+  <scopes>
+    <scope>1:CUDA Advanced Topics</scope>
+    <scope>1:Data-Parallel Algorithms</scope>
+    <scope>1:Performance Strategies</scope>
+  </scopes>
+  <sm-arch>sm30</sm-arch>
+  <sm-arch>sm35</sm-arch>
+  <sm-arch>sm37</sm-arch>
+  <sm-arch>sm50</sm-arch>
+  <sm-arch>sm52</sm-arch>
+  <sm-arch>sm60</sm-arch>
+  <sm-arch>sm61</sm-arch>
+  <sm-arch>sm70</sm-arch>
+  <supported_envs>
+    <env>
+      <arch>x86_64</arch>
+      <platform>linux</platform>
+    </env>
+    <env>
+      <platform>windows7</platform>
+    </env>
+    <env>
+      <arch>x86_64</arch>
+      <platform>macosx</platform>
+    </env>
+    <env>
+      <arch>arm</arch>
+    </env>
+    <env>
+      <arch>aarch64</arch>
+    </env>
+    <env>
+      <arch>ppc64le</arch>
+      <platform>linux</platform>
+    </env>
+  </supported_envs>
+  <supported_sm_architectures>
+    <from>3.0</from>
+  </supported_sm_architectures>
+  <title>CUDA Parallel Prefix Sum with Shuffle Intrinsics (SHFL_Scan)</title>
+  <type>exe</type>
+</entry>
--- a/Samples/shfl_scan/README.md
+++ b/Samples/shfl_scan/README.md
@@ -0,0 +1,91 @@
+# shfl_scan - CUDA Parallel Prefix Sum with Shuffle Intrinsics (SHFL_Scan)
+
+## Description
+
+This example demonstrates how to use the shuffle intrinsic __shfl_up_sync to perform a scan operation across a thread block.
+
+## Key Concepts
+
+Data-Parallel Algorithms, Performance Strategies
+
+## Supported SM Architectures
+
+[SM 3.0 ](https://developer.nvidia.com/cuda-gpus)  [SM 3.5 ](https://developer.nvidia.com/cuda-gpus)  [SM 3.7 ](https://developer.nvidia.com/cuda-gpus)  [SM 5.0 ](https://developer.nvidia.com/cuda-gpus)  [SM 5.2 ](https://developer.nvidia.com/cuda-gpus)  [SM 6.0 ](https://developer.nvidia.com/cuda-gpus)  [SM 6.1 ](https://developer.nvidia.com/cuda-gpus)  [SM 7.0 ](https://developer.nvidia.com/cuda-gpus)
+
+## Supported OSes
+
+Linux, Windows, MacOSX
+
+## Supported CPU Architecture
+
+x86_64, ppc64le, armv7l, aarch64
+
+## CUDA APIs involved
+
+## Prerequisites
+
+Download and install the [CUDA Toolkit 9.2](https://developer.nvidia.com/cuda-downloads) for your corresponding platform.
+
+## Build and Run
+
+### Windows
+The Windows samples are built using the Visual Studio IDE. Solution files (.sln) are provided for each supported version of Visual Studio, using the format:
+```
+*_vs<version>.sln - for Visual Studio <version>
+```
+Each individual sample has its own set of solution files in its directory:
+
+To build/examine all the samples at once, the complete solution files should be used. To build/examine a single sample, the individual sample solution files should be used.
+> **Note:** Some samples require that the Microsoft DirectX SDK (June 2010 or newer) be installed and that the VC++ directory paths are properly set up (**Tools > Options...**). Check DirectX Dependencies section for details."
+
+### Linux
+The Linux samples are built using makefiles. To use the makefiles, change the current directory to the sample directory you wish to build, and run make:
+```
+$ cd <sample_dir>
+$ make
+```
+The samples makefiles can take advantage of certain options:
+*  **TARGET_ARCH=<arch>** - cross-compile targeting a specific architecture. Allowed architectures are x86_64, ppc64le, armv7l, aarch64.
+    By default, TARGET_ARCH is set to HOST_ARCH. On a x86_64 machine, not setting TARGET_ARCH is the equivalent of setting TARGET_ARCH=x86_64.<br/>
+`$ make TARGET_ARCH=x86_64` <br/> `$ make TARGET_ARCH=ppc64le` <br/> `$ make TARGET_ARCH=armv7l` <br/> `$ make TARGET_ARCH=aarch64` <br/>
+    See [here](http://docs.nvidia.com/cuda/cuda-samples/index.html#cross-samples) for more details.
+*   **dbg=1** - build with debug symbols
+    ```
+    $ make dbg=1
+    ```
+*   **SMS="A B ..."** - override the SM architectures for which the sample will be built, where `"A B ..."` is a space-delimited list of SM architectures. For example, to generate SASS for SM 50 and SM 60, use `SMS="50 60"`.
+    ```
+    $ make SMS="50 60"
+    ```
+
+*  **HOST_COMPILER=<host_compiler>** - override the default g++ host compiler. See the [Linux Installation Guide](http://docs.nvidia.com/cuda/cuda-installation-guide-linux/index.html#system-requirements) for a list of supported host compilers.
+```
+    $ make HOST_COMPILER=g++
+```
+
+### Mac
+The Mac samples are built using makefiles. To use the makefiles, change directory into the sample directory you wish to build, and run make:
+```
+$ cd <sample_dir>
+$ make
+```
+
+The samples makefiles can take advantage of certain options:
+
+*  **dbg=1** - build with debug symbols
+    ```
+    $ make dbg=1
+    ```
+
+*  **SMS="A B ..."** - override the SM architectures for which the sample will be built, where "A B ..." is a space-delimited list of SM architectures. For example, to generate SASS for SM 50 and SM 60, use SMS="50 60".
+    ```
+    $ make SMS="A B ..."
+    ```
+
+*  **HOST_COMPILER=<host_compiler>** - override the default clang host compiler. See the [Mac Installation Guide](http://docs.nvidia.com/cuda/cuda-installation-guide-mac-os-x/index.html#system-requirements) for a list of supported host compilers.
+    ```
+    $ make HOST_COMPILER=clang
+    ```
+
+## References (for more details)
+
--- a/Samples/shfl_scan/shfl_integral_image.cuh
+++ b/Samples/shfl_scan/shfl_integral_image.cuh
@@ -0,0 +1,359 @@
+/* Copyright (c) 2018, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions
+ * are met:
+ *  * Redistributions of source code must retain the above copyright
+ *    notice, this list of conditions and the following disclaimer.
+ *  * Redistributions in binary form must reproduce the above copyright
+ *    notice, this list of conditions and the following disclaimer in the
+ *    documentation and/or other materials provided with the distribution.
+ *  * Neither the name of NVIDIA CORPORATION nor the names of its
+ *    contributors may be used to endorse or promote products derived
+ *    from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
+ * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
+ * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+// Utility function to extract unsigned chars from an
+// unsigned integer
+
+__device__ uchar4 int_to_uchar4(unsigned int in) {
+  uchar4 bytes;
+  bytes.x = in & 0x000000ff >> 0;
+  bytes.y = in & 0x0000ff00 >> 8;
+  bytes.z = in & 0x00ff0000 >> 16;
+  bytes.w = in & 0xff000000 >> 24;
+  return bytes;
+}
+
+// This function demonstrates some uses of the shuffle instruction
+// in the generation of an integral image (also
+// called a summed area table)
+// The approach is two pass, a horizontal (scanline) then a vertical
+// (column) pass.
+// This is the horizontal pass kernel.
+__global__ void shfl_intimage_rows(uint4 *img, uint4 *integral_image) {
+  __shared__ int sums[128];
+
+  int id = threadIdx.x;
+  // pointer to head of current scanline
+  uint4 *scanline = &img[blockIdx.x * 120];
+  uint4 data;
+  data = scanline[id];
+  int result[16];
+  int sum;
+  unsigned int lane_id = id % warpSize;
+  int warp_id = threadIdx.x / warpSize;
+
+  uchar4 a = int_to_uchar4(data.x);
+  uchar4 b = int_to_uchar4(data.y);
+  uchar4 c = int_to_uchar4(data.z);
+  uchar4 d = int_to_uchar4(data.w);
+
+  result[0] = a.x;
+  result[1] = a.x + a.y;
+  result[2] = a.x + a.y + a.z;
+  result[3] = a.x + a.y + a.z + a.w;
+
+  result[4] = b.x;
+  result[5] = b.x + b.y;
+  result[6] = b.x + b.y + b.z;
+  result[7] = b.x + b.y + b.z + b.w;
+
+  result[8] = c.x;
+  result[9] = c.x + c.y;
+  result[10] = c.x + c.y + c.z;
+  result[11] = c.x + c.y + c.z + c.w;
+
+  result[12] = d.x;
+  result[13] = d.x + d.y;
+  result[14] = d.x + d.y + d.z;
+  result[15] = d.x + d.y + d.z + d.w;
+
+#pragma unroll
+
+  for (int i = 4; i <= 7; i++) result[i] += result[3];
+
+#pragma unroll
+
+  for (int i = 8; i <= 11; i++) result[i] += result[7];
+
+#pragma unroll
+
+  for (int i = 12; i <= 15; i++) result[i] += result[11];
+
+  sum = result[15];
+
+  // the prefix sum for each thread's 16 value is computed,
+  // now the final sums (result[15]) need to be shared
+  // with the other threads and add.  To do this,
+  // the __shfl_up() instruction is used and a shuffle scan
+  // operation is performed to distribute the sums to the correct
+  // threads
+#pragma unroll
+
+  for (int i = 1; i < 32; i *= 2) {
+    unsigned int mask = 0xffffffff;
+    int n = __shfl_up_sync(mask, sum, i, 32);
+
+    if (lane_id >= i) {
+#pragma unroll
+
+      for (int i = 0; i < 16; i++) {
+        result[i] += n;
+      }
+
+      sum += n;
+    }
+  }
+
+  // Now the final sum for the warp must be shared
+  // between warps.  This is done by each warp
+  // having a thread store to shared memory, then
+  // having some other warp load the values and
+  // compute a prefix sum, again by using __shfl_up.
+  // The results are uniformly added back to the warps.
+  // last thread in the warp holding sum of the warp
+  // places that in shared
+  if (threadIdx.x % warpSize == warpSize - 1) {
+    sums[warp_id] = result[15];
+  }
+
+  __syncthreads();
+
+  if (warp_id == 0) {
+    int warp_sum = sums[lane_id];
+#pragma unroll
+
+    for (int i = 1; i <= 32; i *= 2) {
+      unsigned int mask = 0xffffffff;
+      int n = __shfl_up_sync(mask, warp_sum, i, 32);
+
+      if (lane_id >= i) warp_sum += n;
+    }
+
+    sums[lane_id] = warp_sum;
+  }
+
+  __syncthreads();
+
+  int blockSum = 0;
+
+  // fold in unused warp
+  if (warp_id > 0) {
+    blockSum = sums[warp_id - 1];
+#pragma unroll
+
+    for (int i = 0; i < 16; i++) {
+      result[i] += blockSum;
+    }
+  }
+
+  // assemble result
+  // Each thread has 16 values to write, which are
+  // now integer data (to avoid overflow).  Instead of
+  // each thread writing consecutive uint4s, the
+  // approach shown here experiments using
+  // the shuffle command to reformat the data
+  // inside the registers so that each thread holds
+  // consecutive data to be written so larger contiguous
+  // segments can be assembled for writing.
+  /*
+      For example data that needs to be written as
+
+      GMEM[16] <- x0 x1 x2 x3 y0 y1 y2 y3 z0 z1 z2 z3 w0 w1 w2 w3
+      but is stored in registers (r0..r3), in four threads (0..3) as:
+
+      threadId   0  1  2  3
+        r0      x0 y0 z0 w0
+        r1      x1 y1 z1 w1
+        r2      x2 y2 z2 w2
+        r3      x3 y3 z3 w3
+
+        after apply __shfl_xor operations to move data between registers r1..r3:
+
+      threadId  00 01 10 11
+                x0 y0 z0 w0
+       xor(01)->y1 x1 w1 z1
+       xor(10)->z2 w2 x2 y2
+       xor(11)->w3 z3 y3 x3
+
+       and now x0..x3, and z0..z3 can be written out in order by all threads.
+
+       In the current code, each register above is actually representing
+       four integers to be written as uint4's to GMEM.
+  */
+
+  unsigned int mask = 0xffffffff;
+  uint4 output;
+  result[4] = __shfl_xor_sync(mask, result[4], 1, 32);
+  result[5] = __shfl_xor_sync(mask, result[5], 1, 32);
+  result[6] = __shfl_xor_sync(mask, result[6], 1, 32);
+  result[7] = __shfl_xor_sync(mask, result[7], 1, 32);
+
+  result[8] = __shfl_xor_sync(mask, result[8], 2, 32);
+  result[9] = __shfl_xor_sync(mask, result[9], 2, 32);
+  result[10] = __shfl_xor_sync(mask, result[10], 2, 32);
+  result[11] = __shfl_xor_sync(mask, result[11], 2, 32);
+
+  result[12] = __shfl_xor_sync(mask, result[12], 3, 32);
+  result[13] = __shfl_xor_sync(mask, result[13], 3, 32);
+  result[14] = __shfl_xor_sync(mask, result[14], 3, 32);
+  result[15] = __shfl_xor_sync(mask, result[15], 3, 32);
+
+  if (threadIdx.x % 4 == 0) {
+    output = make_uint4(result[0], result[1], result[2], result[3]);
+  }
+
+  if (threadIdx.x % 4 == 1) {
+    output = make_uint4(result[4], result[5], result[6], result[7]);
+  }
+
+  if (threadIdx.x % 4 == 2) {
+    output = make_uint4(result[8], result[9], result[10], result[11]);
+  }
+
+  if (threadIdx.x % 4 == 3) {
+    output = make_uint4(result[12], result[13], result[14], result[15]);
+  }
+
+  integral_image[blockIdx.x * 480 + threadIdx.x % 4 + (threadIdx.x / 4) * 16] =
+      output;
+
+  if (threadIdx.x % 4 == 2) {
+    output = make_uint4(result[0], result[1], result[2], result[3]);
+  }
+
+  if (threadIdx.x % 4 == 3) {
+    output = make_uint4(result[4], result[5], result[6], result[7]);
+  }
+
+  if (threadIdx.x % 4 == 0) {
+    output = make_uint4(result[8], result[9], result[10], result[11]);
+  }
+
+  if (threadIdx.x % 4 == 1) {
+    output = make_uint4(result[12], result[13], result[14], result[15]);
+  }
+
+  integral_image[blockIdx.x * 480 + (threadIdx.x + 2) % 4 +
+                 (threadIdx.x / 4) * 16 + 8] = output;
+  // continuing from the above example,
+  // this use of __shfl_xor() places the y0..y3 and w0..w3 data
+  // in order.
+#pragma unroll
+
+  for (int i = 0; i < 16; i++) {
+    result[i] = __shfl_xor_sync(mask, result[i], 1, 32);
+  }
+
+  if (threadIdx.x % 4 == 0) {
+    output = make_uint4(result[0], result[1], result[2], result[3]);
+  }
+
+  if (threadIdx.x % 4 == 1) {
+    output = make_uint4(result[4], result[5], result[6], result[7]);
+  }
+
+  if (threadIdx.x % 4 == 2) {
+    output = make_uint4(result[8], result[9], result[10], result[11]);
+  }
+
+  if (threadIdx.x % 4 == 3) {
+    output = make_uint4(result[12], result[13], result[14], result[15]);
+  }
+
+  integral_image[blockIdx.x * 480 + threadIdx.x % 4 + (threadIdx.x / 4) * 16 +
+                 4] = output;
+
+  if (threadIdx.x % 4 == 2) {
+    output = make_uint4(result[0], result[1], result[2], result[3]);
+  }
+
+  if (threadIdx.x % 4 == 3) {
+    output = make_uint4(result[4], result[5], result[6], result[7]);
+  }
+
+  if (threadIdx.x % 4 == 0) {
+    output = make_uint4(result[8], result[9], result[10], result[11]);
+  }
+
+  if (threadIdx.x % 4 == 1) {
+    output = make_uint4(result[12], result[13], result[14], result[15]);
+  }
+
+  integral_image[blockIdx.x * 480 + (threadIdx.x + 2) % 4 +
+                 (threadIdx.x / 4) * 16 + 12] = output;
+}
+
+// This kernel computes columnwise prefix sums.  When the data input is
+// the row sums from above, this completes the integral image.
+// The approach here is to have each block compute a local set of sums.
+// First , the data covered by the block is loaded into shared memory,
+// then instead of performing a sum in shared memory using __syncthreads
+// between stages, the data is reformatted so that the necessary sums
+// occur inside warps and the shuffle scan operation is used.
+// The final set of sums from the block is then propagated, with the block
+// computing "down" the image and adding the running sum to the local
+// block sums.
+__global__ void shfl_vertical_shfl(unsigned int *img, int width, int height) {
+  __shared__ unsigned int sums[32][9];
+  int tidx = blockIdx.x * blockDim.x + threadIdx.x;
+  // int warp_id = threadIdx.x / warpSize ;
+  unsigned int lane_id = tidx % 8;
+  // int rows_per_thread = (height / blockDim. y) ;
+  // int start_row = rows_per_thread * threadIdx.y;
+  unsigned int stepSum = 0;
+  unsigned int mask = 0xffffffff;
+
+  sums[threadIdx.x][threadIdx.y] = 0;
+  __syncthreads();
+
+  for (int step = 0; step < 135; step++) {
+    unsigned int sum = 0;
+    unsigned int *p = img + (threadIdx.y + step * 8) * width + tidx;
+
+    sum = *p;
+    sums[threadIdx.x][threadIdx.y] = sum;
+    __syncthreads();
+
+    // place into SMEM
+    // shfl scan reduce the SMEM, reformating so the column
+    // sums are computed in a warp
+    // then read out properly
+    int partial_sum = 0;
+    int j = threadIdx.x % 8;
+    int k = threadIdx.x / 8 + threadIdx.y * 4;
+
+    partial_sum = sums[k][j];
+
+    for (int i = 1; i <= 8; i *= 2) {
+      int n = __shfl_up_sync(mask, partial_sum, i, 32);
+
+      if (lane_id >= i) partial_sum += n;
+    }
+
+    sums[k][j] = partial_sum;
+    __syncthreads();
+
+    if (threadIdx.y > 0) {
+      sum += sums[threadIdx.x][threadIdx.y - 1];
+    }
+
+    sum += stepSum;
+    stepSum += sums[threadIdx.x][blockDim.y - 1];
+    __syncthreads();
+    *p = sum;
+  }
+}
--- a/Samples/shfl_scan/shfl_scan.cu
+++ b/Samples/shfl_scan/shfl_scan.cu
@@ -0,0 +1,418 @@
+/* Copyright (c) 2018, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions
+ * are met:
+ *  * Redistributions of source code must retain the above copyright
+ *    notice, this list of conditions and the following disclaimer.
+ *  * Redistributions in binary form must reproduce the above copyright
+ *    notice, this list of conditions and the following disclaimer in the
+ *    documentation and/or other materials provided with the distribution.
+ *  * Neither the name of NVIDIA CORPORATION nor the names of its
+ *    contributors may be used to endorse or promote products derived
+ *    from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
+ * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
+ * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+// Shuffle intrinsics CUDA Sample
+// This sample demonstrates the use of the shuffle intrinsic
+// First, a simple example of a prefix sum using the shuffle to
+// perform a scan operation is provided.
+// Secondly, a more involved example of computing an integral image
+// using the shuffle intrinsic is provided, where the shuffle
+// scan operation and shuffle xor operations are used
+
+#include <stdio.h>
+
+#include <cuda_runtime.h>
+
+#include <helper_cuda.h>
+#include <helper_functions.h>
+#include "shfl_integral_image.cuh"
+
+// Scan using shfl - takes log2(n) steps
+// This function demonstrates basic use of the shuffle intrinsic, __shfl_up,
+// to perform a scan operation across a block.
+// First, it performs a scan (prefix sum in this case) inside a warp
+// Then to continue the scan operation across the block,
+// each warp's sum is placed into shared memory.  A single warp
+// then performs a shuffle scan on that shared memory.  The results
+// are then uniformly added to each warp's threads.
+// This pyramid type approach is continued by placing each block's
+// final sum in global memory and prefix summing that via another kernel call,
+// then uniformly adding across the input data via the uniform_add<<<>>> kernel.
+
+__global__ void shfl_scan_test(int *data, int width, int *partial_sums = NULL) {
+  extern __shared__ int sums[];
+  int id = ((blockIdx.x * blockDim.x) + threadIdx.x);
+  int lane_id = id % warpSize;
+  // determine a warp_id within a block
+  int warp_id = threadIdx.x / warpSize;
+
+  // Below is the basic structure of using a shfl instruction
+  // for a scan.
+  // Record "value" as a variable - we accumulate it along the way
+  int value = data[id];
+
+  // Now accumulate in log steps up the chain
+  // compute sums, with another thread's value who is
+  // distance delta away (i).  Note
+  // those threads where the thread 'i' away would have
+  // been out of bounds of the warp are unaffected.  This
+  // creates the scan sum.
+
+#pragma unroll
+  for (int i = 1; i <= width; i *= 2) {
+    unsigned int mask = 0xffffffff;
+    int n = __shfl_up_sync(mask, value, i, width);
+
+    if (lane_id >= i) value += n;
+  }
+
+  // value now holds the scan value for the individual thread
+  // next sum the largest values for each warp
+
+  // write the sum of the warp to smem
+  if (threadIdx.x % warpSize == warpSize - 1) {
+    sums[warp_id] = value;
+  }
+
+  __syncthreads();
+
+  //
+  // scan sum the warp sums
+  // the same shfl scan operation, but performed on warp sums
+  //
+  if (warp_id == 0 && lane_id < (blockDim.x / warpSize)) {
+    int warp_sum = sums[lane_id];
+
+    int mask = (1 << (blockDim.x / warpSize)) - 1;
+    for (int i = 1; i <= (blockDim.x / warpSize); i *= 2) {
+      int n = __shfl_up_sync(mask, warp_sum, i, (blockDim.x / warpSize));
+
+      if (lane_id >= i) warp_sum += n;
+    }
+
+    sums[lane_id] = warp_sum;
+  }
+
+  __syncthreads();
+
+  // perform a uniform add across warps in the block
+  // read neighbouring warp's sum and add it to threads value
+  int blockSum = 0;
+
+  if (warp_id > 0) {
+    blockSum = sums[warp_id - 1];
+  }
+
+  value += blockSum;
+
+  // Now write out our result
+  data[id] = value;
+
+  // last thread has sum, write write out the block's sum
+  if (partial_sums != NULL && threadIdx.x == blockDim.x - 1) {
+    partial_sums[blockIdx.x] = value;
+  }
+}
+
+// Uniform add: add partial sums array
+__global__ void uniform_add(int *data, int *partial_sums, int len) {
+  __shared__ int buf;
+  int id = ((blockIdx.x * blockDim.x) + threadIdx.x);
+
+  if (id > len) return;
+
+  if (threadIdx.x == 0) {
+    buf = partial_sums[blockIdx.x];
+  }
+
+  __syncthreads();
+  data[id] += buf;
+}
+
+static unsigned int iDivUp(unsigned int dividend, unsigned int divisor) {
+  return ((dividend % divisor) == 0) ? (dividend / divisor)
+                                     : (dividend / divisor + 1);
+}
+
+// This function verifies the shuffle scan result, for the simple
+// prefix sum case.
+bool CPUverify(int *h_data, int *h_result, int n_elements) {
+  // cpu verify
+  for (int i = 0; i < n_elements - 1; i++) {
+    h_data[i + 1] = h_data[i] + h_data[i + 1];
+  }
+
+  int diff = 0;
+
+  for (int i = 0; i < n_elements; i++) {
+    diff += h_data[i] - h_result[i];
+  }
+
+  printf("CPU verify result diff (GPUvsCPU) = %d\n", diff);
+  bool bTestResult = false;
+
+  if (diff == 0) bTestResult = true;
+
+  StopWatchInterface *hTimer = NULL;
+  sdkCreateTimer(&hTimer);
+  sdkResetTimer(&hTimer);
+  sdkStartTimer(&hTimer);
+
+  for (int j = 0; j < 100; j++)
+    for (int i = 0; i < n_elements - 1; i++) {
+      h_data[i + 1] = h_data[i] + h_data[i + 1];
+    }
+
+  sdkStopTimer(&hTimer);
+  double cput = sdkGetTimerValue(&hTimer);
+  printf("CPU sum (naive) took %f ms\n", cput / 100);
+  return bTestResult;
+}
+
+// this verifies the row scan result for synthetic data of all 1's
+unsigned int verifyDataRowSums(unsigned int *h_image, int w, int h) {
+  unsigned int diff = 0;
+
+  for (int j = 0; j < h; j++) {
+    for (int i = 0; i < w; i++) {
+      int gold = i + 1;
+      diff +=
+          abs(static_cast<int>(gold) - static_cast<int>(h_image[j * w + i]));
+    }
+  }
+
+  return diff;
+}
+
+bool shuffle_simple_test(int argc, char **argv) {
+  int *h_data, *h_partial_sums, *h_result;
+  int *d_data, *d_partial_sums;
+  const int n_elements = 65536;
+  int sz = sizeof(int) * n_elements;
+  int cuda_device = 0;
+
+  printf("Starting shfl_scan\n");
+
+  // use command-line specified CUDA device, otherwise use device with highest
+  // Gflops/s
+  cuda_device = findCudaDevice(argc, (const char **)argv);
+
+  cudaDeviceProp deviceProp;
+  checkCudaErrors(cudaGetDevice(&cuda_device));
+
+  checkCudaErrors(cudaGetDeviceProperties(&deviceProp, cuda_device));
+
+  printf("> Detected Compute SM %d.%d hardware with %d multi-processors\n",
+         deviceProp.major, deviceProp.minor, deviceProp.multiProcessorCount);
+
+  // __shfl intrinsic needs SM 3.0 or higher
+  if (deviceProp.major < 3) {
+    printf("> __shfl() intrinsic requires device SM 3.0+\n");
+    printf("> Waiving test.\n");
+    exit(EXIT_WAIVED);
+  }
+
+  checkCudaErrors(cudaMallocHost(reinterpret_cast<void **>(&h_data),
+                                 sizeof(int) * n_elements));
+  checkCudaErrors(cudaMallocHost(reinterpret_cast<void **>(&h_result),
+                                 sizeof(int) * n_elements));
+
+  // initialize data:
+  printf("Computing Simple Sum test\n");
+  printf("---------------------------------------------------\n");
+
+  printf("Initialize test data [1, 1, 1...]\n");
+
+  for (int i = 0; i < n_elements; i++) {
+    h_data[i] = 1;
+  }
+
+  int blockSize = 256;
+  int gridSize = n_elements / blockSize;
+  int nWarps = blockSize / 32;
+  int shmem_sz = nWarps * sizeof(int);
+  int n_partialSums = n_elements / blockSize;
+  int partial_sz = n_partialSums * sizeof(int);
+
+  printf("Scan summation for %d elements, %d partial sums\n", n_elements,
+         n_elements / blockSize);
+
+  int p_blockSize = min(n_partialSums, blockSize);
+  int p_gridSize = iDivUp(n_partialSums, p_blockSize);
+  printf("Partial summing %d elements with %d blocks of size %d\n",
+         n_partialSums, p_gridSize, p_blockSize);
+
+  // initialize a timer
+  cudaEvent_t start, stop;
+  checkCudaErrors(cudaEventCreate(&start));
+  checkCudaErrors(cudaEventCreate(&stop));
+  float et = 0;
+  float inc = 0;
+
+  checkCudaErrors(cudaMalloc(reinterpret_cast<void **>(&d_data), sz));
+  checkCudaErrors(
+      cudaMalloc(reinterpret_cast<void **>(&d_partial_sums), partial_sz));
+  checkCudaErrors(cudaMemset(d_partial_sums, 0, partial_sz));
+
+  checkCudaErrors(
+      cudaMallocHost(reinterpret_cast<void **>(&h_partial_sums), partial_sz));
+  checkCudaErrors(cudaMemcpy(d_data, h_data, sz, cudaMemcpyHostToDevice));
+
+  checkCudaErrors(cudaEventRecord(start, 0));
+  shfl_scan_test<<<gridSize, blockSize, shmem_sz>>>(d_data, 32, d_partial_sums);
+  shfl_scan_test<<<p_gridSize, p_blockSize, shmem_sz>>>(d_partial_sums, 32);
+  uniform_add<<<gridSize - 1, blockSize>>>(d_data + blockSize, d_partial_sums,
+                                           n_elements);
+  checkCudaErrors(cudaEventRecord(stop, 0));
+  checkCudaErrors(cudaEventSynchronize(stop));
+  checkCudaErrors(cudaEventElapsedTime(&inc, start, stop));
+  et += inc;
+
+  checkCudaErrors(cudaMemcpy(h_result, d_data, sz, cudaMemcpyDeviceToHost));
+  checkCudaErrors(cudaMemcpy(h_partial_sums, d_partial_sums, partial_sz,
+                             cudaMemcpyDeviceToHost));
+
+  printf("Test Sum: %d\n", h_partial_sums[n_partialSums - 1]);
+  printf("Time (ms): %f\n", et);
+  printf("%d elements scanned in %f ms -> %f MegaElements/s\n", n_elements, et,
+         n_elements / (et / 1000.0f) / 1000000.0f);
+
+  bool bTestResult = CPUverify(h_data, h_result, n_elements);
+
+  checkCudaErrors(cudaFreeHost(h_data));
+  checkCudaErrors(cudaFreeHost(h_result));
+  checkCudaErrors(cudaFreeHost(h_partial_sums));
+  checkCudaErrors(cudaFree(d_data));
+  checkCudaErrors(cudaFree(d_partial_sums));
+
+  return bTestResult;
+}
+
+// This function tests creation of an integral image using
+// synthetic data, of size 1920x1080 pixels greyscale.
+bool shuffle_integral_image_test() {
+  char *d_data;
+  unsigned int *h_image;
+  unsigned int *d_integral_image;
+  int w = 1920;
+  int h = 1080;
+  int n_elements = w * h;
+  int sz = sizeof(unsigned int) * n_elements;
+
+  printf("\nComputing Integral Image Test on size %d x %d synthetic data\n", w,
+         h);
+  printf("---------------------------------------------------\n");
+  checkCudaErrors(cudaMallocHost(reinterpret_cast<void **>(&h_image), sz));
+  // fill test "image" with synthetic 1's data
+  memset(h_image, 0, sz);
+
+  // each thread handles 16 values, use 1 block/row
+  int blockSize = iDivUp(w, 16);
+  // launch 1 block / row
+  int gridSize = h;
+
+  // Create a synthetic image for testing
+  checkCudaErrors(cudaMalloc(reinterpret_cast<void **>(&d_data), sz));
+  checkCudaErrors(cudaMalloc(reinterpret_cast<void **>(&d_integral_image),
+                             n_elements * sizeof(int) * 4));
+  checkCudaErrors(cudaMemset(d_data, 1, sz));
+  checkCudaErrors(cudaMemset(d_integral_image, 0, sz));
+
+  cudaEvent_t start, stop;
+  cudaEventCreate(&start);
+  cudaEventCreate(&stop);
+  float et = 0;
+  unsigned int err;
+
+  // Execute scan line prefix sum kernel, and time it
+  cudaEventRecord(start);
+  shfl_intimage_rows<<<gridSize, blockSize>>>(
+      reinterpret_cast<uint4 *>(d_data),
+      reinterpret_cast<uint4 *>(d_integral_image));
+  cudaEventRecord(stop);
+  checkCudaErrors(cudaEventSynchronize(stop));
+  checkCudaErrors(cudaEventElapsedTime(&et, start, stop));
+  printf("Method: Fast  Time (GPU Timer): %f ms ", et);
+
+  // verify the scan line results
+  checkCudaErrors(
+      cudaMemcpy(h_image, d_integral_image, sz, cudaMemcpyDeviceToHost));
+  err = verifyDataRowSums(h_image, w, h);
+  printf("Diff = %d\n", err);
+
+  // Execute column prefix sum kernel and time it
+  dim3 blockSz(32, 8);
+  dim3 testGrid(w / blockSz.x, 1);
+
+  cudaEventRecord(start);
+  shfl_vertical_shfl<<<testGrid, blockSz>>>((unsigned int *)d_integral_image, w,
+                                            h);
+  cudaEventRecord(stop);
+  checkCudaErrors(cudaEventSynchronize(stop));
+  checkCudaErrors(cudaEventElapsedTime(&et, start, stop));
+  printf("Method: Vertical Scan  Time (GPU Timer): %f ms ", et);
+
+  // Verify the column results
+  checkCudaErrors(
+      cudaMemcpy(h_image, d_integral_image, sz, cudaMemcpyDeviceToHost));
+  printf("\n");
+
+  int finalSum = h_image[w * h - 1];
+  printf("CheckSum: %d, (expect %dx%d=%d)\n", finalSum, w, h, w * h);
+
+  checkCudaErrors(cudaFree(d_data));
+  checkCudaErrors(cudaFree(d_integral_image));
+  checkCudaErrors(cudaFreeHost(h_image));
+  // verify final sum: if the final value in the corner is the same as the size
+  // of the buffer (all 1's) then the integral image was generated successfully
+  return (finalSum == w * h) ? true : false;
+}
+
+int main(int argc, char *argv[]) {
+  // Initialization.  The shuffle intrinsic is not available on SM < 3.0
+  // so waive the test if the hardware is not present.
+  int cuda_device = 0;
+
+  printf("Starting shfl_scan\n");
+
+  // use command-line specified CUDA device, otherwise use device with highest
+  // Gflops/s
+  cuda_device = findCudaDevice(argc, (const char **)argv);
+
+  cudaDeviceProp deviceProp;
+  checkCudaErrors(cudaGetDevice(&cuda_device));
+
+  checkCudaErrors(cudaGetDeviceProperties(&deviceProp, cuda_device));
+
+  printf("> Detected Compute SM %d.%d hardware with %d multi-processors\n",
+         deviceProp.major, deviceProp.minor, deviceProp.multiProcessorCount);
+
+  // __shfl intrinsic needs SM 3.0 or higher
+  if (deviceProp.major < 3) {
+    printf("> __shfl() intrinsic requires device SM 3.0+\n");
+    printf("> Waiving test.\n");
+    exit(EXIT_WAIVED);
+  }
+
+  bool bTestResult = true;
+  bool simpleTest = shuffle_simple_test(argc, argv);
+  bool intTest = shuffle_integral_image_test();
+
+  bTestResult = simpleTest & intTest;
+
+  exit((bTestResult) ? EXIT_SUCCESS : EXIT_FAILURE);
+}
--- a/Samples/shfl_scan/shfl_scan_vs2010.sln
+++ b/Samples/shfl_scan/shfl_scan_vs2010.sln
@@ -0,0 +1,20 @@
+
+Microsoft Visual Studio Solution File, Format Version 11.00
+# Visual Studio 2010
+Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "shfl_scan", "shfl_scan_vs2010.vcxproj", "{997E0757-EA74-4A4E-A0FC-47D8C8831A15}"
+EndProject
+Global
+	GlobalSection(SolutionConfigurationPlatforms) = preSolution
+		Debug|x64 = Debug|x64
+		Release|x64 = Release|x64
+	EndGlobalSection
+	GlobalSection(ProjectConfigurationPlatforms) = postSolution
+		{997E0757-EA74-4A4E-A0FC-47D8C8831A15}.Debug|x64.ActiveCfg = Debug|x64
+		{997E0757-EA74-4A4E-A0FC-47D8C8831A15}.Debug|x64.Build.0 = Debug|x64
+		{997E0757-EA74-4A4E-A0FC-47D8C8831A15}.Release|x64.ActiveCfg = Release|x64
+		{997E0757-EA74-4A4E-A0FC-47D8C8831A15}.Release|x64.Build.0 = Release|x64
+	EndGlobalSection
+	GlobalSection(SolutionProperties) = preSolution
+		HideSolutionNode = FALSE
+	EndGlobalSection
+EndGlobal
--- a/Samples/shfl_scan/shfl_scan_vs2010.vcxproj
+++ b/Samples/shfl_scan/shfl_scan_vs2010.vcxproj
@@ -0,0 +1,107 @@
+<?xml version="1.0" encoding="utf-8"?>
+<Project DefaultTargets="Build" ToolsVersion="4.0" xmlns="http://schemas.microsoft.com/developer/msbuild/2003">
+  <PropertyGroup>
+    <CUDAPropsPath Condition="'$(CUDAPropsPath)'==''">$(VCTargetsPath)\BuildCustomizations</CUDAPropsPath>
+  </PropertyGroup>
+  <ItemGroup Label="ProjectConfigurations">
+    <ProjectConfiguration Include="Debug|x64">
+      <Configuration>Debug</Configuration>
+      <Platform>x64</Platform>
+    </ProjectConfiguration>
+    <ProjectConfiguration Include="Release|x64">
+      <Configuration>Release</Configuration>
+      <Platform>x64</Platform>
+    </ProjectConfiguration>
+  </ItemGroup>
+  <PropertyGroup Label="Globals">
+    <ProjectGuid>{997E0757-EA74-4A4E-A0FC-47D8C8831A15}</ProjectGuid>
+    <RootNamespace>shfl_scan_vs2010</RootNamespace>
+    <ProjectName>shfl_scan</ProjectName>
+    <CudaToolkitCustomDir />
+  </PropertyGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.Default.props" />
+  <PropertyGroup>
+    <ConfigurationType>Application</ConfigurationType>
+    <CharacterSet>MultiByte</CharacterSet>
+  </PropertyGroup>
+  <PropertyGroup Condition="'$(Configuration)'=='Debug'">
+    <UseDebugLibraries>true</UseDebugLibraries>
+  </PropertyGroup>
+  <PropertyGroup Condition="'$(Configuration)'=='Release'">
+    <WholeProgramOptimization>true</WholeProgramOptimization>
+  </PropertyGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.props" />
+  <ImportGroup Label="ExtensionSettings">
+    <Import Project="$(CUDAPropsPath)\CUDA 9.2.props" />
+  </ImportGroup>
+  <ImportGroup Label="PropertySheets">
+    <Import Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" />
+  </ImportGroup>
+  <PropertyGroup Label="UserMacros" />
+  <PropertyGroup>
+    <IntDir>$(Platform)/$(Configuration)/</IntDir>
+    <IncludePath>$(IncludePath)</IncludePath>
+    <CodeAnalysisRuleSet>AllRules.ruleset</CodeAnalysisRuleSet>
+    <CodeAnalysisRules />
+    <CodeAnalysisRuleAssemblies />
+  </PropertyGroup>
+  <PropertyGroup Condition="'$(Platform)'=='x64'">
+    <OutDir>../../bin/win64/$(Configuration)/</OutDir>
+  </PropertyGroup>
+  <ItemDefinitionGroup>
+    <ClCompile>
+      <WarningLevel>Level3</WarningLevel>
+      <PreprocessorDefinitions>WIN32;_MBCS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
+      <AdditionalIncludeDirectories>./;$(CudaToolkitDir)/include;../../Common;</AdditionalIncludeDirectories>
+    </ClCompile>
+    <Link>
+      <SubSystem>Console</SubSystem>
+      <AdditionalDependencies>cudart_static.lib;kernel32.lib;user32.lib;gdi32.lib;winspool.lib;comdlg32.lib;advapi32.lib;shell32.lib;ole32.lib;oleaut32.lib;uuid.lib;odbc32.lib;odbccp32.lib;%(AdditionalDependencies)</AdditionalDependencies>
+      <AdditionalLibraryDirectories>$(CudaToolkitLibDir);</AdditionalLibraryDirectories>
+      <OutputFile>$(OutDir)/shfl_scan.exe</OutputFile>
+    </Link>
+    <CudaCompile>
+      <CodeGeneration>compute_30,sm_30;compute_35,sm_35;compute_37,sm_37;compute_50,sm_50;compute_52,sm_52;compute_60,sm_60;compute_61,sm_61;compute_70,sm_70;</CodeGeneration>
+      <AdditionalOptions>-Xcompiler "/wd 4819" %(AdditionalOptions)</AdditionalOptions>
+      <Include>./;../../Common</Include>
+      <Defines>WIN32</Defines>
+    </CudaCompile>
+  </ItemDefinitionGroup>
+  <ItemDefinitionGroup Condition="'$(Configuration)'=='Debug'">
+    <ClCompile>
+      <Optimization>Disabled</Optimization>
+      <RuntimeLibrary>MultiThreadedDebug</RuntimeLibrary>
+    </ClCompile>
+    <Link>
+      <GenerateDebugInformation>true</GenerateDebugInformation>
+      <LinkTimeCodeGeneration>Default</LinkTimeCodeGeneration>
+    </Link>
+    <CudaCompile>
+      <Runtime>MTd</Runtime>
+      <TargetMachinePlatform>64</TargetMachinePlatform>
+    </CudaCompile>
+  </ItemDefinitionGroup>
+  <ItemDefinitionGroup Condition="'$(Configuration)'=='Release'">
+    <ClCompile>
+      <Optimization>MaxSpeed</Optimization>
+      <RuntimeLibrary>MultiThreaded</RuntimeLibrary>
+    </ClCompile>
+    <Link>
+      <GenerateDebugInformation>false</GenerateDebugInformation>
+      <LinkTimeCodeGeneration>UseLinkTimeCodeGeneration</LinkTimeCodeGeneration>
+    </Link>
+    <CudaCompile>
+      <Runtime>MT</Runtime>
+      <TargetMachinePlatform>64</TargetMachinePlatform>
+    </CudaCompile>
+  </ItemDefinitionGroup>
+  <ItemGroup>
+    <CudaCompile Include="shfl_scan.cu" />
+    <ClInclude Include="util.h" />
+    <None Include="shfl_integral_image.cuh" />
+  </ItemGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" />
+  <ImportGroup Label="ExtensionTargets">
+    <Import Project="$(CUDAPropsPath)\CUDA 9.2.targets" />
+  </ImportGroup>
+</Project>
--- a/Samples/shfl_scan/shfl_scan_vs2012.sln
+++ b/Samples/shfl_scan/shfl_scan_vs2012.sln
@@ -0,0 +1,20 @@
+
+Microsoft Visual Studio Solution File, Format Version 12.00
+# Visual Studio 2012
+Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "shfl_scan", "shfl_scan_vs2012.vcxproj", "{997E0757-EA74-4A4E-A0FC-47D8C8831A15}"
+EndProject
+Global
+	GlobalSection(SolutionConfigurationPlatforms) = preSolution
+		Debug|x64 = Debug|x64
+		Release|x64 = Release|x64
+	EndGlobalSection
+	GlobalSection(ProjectConfigurationPlatforms) = postSolution
+		{997E0757-EA74-4A4E-A0FC-47D8C8831A15}.Debug|x64.ActiveCfg = Debug|x64
+		{997E0757-EA74-4A4E-A0FC-47D8C8831A15}.Debug|x64.Build.0 = Debug|x64
+		{997E0757-EA74-4A4E-A0FC-47D8C8831A15}.Release|x64.ActiveCfg = Release|x64
+		{997E0757-EA74-4A4E-A0FC-47D8C8831A15}.Release|x64.Build.0 = Release|x64
+	EndGlobalSection
+	GlobalSection(SolutionProperties) = preSolution
+		HideSolutionNode = FALSE
+	EndGlobalSection
+EndGlobal
--- a/Samples/shfl_scan/shfl_scan_vs2012.vcxproj
+++ b/Samples/shfl_scan/shfl_scan_vs2012.vcxproj
@@ -0,0 +1,108 @@
+<?xml version="1.0" encoding="utf-8"?>
+<Project DefaultTargets="Build" ToolsVersion="4.0" xmlns="http://schemas.microsoft.com/developer/msbuild/2003">
+  <PropertyGroup>
+    <CUDAPropsPath Condition="'$(CUDAPropsPath)'==''">$(VCTargetsPath)\BuildCustomizations</CUDAPropsPath>
+  </PropertyGroup>
+  <ItemGroup Label="ProjectConfigurations">
+    <ProjectConfiguration Include="Debug|x64">
+      <Configuration>Debug</Configuration>
+      <Platform>x64</Platform>
+    </ProjectConfiguration>
+    <ProjectConfiguration Include="Release|x64">
+      <Configuration>Release</Configuration>
+      <Platform>x64</Platform>
+    </ProjectConfiguration>
+  </ItemGroup>
+  <PropertyGroup Label="Globals">
+    <ProjectGuid>{997E0757-EA74-4A4E-A0FC-47D8C8831A15}</ProjectGuid>
+    <RootNamespace>shfl_scan_vs2012</RootNamespace>
+    <ProjectName>shfl_scan</ProjectName>
+    <CudaToolkitCustomDir />
+  </PropertyGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.Default.props" />
+  <PropertyGroup>
+    <ConfigurationType>Application</ConfigurationType>
+    <CharacterSet>MultiByte</CharacterSet>
+    <PlatformToolset>v110</PlatformToolset>
+  </PropertyGroup>
+  <PropertyGroup Condition="'$(Configuration)'=='Debug'">
+    <UseDebugLibraries>true</UseDebugLibraries>
+  </PropertyGroup>
+  <PropertyGroup Condition="'$(Configuration)'=='Release'">
+    <WholeProgramOptimization>true</WholeProgramOptimization>
+  </PropertyGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.props" />
+  <ImportGroup Label="ExtensionSettings">
+    <Import Project="$(CUDAPropsPath)\CUDA 9.2.props" />
+  </ImportGroup>
+  <ImportGroup Label="PropertySheets">
+    <Import Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" />
+  </ImportGroup>
+  <PropertyGroup Label="UserMacros" />
+  <PropertyGroup>
+    <IntDir>$(Platform)/$(Configuration)/</IntDir>
+    <IncludePath>$(IncludePath)</IncludePath>
+    <CodeAnalysisRuleSet>AllRules.ruleset</CodeAnalysisRuleSet>
+    <CodeAnalysisRules />
+    <CodeAnalysisRuleAssemblies />
+  </PropertyGroup>
+  <PropertyGroup Condition="'$(Platform)'=='x64'">
+    <OutDir>../../bin/win64/$(Configuration)/</OutDir>
+  </PropertyGroup>
+  <ItemDefinitionGroup>
+    <ClCompile>
+      <WarningLevel>Level3</WarningLevel>
+      <PreprocessorDefinitions>WIN32;_MBCS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
+      <AdditionalIncludeDirectories>./;$(CudaToolkitDir)/include;../../Common;</AdditionalIncludeDirectories>
+    </ClCompile>
+    <Link>
+      <SubSystem>Console</SubSystem>
+      <AdditionalDependencies>cudart_static.lib;kernel32.lib;user32.lib;gdi32.lib;winspool.lib;comdlg32.lib;advapi32.lib;shell32.lib;ole32.lib;oleaut32.lib;uuid.lib;odbc32.lib;odbccp32.lib;%(AdditionalDependencies)</AdditionalDependencies>
+      <AdditionalLibraryDirectories>$(CudaToolkitLibDir);</AdditionalLibraryDirectories>
+      <OutputFile>$(OutDir)/shfl_scan.exe</OutputFile>
+    </Link>
+    <CudaCompile>
+      <CodeGeneration>compute_30,sm_30;compute_35,sm_35;compute_37,sm_37;compute_50,sm_50;compute_52,sm_52;compute_60,sm_60;compute_61,sm_61;compute_70,sm_70;</CodeGeneration>
+      <AdditionalOptions>-Xcompiler "/wd 4819" %(AdditionalOptions)</AdditionalOptions>
+      <Include>./;../../Common</Include>
+      <Defines>WIN32</Defines>
+    </CudaCompile>
+  </ItemDefinitionGroup>
+  <ItemDefinitionGroup Condition="'$(Configuration)'=='Debug'">
+    <ClCompile>
+      <Optimization>Disabled</Optimization>
+      <RuntimeLibrary>MultiThreadedDebug</RuntimeLibrary>
+    </ClCompile>
+    <Link>
+      <GenerateDebugInformation>true</GenerateDebugInformation>
+      <LinkTimeCodeGeneration>Default</LinkTimeCodeGeneration>
+    </Link>
+    <CudaCompile>
+      <Runtime>MTd</Runtime>
+      <TargetMachinePlatform>64</TargetMachinePlatform>
+    </CudaCompile>
+  </ItemDefinitionGroup>
+  <ItemDefinitionGroup Condition="'$(Configuration)'=='Release'">
+    <ClCompile>
+      <Optimization>MaxSpeed</Optimization>
+      <RuntimeLibrary>MultiThreaded</RuntimeLibrary>
+    </ClCompile>
+    <Link>
+      <GenerateDebugInformation>false</GenerateDebugInformation>
+      <LinkTimeCodeGeneration>UseLinkTimeCodeGeneration</LinkTimeCodeGeneration>
+    </Link>
+    <CudaCompile>
+      <Runtime>MT</Runtime>
+      <TargetMachinePlatform>64</TargetMachinePlatform>
+    </CudaCompile>
+  </ItemDefinitionGroup>
+  <ItemGroup>
+    <CudaCompile Include="shfl_scan.cu" />
+    <ClInclude Include="util.h" />
+    <None Include="shfl_integral_image.cuh" />
+  </ItemGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" />
+  <ImportGroup Label="ExtensionTargets">
+    <Import Project="$(CUDAPropsPath)\CUDA 9.2.targets" />
+  </ImportGroup>
+</Project>
--- a/Samples/shfl_scan/shfl_scan_vs2013.sln
+++ b/Samples/shfl_scan/shfl_scan_vs2013.sln
@@ -0,0 +1,20 @@
+
+Microsoft Visual Studio Solution File, Format Version 13.00
+# Visual Studio 2013
+Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "shfl_scan", "shfl_scan_vs2013.vcxproj", "{997E0757-EA74-4A4E-A0FC-47D8C8831A15}"
+EndProject
+Global
+	GlobalSection(SolutionConfigurationPlatforms) = preSolution
+		Debug|x64 = Debug|x64
+		Release|x64 = Release|x64
+	EndGlobalSection
+	GlobalSection(ProjectConfigurationPlatforms) = postSolution
+		{997E0757-EA74-4A4E-A0FC-47D8C8831A15}.Debug|x64.ActiveCfg = Debug|x64
+		{997E0757-EA74-4A4E-A0FC-47D8C8831A15}.Debug|x64.Build.0 = Debug|x64
+		{997E0757-EA74-4A4E-A0FC-47D8C8831A15}.Release|x64.ActiveCfg = Release|x64
+		{997E0757-EA74-4A4E-A0FC-47D8C8831A15}.Release|x64.Build.0 = Release|x64
+	EndGlobalSection
+	GlobalSection(SolutionProperties) = preSolution
+		HideSolutionNode = FALSE
+	EndGlobalSection
+EndGlobal
--- a/Samples/shfl_scan/shfl_scan_vs2013.vcxproj
+++ b/Samples/shfl_scan/shfl_scan_vs2013.vcxproj
@@ -0,0 +1,108 @@
+<?xml version="1.0" encoding="utf-8"?>
+<Project DefaultTargets="Build" ToolsVersion="4.0" xmlns="http://schemas.microsoft.com/developer/msbuild/2003">
+  <PropertyGroup>
+    <CUDAPropsPath Condition="'$(CUDAPropsPath)'==''">$(VCTargetsPath)\BuildCustomizations</CUDAPropsPath>
+  </PropertyGroup>
+  <ItemGroup Label="ProjectConfigurations">
+    <ProjectConfiguration Include="Debug|x64">
+      <Configuration>Debug</Configuration>
+      <Platform>x64</Platform>
+    </ProjectConfiguration>
+    <ProjectConfiguration Include="Release|x64">
+      <Configuration>Release</Configuration>
+      <Platform>x64</Platform>
+    </ProjectConfiguration>
+  </ItemGroup>
+  <PropertyGroup Label="Globals">
+    <ProjectGuid>{997E0757-EA74-4A4E-A0FC-47D8C8831A15}</ProjectGuid>
+    <RootNamespace>shfl_scan_vs2013</RootNamespace>
+    <ProjectName>shfl_scan</ProjectName>
+    <CudaToolkitCustomDir />
+  </PropertyGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.Default.props" />
+  <PropertyGroup>
+    <ConfigurationType>Application</ConfigurationType>
+    <CharacterSet>MultiByte</CharacterSet>
+    <PlatformToolset>v120</PlatformToolset>
+  </PropertyGroup>
+  <PropertyGroup Condition="'$(Configuration)'=='Debug'">
+    <UseDebugLibraries>true</UseDebugLibraries>
+  </PropertyGroup>
+  <PropertyGroup Condition="'$(Configuration)'=='Release'">
+    <WholeProgramOptimization>true</WholeProgramOptimization>
+  </PropertyGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.props" />
+  <ImportGroup Label="ExtensionSettings">
+    <Import Project="$(CUDAPropsPath)\CUDA 9.2.props" />
+  </ImportGroup>
+  <ImportGroup Label="PropertySheets">
+    <Import Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" />
+  </ImportGroup>
+  <PropertyGroup Label="UserMacros" />
+  <PropertyGroup>
+    <IntDir>$(Platform)/$(Configuration)/</IntDir>
+    <IncludePath>$(IncludePath)</IncludePath>
+    <CodeAnalysisRuleSet>AllRules.ruleset</CodeAnalysisRuleSet>
+    <CodeAnalysisRules />
+    <CodeAnalysisRuleAssemblies />
+  </PropertyGroup>
+  <PropertyGroup Condition="'$(Platform)'=='x64'">
+    <OutDir>../../bin/win64/$(Configuration)/</OutDir>
+  </PropertyGroup>
+  <ItemDefinitionGroup>
+    <ClCompile>
+      <WarningLevel>Level3</WarningLevel>
+      <PreprocessorDefinitions>WIN32;_MBCS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
+      <AdditionalIncludeDirectories>./;$(CudaToolkitDir)/include;../../Common;</AdditionalIncludeDirectories>
+    </ClCompile>
+    <Link>
+      <SubSystem>Console</SubSystem>
+      <AdditionalDependencies>cudart_static.lib;kernel32.lib;user32.lib;gdi32.lib;winspool.lib;comdlg32.lib;advapi32.lib;shell32.lib;ole32.lib;oleaut32.lib;uuid.lib;odbc32.lib;odbccp32.lib;%(AdditionalDependencies)</AdditionalDependencies>
+      <AdditionalLibraryDirectories>$(CudaToolkitLibDir);</AdditionalLibraryDirectories>
+      <OutputFile>$(OutDir)/shfl_scan.exe</OutputFile>
+    </Link>
+    <CudaCompile>
+      <CodeGeneration>compute_30,sm_30;compute_35,sm_35;compute_37,sm_37;compute_50,sm_50;compute_52,sm_52;compute_60,sm_60;compute_61,sm_61;compute_70,sm_70;</CodeGeneration>
+      <AdditionalOptions>-Xcompiler "/wd 4819" %(AdditionalOptions)</AdditionalOptions>
+      <Include>./;../../Common</Include>
+      <Defines>WIN32</Defines>
+    </CudaCompile>
+  </ItemDefinitionGroup>
+  <ItemDefinitionGroup Condition="'$(Configuration)'=='Debug'">
+    <ClCompile>
+      <Optimization>Disabled</Optimization>
+      <RuntimeLibrary>MultiThreadedDebug</RuntimeLibrary>
+    </ClCompile>
+    <Link>
+      <GenerateDebugInformation>true</GenerateDebugInformation>
+      <LinkTimeCodeGeneration>Default</LinkTimeCodeGeneration>
+    </Link>
+    <CudaCompile>
+      <Runtime>MTd</Runtime>
+      <TargetMachinePlatform>64</TargetMachinePlatform>
+    </CudaCompile>
+  </ItemDefinitionGroup>
+  <ItemDefinitionGroup Condition="'$(Configuration)'=='Release'">
+    <ClCompile>
+      <Optimization>MaxSpeed</Optimization>
+      <RuntimeLibrary>MultiThreaded</RuntimeLibrary>
+    </ClCompile>
+    <Link>
+      <GenerateDebugInformation>false</GenerateDebugInformation>
+      <LinkTimeCodeGeneration>UseLinkTimeCodeGeneration</LinkTimeCodeGeneration>
+    </Link>
+    <CudaCompile>
+      <Runtime>MT</Runtime>
+      <TargetMachinePlatform>64</TargetMachinePlatform>
+    </CudaCompile>
+  </ItemDefinitionGroup>
+  <ItemGroup>
+    <CudaCompile Include="shfl_scan.cu" />
+    <ClInclude Include="util.h" />
+    <None Include="shfl_integral_image.cuh" />
+  </ItemGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" />
+  <ImportGroup Label="ExtensionTargets">
+    <Import Project="$(CUDAPropsPath)\CUDA 9.2.targets" />
+  </ImportGroup>
+</Project>
--- a/Samples/shfl_scan/shfl_scan_vs2015.sln
+++ b/Samples/shfl_scan/shfl_scan_vs2015.sln
@@ -0,0 +1,20 @@
+
+Microsoft Visual Studio Solution File, Format Version 14.00
+# Visual Studio 2015
+Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "shfl_scan", "shfl_scan_vs2015.vcxproj", "{997E0757-EA74-4A4E-A0FC-47D8C8831A15}"
+EndProject
+Global
+	GlobalSection(SolutionConfigurationPlatforms) = preSolution
+		Debug|x64 = Debug|x64
+		Release|x64 = Release|x64
+	EndGlobalSection
+	GlobalSection(ProjectConfigurationPlatforms) = postSolution
+		{997E0757-EA74-4A4E-A0FC-47D8C8831A15}.Debug|x64.ActiveCfg = Debug|x64
+		{997E0757-EA74-4A4E-A0FC-47D8C8831A15}.Debug|x64.Build.0 = Debug|x64
+		{997E0757-EA74-4A4E-A0FC-47D8C8831A15}.Release|x64.ActiveCfg = Release|x64
+		{997E0757-EA74-4A4E-A0FC-47D8C8831A15}.Release|x64.Build.0 = Release|x64
+	EndGlobalSection
+	GlobalSection(SolutionProperties) = preSolution
+		HideSolutionNode = FALSE
+	EndGlobalSection
+EndGlobal
--- a/Samples/shfl_scan/shfl_scan_vs2015.vcxproj
+++ b/Samples/shfl_scan/shfl_scan_vs2015.vcxproj
@@ -0,0 +1,108 @@
+<?xml version="1.0" encoding="utf-8"?>
+<Project DefaultTargets="Build" ToolsVersion="4.0" xmlns="http://schemas.microsoft.com/developer/msbuild/2003">
+  <PropertyGroup>
+    <CUDAPropsPath Condition="'$(CUDAPropsPath)'==''">$(VCTargetsPath)\BuildCustomizations</CUDAPropsPath>
+  </PropertyGroup>
+  <ItemGroup Label="ProjectConfigurations">
+    <ProjectConfiguration Include="Debug|x64">
+      <Configuration>Debug</Configuration>
+      <Platform>x64</Platform>
+    </ProjectConfiguration>
+    <ProjectConfiguration Include="Release|x64">
+      <Configuration>Release</Configuration>
+      <Platform>x64</Platform>
+    </ProjectConfiguration>
+  </ItemGroup>
+  <PropertyGroup Label="Globals">
+    <ProjectGuid>{997E0757-EA74-4A4E-A0FC-47D8C8831A15}</ProjectGuid>
+    <RootNamespace>shfl_scan_vs2015</RootNamespace>
+    <ProjectName>shfl_scan</ProjectName>
+    <CudaToolkitCustomDir />
+  </PropertyGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.Default.props" />
+  <PropertyGroup>
+    <ConfigurationType>Application</ConfigurationType>
+    <CharacterSet>MultiByte</CharacterSet>
+    <PlatformToolset>v140</PlatformToolset>
+  </PropertyGroup>
+  <PropertyGroup Condition="'$(Configuration)'=='Debug'">
+    <UseDebugLibraries>true</UseDebugLibraries>
+  </PropertyGroup>
+  <PropertyGroup Condition="'$(Configuration)'=='Release'">
+    <WholeProgramOptimization>true</WholeProgramOptimization>
+  </PropertyGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.props" />
+  <ImportGroup Label="ExtensionSettings">
+    <Import Project="$(CUDAPropsPath)\CUDA 9.2.props" />
+  </ImportGroup>
+  <ImportGroup Label="PropertySheets">
+    <Import Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" />
+  </ImportGroup>
+  <PropertyGroup Label="UserMacros" />
+  <PropertyGroup>
+    <IntDir>$(Platform)/$(Configuration)/</IntDir>
+    <IncludePath>$(IncludePath)</IncludePath>
+    <CodeAnalysisRuleSet>AllRules.ruleset</CodeAnalysisRuleSet>
+    <CodeAnalysisRules />
+    <CodeAnalysisRuleAssemblies />
+  </PropertyGroup>
+  <PropertyGroup Condition="'$(Platform)'=='x64'">
+    <OutDir>../../bin/win64/$(Configuration)/</OutDir>
+  </PropertyGroup>
+  <ItemDefinitionGroup>
+    <ClCompile>
+      <WarningLevel>Level3</WarningLevel>
+      <PreprocessorDefinitions>WIN32;_MBCS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
+      <AdditionalIncludeDirectories>./;$(CudaToolkitDir)/include;../../Common;</AdditionalIncludeDirectories>
+    </ClCompile>
+    <Link>
+      <SubSystem>Console</SubSystem>
+      <AdditionalDependencies>cudart_static.lib;kernel32.lib;user32.lib;gdi32.lib;winspool.lib;comdlg32.lib;advapi32.lib;shell32.lib;ole32.lib;oleaut32.lib;uuid.lib;odbc32.lib;odbccp32.lib;%(AdditionalDependencies)</AdditionalDependencies>
+      <AdditionalLibraryDirectories>$(CudaToolkitLibDir);</AdditionalLibraryDirectories>
+      <OutputFile>$(OutDir)/shfl_scan.exe</OutputFile>
+    </Link>
+    <CudaCompile>
+      <CodeGeneration>compute_30,sm_30;compute_35,sm_35;compute_37,sm_37;compute_50,sm_50;compute_52,sm_52;compute_60,sm_60;compute_61,sm_61;compute_70,sm_70;</CodeGeneration>
+      <AdditionalOptions>-Xcompiler "/wd 4819" %(AdditionalOptions)</AdditionalOptions>
+      <Include>./;../../Common</Include>
+      <Defines>WIN32</Defines>
+    </CudaCompile>
+  </ItemDefinitionGroup>
+  <ItemDefinitionGroup Condition="'$(Configuration)'=='Debug'">
+    <ClCompile>
+      <Optimization>Disabled</Optimization>
+      <RuntimeLibrary>MultiThreadedDebug</RuntimeLibrary>
+    </ClCompile>
+    <Link>
+      <GenerateDebugInformation>true</GenerateDebugInformation>
+      <LinkTimeCodeGeneration>Default</LinkTimeCodeGeneration>
+    </Link>
+    <CudaCompile>
+      <Runtime>MTd</Runtime>
+      <TargetMachinePlatform>64</TargetMachinePlatform>
+    </CudaCompile>
+  </ItemDefinitionGroup>
+  <ItemDefinitionGroup Condition="'$(Configuration)'=='Release'">
+    <ClCompile>
+      <Optimization>MaxSpeed</Optimization>
+      <RuntimeLibrary>MultiThreaded</RuntimeLibrary>
+    </ClCompile>
+    <Link>
+      <GenerateDebugInformation>false</GenerateDebugInformation>
+      <LinkTimeCodeGeneration>UseLinkTimeCodeGeneration</LinkTimeCodeGeneration>
+    </Link>
+    <CudaCompile>
+      <Runtime>MT</Runtime>
+      <TargetMachinePlatform>64</TargetMachinePlatform>
+    </CudaCompile>
+  </ItemDefinitionGroup>
+  <ItemGroup>
+    <CudaCompile Include="shfl_scan.cu" />
+    <ClInclude Include="util.h" />
+    <None Include="shfl_integral_image.cuh" />
+  </ItemGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" />
+  <ImportGroup Label="ExtensionTargets">
+    <Import Project="$(CUDAPropsPath)\CUDA 9.2.targets" />
+  </ImportGroup>
+</Project>
--- a/Samples/shfl_scan/shfl_scan_vs2017.sln
+++ b/Samples/shfl_scan/shfl_scan_vs2017.sln
@@ -0,0 +1,20 @@
+
+Microsoft Visual Studio Solution File, Format Version 12.00
+# Visual Studio 2017
+Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "shfl_scan", "shfl_scan_vs2017.vcxproj", "{997E0757-EA74-4A4E-A0FC-47D8C8831A15}"
+EndProject
+Global
+	GlobalSection(SolutionConfigurationPlatforms) = preSolution
+		Debug|x64 = Debug|x64
+		Release|x64 = Release|x64
+	EndGlobalSection
+	GlobalSection(ProjectConfigurationPlatforms) = postSolution
+		{997E0757-EA74-4A4E-A0FC-47D8C8831A15}.Debug|x64.ActiveCfg = Debug|x64
+		{997E0757-EA74-4A4E-A0FC-47D8C8831A15}.Debug|x64.Build.0 = Debug|x64
+		{997E0757-EA74-4A4E-A0FC-47D8C8831A15}.Release|x64.ActiveCfg = Release|x64
+		{997E0757-EA74-4A4E-A0FC-47D8C8831A15}.Release|x64.Build.0 = Release|x64
+	EndGlobalSection
+	GlobalSection(SolutionProperties) = preSolution
+		HideSolutionNode = FALSE
+	EndGlobalSection
+EndGlobal
--- a/Samples/shfl_scan/shfl_scan_vs2017.vcxproj
+++ b/Samples/shfl_scan/shfl_scan_vs2017.vcxproj
@@ -0,0 +1,109 @@
+<?xml version="1.0" encoding="utf-8"?>
+<Project DefaultTargets="Build" ToolsVersion="4.0" xmlns="http://schemas.microsoft.com/developer/msbuild/2003">
+  <PropertyGroup>
+    <CUDAPropsPath Condition="'$(CUDAPropsPath)'==''">$(VCTargetsPath)\BuildCustomizations</CUDAPropsPath>
+  </PropertyGroup>
+  <ItemGroup Label="ProjectConfigurations">
+    <ProjectConfiguration Include="Debug|x64">
+      <Configuration>Debug</Configuration>
+      <Platform>x64</Platform>
+    </ProjectConfiguration>
+    <ProjectConfiguration Include="Release|x64">
+      <Configuration>Release</Configuration>
+      <Platform>x64</Platform>
+    </ProjectConfiguration>
+  </ItemGroup>
+  <PropertyGroup Label="Globals">
+    <ProjectGuid>{997E0757-EA74-4A4E-A0FC-47D8C8831A15}</ProjectGuid>
+    <RootNamespace>shfl_scan_vs2017</RootNamespace>
+    <ProjectName>shfl_scan</ProjectName>
+    <CudaToolkitCustomDir />
+  </PropertyGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.Default.props" />
+  <PropertyGroup>
+    <ConfigurationType>Application</ConfigurationType>
+    <CharacterSet>MultiByte</CharacterSet>
+    <PlatformToolset>v141</PlatformToolset>
+	<WindowsTargetPlatformVersion>10.0.15063.0</WindowsTargetPlatformVersion>
+  </PropertyGroup>
+  <PropertyGroup Condition="'$(Configuration)'=='Debug'">
+    <UseDebugLibraries>true</UseDebugLibraries>
+  </PropertyGroup>
+  <PropertyGroup Condition="'$(Configuration)'=='Release'">
+    <WholeProgramOptimization>true</WholeProgramOptimization>
+  </PropertyGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.props" />
+  <ImportGroup Label="ExtensionSettings">
+    <Import Project="$(CUDAPropsPath)\CUDA 9.2.props" />
+  </ImportGroup>
+  <ImportGroup Label="PropertySheets">
+    <Import Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" />
+  </ImportGroup>
+  <PropertyGroup Label="UserMacros" />
+  <PropertyGroup>
+    <IntDir>$(Platform)/$(Configuration)/</IntDir>
+    <IncludePath>$(IncludePath)</IncludePath>
+    <CodeAnalysisRuleSet>AllRules.ruleset</CodeAnalysisRuleSet>
+    <CodeAnalysisRules />
+    <CodeAnalysisRuleAssemblies />
+  </PropertyGroup>
+  <PropertyGroup Condition="'$(Platform)'=='x64'">
+    <OutDir>../../bin/win64/$(Configuration)/</OutDir>
+  </PropertyGroup>
+  <ItemDefinitionGroup>
+    <ClCompile>
+      <WarningLevel>Level3</WarningLevel>
+      <PreprocessorDefinitions>WIN32;_MBCS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
+      <AdditionalIncludeDirectories>./;$(CudaToolkitDir)/include;../../Common;</AdditionalIncludeDirectories>
+    </ClCompile>
+    <Link>
+      <SubSystem>Console</SubSystem>
+      <AdditionalDependencies>cudart_static.lib;kernel32.lib;user32.lib;gdi32.lib;winspool.lib;comdlg32.lib;advapi32.lib;shell32.lib;ole32.lib;oleaut32.lib;uuid.lib;odbc32.lib;odbccp32.lib;%(AdditionalDependencies)</AdditionalDependencies>
+      <AdditionalLibraryDirectories>$(CudaToolkitLibDir);</AdditionalLibraryDirectories>
+      <OutputFile>$(OutDir)/shfl_scan.exe</OutputFile>
+    </Link>
+    <CudaCompile>
+      <CodeGeneration>compute_30,sm_30;compute_35,sm_35;compute_37,sm_37;compute_50,sm_50;compute_52,sm_52;compute_60,sm_60;compute_61,sm_61;compute_70,sm_70;</CodeGeneration>
+      <AdditionalOptions>-Xcompiler "/wd 4819" %(AdditionalOptions)</AdditionalOptions>
+      <Include>./;../../Common</Include>
+      <Defines>WIN32</Defines>
+    </CudaCompile>
+  </ItemDefinitionGroup>
+  <ItemDefinitionGroup Condition="'$(Configuration)'=='Debug'">
+    <ClCompile>
+      <Optimization>Disabled</Optimization>
+      <RuntimeLibrary>MultiThreadedDebug</RuntimeLibrary>
+    </ClCompile>
+    <Link>
+      <GenerateDebugInformation>true</GenerateDebugInformation>
+      <LinkTimeCodeGeneration>Default</LinkTimeCodeGeneration>
+    </Link>
+    <CudaCompile>
+      <Runtime>MTd</Runtime>
+      <TargetMachinePlatform>64</TargetMachinePlatform>
+    </CudaCompile>
+  </ItemDefinitionGroup>
+  <ItemDefinitionGroup Condition="'$(Configuration)'=='Release'">
+    <ClCompile>
+      <Optimization>MaxSpeed</Optimization>
+      <RuntimeLibrary>MultiThreaded</RuntimeLibrary>
+    </ClCompile>
+    <Link>
+      <GenerateDebugInformation>false</GenerateDebugInformation>
+      <LinkTimeCodeGeneration>UseLinkTimeCodeGeneration</LinkTimeCodeGeneration>
+    </Link>
+    <CudaCompile>
+      <Runtime>MT</Runtime>
+      <TargetMachinePlatform>64</TargetMachinePlatform>
+    </CudaCompile>
+  </ItemDefinitionGroup>
+  <ItemGroup>
+    <CudaCompile Include="shfl_scan.cu" />
+    <ClInclude Include="util.h" />
+    <None Include="shfl_integral_image.cuh" />
+  </ItemGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" />
+  <ImportGroup Label="ExtensionTargets">
+    <Import Project="$(CUDAPropsPath)\CUDA 9.2.targets" />
+  </ImportGroup>
+</Project>
--- a/Samples/shfl_scan/util.h
+++ b/Samples/shfl_scan/util.h
@@ -0,0 +1,61 @@
+/* Copyright (c) 2018, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions
+ * are met:
+ *  * Redistributions of source code must retain the above copyright
+ *    notice, this list of conditions and the following disclaimer.
+ *  * Redistributions in binary form must reproduce the above copyright
+ *    notice, this list of conditions and the following disclaimer in the
+ *    documentation and/or other materials provided with the distribution.
+ *  * Neither the name of NVIDIA CORPORATION nor the names of its
+ *    contributors may be used to endorse or promote products derived
+ *    from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
+ * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
+ * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+#ifndef SAMPLES_SHFL_SCAN_UTIL_H_
+#define SAMPLES_SHFL_SCAN_UTIL_H_
+
+// Macro to catch CUDA errors in kernel launches
+#define CHECK_LAUNCH_ERROR()                                                  \
+  do {                                                                        \
+    /* Check synchronous errors, i.e. pre-launch */                           \
+    cudaError_t err = cudaGetLastError();                                     \
+    if (cudaSuccess != err) {                                                 \
+      fprintf(stderr, "Cuda error in file '%s' in line %i : %s.\n", __FILE__, \
+              __LINE__, cudaGetErrorString(err));                             \
+      exit(EXIT_FAILURE);                                                     \
+    }                                                                         \
+    /* Check asynchronous errors, i.e. kernel failed (ULF) */                 \
+    err = cudaDeviceSynchronize();                                            \
+    if (cudaSuccess != err) {                                                 \
+      fprintf(stderr, "Cuda error in file '%s' in line %i : %s!\n", __FILE__, \
+              __LINE__, cudaGetErrorString(err));                             \
+      exit(EXIT_FAILURE);                                                     \
+    }                                                                         \
+  } while (0)
+
+// Macro to catch CUDA errors in CUDA runtime calls
+#define CUDA_CHECK(call)                                                      \
+  do {                                                                        \
+    cudaError_t err = call;                                                   \
+    if (cudaSuccess != err) {                                                 \
+      fprintf(stderr, "Cuda error in file '%s' in line %i : %s.\n", __FILE__, \
+              __LINE__, cudaGetErrorString(err));                             \
+      exit(EXIT_FAILURE);                                                     \
+    }                                                                         \
+  } while (0)
+
+#endif  // SAMPLES_SHFL_SCAN_UTIL_H_