diff --git a/CMakeLists.txt b/CMakeLists.txt
index daba3c4..364455c 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -29,7 +29,6 @@ if (CUDA_FOUND)
 	add_subdirectory(chapter08)
 	add_subdirectory(chapter09)
 	add_subdirectory(chapter10)
-	add_subdirectory(chapter11)
 else()
 	message("CUDA not found!")
 endif()
diff --git a/chapter11/CMakeLists.txt b/chapter11/CMakeLists.txt
deleted file mode 100644
index 321f69b..0000000
--- a/chapter11/CMakeLists.txt
+++ /dev/null
@@ -1,11 +0,0 @@
-CMAKE_MINIMUM_REQUIRED(VERSION 3.1)
-
-
-CUDA_ADD_EXECUTABLE(multidevice multidevice.cu)
-SET_PROPERTY(TARGET multidevice  PROPERTY FOLDER chapter11)
-
-CUDA_ADD_EXECUTABLE(portable portable.cu)
-SET_PROPERTY(TARGET portable  PROPERTY FOLDER chapter11)
-
-CUDA_ADD_EXECUTABLE(time_zero2 time_zero2.cu)
-SET_PROPERTY(TARGET time_zero2  PROPERTY FOLDER chapter11)
\ No newline at end of file
diff --git a/chapter11/multidevice.cu b/chapter11/multidevice.cu
deleted file mode 100644
index a6d2d68..0000000
--- a/chapter11/multidevice.cu
+++ /dev/null
@@ -1,167 +0,0 @@
-/*
- * Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.
- *
- * NVIDIA Corporation and its licensors retain all intellectual property and 
- * proprietary rights in and to this software and related documentation. 
- * Any use, reproduction, disclosure, or distribution of this software 
- * and related documentation without an express license agreement from
- * NVIDIA Corporation is strictly prohibited.
- *
- * Please refer to the applicable NVIDIA end user license agreement (EULA) 
- * associated with this source code for terms and conditions that govern 
- * your use of this NVIDIA software.
- * 
- */
-
-
-#include "../common/book.h"
-
-
-#define imin(a,b) (a<b?a:b)
-
-#define     N    (33*1024*1024)
-const int threadsPerBlock = 256;
-const int blocksPerGrid =
-            imin( 32, (N/2+threadsPerBlock-1) / threadsPerBlock );
-
-
-__global__ void dot( int size, float *a, float *b, float *c ) {
-    __shared__ float cache[threadsPerBlock];
-    int tid = threadIdx.x + blockIdx.x * blockDim.x;
-    int cacheIndex = threadIdx.x;
-
-    float   temp = 0;
-    while (tid < size) {
-        temp += a[tid] * b[tid];
-        tid += blockDim.x * gridDim.x;
-    }
-    
-    // set the cache values
-    cache[cacheIndex] = temp;
-    
-    // synchronize threads in this block
-    __syncthreads();
-
-    // for reductions, threadsPerBlock must be a power of 2
-    // because of the following code
-    int i = blockDim.x/2;
-    while (i != 0) {
-        if (cacheIndex < i)
-            cache[cacheIndex] += cache[cacheIndex + i];
-        __syncthreads();
-        i /= 2;
-    }
-
-    if (cacheIndex == 0)
-        c[blockIdx.x] = cache[0];
-}
-
-
-struct DataStruct {
-    int     deviceID;
-    int     size;
-    float   *a;
-    float   *b;
-    float   returnValue;
-};
-
-
-CUT_THREADROUTINE routine( void *pvoidData ) {
-    DataStruct  *data = (DataStruct*)pvoidData;
-    HANDLE_ERROR( cudaSetDevice( data->deviceID ) );
-
-    int     size = data->size;
-    float   *a, *b, c, *partial_c;
-    float   *dev_a, *dev_b, *dev_partial_c;
-
-    // allocate memory on the CPU side
-    a = data->a;
-    b = data->b;
-    partial_c = (float*)malloc( blocksPerGrid*sizeof(float) );
-
-    // allocate the memory on the GPU
-    HANDLE_ERROR( cudaMalloc( (void**)&dev_a,
-                              size*sizeof(float) ) );
-    HANDLE_ERROR( cudaMalloc( (void**)&dev_b,
-                              size*sizeof(float) ) );
-    HANDLE_ERROR( cudaMalloc( (void**)&dev_partial_c,
-                              blocksPerGrid*sizeof(float) ) );
-
-    // copy the arrays 'a' and 'b' to the GPU
-    HANDLE_ERROR( cudaMemcpy( dev_a, a, size*sizeof(float),
-                              cudaMemcpyHostToDevice ) );
-    HANDLE_ERROR( cudaMemcpy( dev_b, b, size*sizeof(float),
-                              cudaMemcpyHostToDevice ) ); 
-
-    dot<<<blocksPerGrid,threadsPerBlock>>>( size, dev_a, dev_b,
-                                            dev_partial_c );
-    // copy the array 'c' back from the GPU to the CPU
-    HANDLE_ERROR( cudaMemcpy( partial_c, dev_partial_c,
-                              blocksPerGrid*sizeof(float),
-                              cudaMemcpyDeviceToHost ) );
-
-    // finish up on the CPU side
-    c = 0;
-    for (int i=0; i<blocksPerGrid; i++) {
-        c += partial_c[i];
-    }
-
-    HANDLE_ERROR( cudaFree( dev_a ) );
-    HANDLE_ERROR( cudaFree( dev_b ) );
-    HANDLE_ERROR( cudaFree( dev_partial_c ) );
-
-    // free memory on the CPU side
-    free( partial_c );
-
-    data->returnValue = c;
-    return 0;
-}
-
-
-int main( void ) {
-    int deviceCount;
-    HANDLE_ERROR( cudaGetDeviceCount( &deviceCount ) );
-    if (deviceCount < 2) {
-        printf( "We need at least two compute 1.0 or greater "
-                "devices, but only found %d\n", deviceCount );
-        return 0;
-    }
-
-    float   *a = (float*)malloc( sizeof(float) * N );
-    HANDLE_NULL( a );
-    float   *b = (float*)malloc( sizeof(float) * N );
-    HANDLE_NULL( b );
-
-    // fill in the host memory with data
-    for (int i=0; i<N; i++) {
-        a[i] = i;
-        b[i] = i*2;
-    }
-
-    // prepare for multithread
-    DataStruct  data[2];
-    data[0].deviceID = 0;
-    data[0].size = N/2;
-    data[0].a = a;
-    data[0].b = b;
-
-    data[1].deviceID = 1;
-    data[1].size = N/2;
-    data[1].a = a + N/2;
-    data[1].b = b + N/2;
-
-    CUTThread   thread = start_thread( routine, &(data[0]) );
-    routine( &(data[1]) );
-    end_thread( thread );
-
-
-    // free memory on the CPU side
-    free( a );
-    free( b );
-
-    printf( "Value calculated:  %f\n",
-            data[0].returnValue + data[1].returnValue );
-
-    return 0;
-}
-
diff --git a/chapter11/portable.cu b/chapter11/portable.cu
deleted file mode 100644
index 49609c2..0000000
--- a/chapter11/portable.cu
+++ /dev/null
@@ -1,184 +0,0 @@
-/*
- * Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.
- *
- * NVIDIA Corporation and its licensors retain all intellectual property and 
- * proprietary rights in and to this software and related documentation. 
- * Any use, reproduction, disclosure, or distribution of this software 
- * and related documentation without an express license agreement from
- * NVIDIA Corporation is strictly prohibited.
- *
- * Please refer to the applicable NVIDIA end user license agreement (EULA) 
- * associated with this source code for terms and conditions that govern 
- * your use of this NVIDIA software.
- * 
- */
-
-
-#include "../common/book.h"
-
-
-#define imin(a,b) (a<b?a:b)
-
-#define     N    (33*1024*1024)
-const int threadsPerBlock = 256;
-const int blocksPerGrid =
-            imin( 32, (N/2+threadsPerBlock-1) / threadsPerBlock );
-
-
-__global__ void dot( int size, float *a, float *b, float *c ) {
-    __shared__ float cache[threadsPerBlock];
-    int tid = threadIdx.x + blockIdx.x * blockDim.x;
-    int cacheIndex = threadIdx.x;
-
-    float   temp = 0;
-    while (tid < size) {
-        temp += a[tid] * b[tid];
-        tid += blockDim.x * gridDim.x;
-    }
-    
-    // set the cache values
-    cache[cacheIndex] = temp;
-    
-    // synchronize threads in this block
-    __syncthreads();
-
-    // for reductions, threadsPerBlock must be a power of 2
-    // because of the following code
-    int i = blockDim.x/2;
-    while (i != 0) {
-        if (cacheIndex < i)
-            cache[cacheIndex] += cache[cacheIndex + i];
-        __syncthreads();
-        i /= 2;
-    }
-
-    if (cacheIndex == 0)
-        c[blockIdx.x] = cache[0];
-}
-
-
-struct DataStruct {
-    int     deviceID;
-    int     size;
-    int     offset;
-    float   *a;
-    float   *b;
-    float   returnValue;
-};
-
-
-CUT_THREADPROC  routine( void *pvoidData ) 
-{
-    DataStruct  *data = (DataStruct*)pvoidData;
-    if (data->deviceID != 0) {
-        HANDLE_ERROR( cudaSetDevice( data->deviceID ) );
-        HANDLE_ERROR( cudaSetDeviceFlags( cudaDeviceMapHost ) );
-    }
-
-    int     size = data->size;
-    float   *a, *b, c, *partial_c;
-    float   *dev_a, *dev_b, *dev_partial_c;
-
-    // allocate memory on the CPU side
-    a = data->a;
-    b = data->b;
-    partial_c = (float*)malloc( blocksPerGrid*sizeof(float) );
-
-    // allocate the memory on the GPU
-    HANDLE_ERROR( cudaHostGetDevicePointer( &dev_a, a, 0 ) );
-    HANDLE_ERROR( cudaHostGetDevicePointer( &dev_b, b, 0 ) );
-    HANDLE_ERROR( cudaMalloc( (void**)&dev_partial_c,
-                              blocksPerGrid*sizeof(float) ) );
-
-    // offset 'a' and 'b' to where this GPU is gets it data
-    dev_a += data->offset;
-    dev_b += data->offset;
-
-    dot<<<blocksPerGrid,threadsPerBlock>>>( size, dev_a, dev_b,
-                                            dev_partial_c );
-    // copy the array 'c' back from the GPU to the CPU
-    HANDLE_ERROR( cudaMemcpy( partial_c, dev_partial_c,
-                              blocksPerGrid*sizeof(float),
-                              cudaMemcpyDeviceToHost ) );
-
-    // finish up on the CPU side
-    c = 0;
-    for (int i=0; i<blocksPerGrid; i++) {
-        c += partial_c[i];
-    }
-
-    HANDLE_ERROR( cudaFree( dev_partial_c ) );
-
-    // free memory on the CPU side
-    free( partial_c );
-
-    data->returnValue = c;
-    return 0;
-}
-
-
-int main( void ) {
-    int deviceCount;
-    HANDLE_ERROR( cudaGetDeviceCount( &deviceCount ) );
-    if (deviceCount < 2) {
-        printf( "We need at least two compute 1.0 or greater "
-                "devices, but only found %d\n", deviceCount );
-        return 0;
-    }
-
-    cudaDeviceProp  prop;
-    for (int i=0; i<2; i++) {
-        HANDLE_ERROR( cudaGetDeviceProperties( &prop, i ) );
-        if (prop.canMapHostMemory != 1) {
-            printf( "Device %d can not map memory.\n", i );
-            return 0;
-        }
-    }
-
-    float *a, *b;
-    HANDLE_ERROR( cudaSetDevice( 0 ) );
-    HANDLE_ERROR( cudaSetDeviceFlags( cudaDeviceMapHost ) );
-    HANDLE_ERROR( cudaHostAlloc( (void**)&a, N*sizeof(float),
-                              cudaHostAllocWriteCombined |
-                              cudaHostAllocPortable |
-                              cudaHostAllocMapped ) );
-    HANDLE_ERROR( cudaHostAlloc( (void**)&b, N*sizeof(float),
-                              cudaHostAllocWriteCombined |
-                              cudaHostAllocPortable      |
-                              cudaHostAllocMapped ) );
-
-    // fill in the host memory with data
-    for (int i=0; i<N; i++) {
-        a[i] = i;
-        b[i] = i*2;
-    }
-
-    // prepare for multithread
-    DataStruct  data[2];
-    data[0].deviceID = 0;
-    data[0].offset = 0;
-    data[0].size = N/2;
-    data[0].a = a;
-    data[0].b = b;
-
-    data[1].deviceID = 1;
-    data[1].offset = N/2;
-    data[1].size = N/2;
-    data[1].a = a;
-    data[1].b = b;
-
-    CUTThread   thread = start_thread( routine, &(data[1]) );
-    routine( &(data[0]) );
-    end_thread( thread );
-
-
-    // free memory on the CPU side
-    HANDLE_ERROR( cudaFreeHost( a ) );
-    HANDLE_ERROR( cudaFreeHost( b ) );
-
-    printf( "Value calculated:  %f\n",
-            data[0].returnValue + data[1].returnValue );
-
-    return 0;
-}
-
diff --git a/chapter11/time_zero2.cu b/chapter11/time_zero2.cu
deleted file mode 100644
index 2e1da56..0000000
--- a/chapter11/time_zero2.cu
+++ /dev/null
@@ -1,219 +0,0 @@
-/*
- * Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.
- *
- * NVIDIA Corporation and its licensors retain all intellectual property and 
- * proprietary rights in and to this software and related documentation. 
- * Any use, reproduction, disclosure, or distribution of this software 
- * and related documentation without an express license agreement from
- * NVIDIA Corporation is strictly prohibited.
- *
- * Please refer to the applicable NVIDIA end user license agreement (EULA) 
- * associated with this source code for terms and conditions that govern 
- * your use of this NVIDIA software.
- * 
- */
-
-
-#include "../common/book.h"
-
-#define imin(a,b) (a<b?a:b)
-
-const int N = 33 * 1024 * 1024;
-const int threadsPerBlock = 256;
-const int blocksPerGrid =
-            imin( 32, (N+threadsPerBlock-1) / threadsPerBlock );
-
-
-__global__ void dot( int size, float *a, float *b, float *c ) {
-    __shared__ float cache[threadsPerBlock];
-    int tid = threadIdx.x + blockIdx.x * blockDim.x;
-    int cacheIndex = threadIdx.x;
-
-    float   temp = 0;
-    while (tid < size) {
-        temp += a[tid] * b[tid];
-        tid += blockDim.x * gridDim.x;
-    }
-    
-    // set the cache values
-    cache[cacheIndex] = temp;
-    
-    // synchronize threads in this block
-    __syncthreads();
-
-    // for reductions, threadsPerBlock must be a power of 2
-    // because of the following code
-    int i = blockDim.x/2;
-    while (i != 0) {
-        if (cacheIndex < i)
-            cache[cacheIndex] += cache[cacheIndex + i];
-        __syncthreads();
-        i /= 2;
-    }
-
-    if (cacheIndex == 0)
-        c[blockIdx.x] = cache[0];
-}
-
-
-float malloc_test( int size ) {
-    cudaEvent_t     start, stop;
-    float           *a, *b, c, *partial_c;
-    float           *dev_a, *dev_b, *dev_partial_c;
-    float           elapsedTime;
-
-    HANDLE_ERROR( cudaEventCreate( &start ) );
-    HANDLE_ERROR( cudaEventCreate( &stop ) );
-
-    // allocate memory on the CPU side
-    a = (float*)malloc( size*sizeof(float) );
-    b = (float*)malloc( size*sizeof(float) );
-    partial_c = (float*)malloc( blocksPerGrid*sizeof(float) );
-
-    // allocate the memory on the GPU
-    HANDLE_ERROR( cudaMalloc( (void**)&dev_a,
-                              size*sizeof(float) ) );
-    HANDLE_ERROR( cudaMalloc( (void**)&dev_b,
-                              size*sizeof(float) ) );
-    HANDLE_ERROR( cudaMalloc( (void**)&dev_partial_c,
-                              blocksPerGrid*sizeof(float) ) );
-
-    // fill in the host memory with data
-    for (int i=0; i<size; i++) {
-        a[i] = i;
-        b[i] = i*2;
-    }
-
-    HANDLE_ERROR( cudaEventRecord( start, 0 ) );
-    // copy the arrays 'a' and 'b' to the GPU
-    HANDLE_ERROR( cudaMemcpy( dev_a, a, size*sizeof(float),
-                              cudaMemcpyHostToDevice ) );
-    HANDLE_ERROR( cudaMemcpy( dev_b, b, size*sizeof(float),
-                              cudaMemcpyHostToDevice ) ); 
-
-    dot<<<blocksPerGrid,threadsPerBlock>>>( size, dev_a, dev_b,
-                                            dev_partial_c );
-    // copy the array 'c' back from the GPU to the CPU
-    HANDLE_ERROR( cudaMemcpy( partial_c, dev_partial_c,
-                              blocksPerGrid*sizeof(float),
-                              cudaMemcpyDeviceToHost ) );
-
-    HANDLE_ERROR( cudaEventRecord( stop, 0 ) );
-    HANDLE_ERROR( cudaEventSynchronize( stop ) );
-    HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime,
-                                        start, stop ) );
-
-    // finish up on the CPU side
-    c = 0;
-    for (int i=0; i<blocksPerGrid; i++) {
-        c += partial_c[i];
-    }
-
-    HANDLE_ERROR( cudaFree( dev_a ) );
-    HANDLE_ERROR( cudaFree( dev_b ) );
-    HANDLE_ERROR( cudaFree( dev_partial_c ) );
-
-    // free memory on the CPU side
-    free( a );
-    free( b );
-    free( partial_c );
-
-    // free events
-    HANDLE_ERROR( cudaEventDestroy( start ) );
-    HANDLE_ERROR( cudaEventDestroy( stop ) );
-
-    printf( "Value calculated:  %f\n", c );
-
-    return elapsedTime;
-}
-
-
-float cuda_host_alloc_test( int size ) {
-    cudaEvent_t     start, stop;
-    float           *a, *b, c, *partial_c;
-    float           *dev_a, *dev_b, *dev_partial_c;
-    float           elapsedTime;
-
-    HANDLE_ERROR( cudaEventCreate( &start ) );
-    HANDLE_ERROR( cudaEventCreate( &stop ) );
-
-    // allocate the memory on the CPU
-    HANDLE_ERROR( cudaHostAlloc( (void**)&a,
-                              size*sizeof(float),
-                              cudaHostAllocWriteCombined |
-                                     cudaHostAllocMapped ) );
-    HANDLE_ERROR( cudaHostAlloc( (void**)&b,
-                              size*sizeof(float),
-                              cudaHostAllocWriteCombined |
-                                     cudaHostAllocMapped ) );
-    HANDLE_ERROR( cudaHostAlloc( (void**)&partial_c,
-                              blocksPerGrid*sizeof(float),
-                              cudaHostAllocMapped ) );
-
-    // find out the GPU pointers
-    HANDLE_ERROR( cudaHostGetDevicePointer( &dev_a, a, 0 ) );
-    HANDLE_ERROR( cudaHostGetDevicePointer( &dev_b, b, 0 ) );
-    HANDLE_ERROR( cudaHostGetDevicePointer( &dev_partial_c,
-                                            partial_c, 0 ) );
-
-    // fill in the host memory with data
-    for (int i=0; i<size; i++) {
-        a[i] = i;
-        b[i] = i*2;
-    }
-
-    HANDLE_ERROR( cudaEventRecord( start, 0 ) );
-
-    dot<<<blocksPerGrid,threadsPerBlock>>>( size, dev_a, dev_b,
-                                            dev_partial_c );
-
-    HANDLE_ERROR( cudaThreadSynchronize() );
-    HANDLE_ERROR( cudaEventRecord( stop, 0 ) );
-    HANDLE_ERROR( cudaEventSynchronize( stop ) );
-    HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime,
-                                        start, stop ) );
-
-    // finish up on the CPU side
-    c = 0;
-    for (int i=0; i<blocksPerGrid; i++) {
-        c += partial_c[i];
-    }
-
-    HANDLE_ERROR( cudaFreeHost( a ) );
-    HANDLE_ERROR( cudaFreeHost( b ) );
-    HANDLE_ERROR( cudaFreeHost( partial_c ) );
-
-    // free events
-    HANDLE_ERROR( cudaEventDestroy( start ) );
-    HANDLE_ERROR( cudaEventDestroy( stop ) );
-
-    printf( "Value calculated:  %f\n", c );
-
-    return elapsedTime;
-}
-
-
-int main( void ) {
-    cudaDeviceProp  prop;
-    int whichDevice;
-    HANDLE_ERROR( cudaGetDevice( &whichDevice ) );
-    HANDLE_ERROR( cudaGetDeviceProperties( &prop, whichDevice ) );
-    if (prop.canMapHostMemory != 1) {
-        printf( "Device can not map memory.\n" );
-        return 0;
-    }
-
-    float           elapsedTime;
-
-    HANDLE_ERROR( cudaSetDeviceFlags( cudaDeviceMapHost ) );
-
-    // try it with malloc
-    elapsedTime = malloc_test( N );
-    printf( "Time using cudaMalloc:  %3.1f ms\n",
-            elapsedTime );
-
-    // now try it with cudaHostAlloc
-    elapsedTime = cuda_host_alloc_test( N );
-    printf( "Time using cudaHostAlloc:  %3.1f ms\n",
-            elapsedTime );
-}
diff --git a/common/book.h b/common/book.h
index 065bc0b..ce4a5c4 100644
--- a/common/book.h
+++ b/common/book.h
@@ -148,7 +148,7 @@ __global__ void float_to_color( uchar4 *optr,
 #endif
 
 //Create thread.
-CUTThread start_thread( CUT_THREADROUTINE, void *data );
+CUTThread start_thread( CUT_THREADROUTINE func, void *data );
 
 //Wait for thread to finish.
 void end_thread( CUTThread thread );