8.2.2 Device Level . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 140
8.2.3 Multiprocessor Level . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 140
8.2.3.1 Occupancy Calculator . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 142
8.3 Maximize Memory Throughput . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 144
8.3.1 Data Transfer between Host and Device . . . . . . . . . . . . . . . . . . . . . . . . . . . . 145
8.3.2 Device Memory Accesses . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 146
8.4 Maximize Instruction Throughput . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 149
8.4.1 Arithmetic Instructions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 149
8.4.2 Control Flow Instructions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 156
8.4.3 Synchronization Instruction . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 156
8.5 Minimize Memory Thrashing . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 156
9 CUDA-Enabled GPUs 159
10 C++ Language Extensions 161
10.1 Function Execution Space Speciers . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 161
10.1.1 __global__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 161
10.1.2 __device__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 161
10.1.3 __host__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 162
10.1.4 Undened behavior . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 162
10.1.5 __noinline__ and __forceinline__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 162
10.1.6 __inline_hint__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 163
10.2 Variable Memory Space Speciers . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 163
10.2.1 __device__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 163
10.2.2 __constant__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 163
10.2.3 __shared__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 164
10.2.4 __grid_constant__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 165
10.2.5 __managed__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 165
10.2.6 __restrict__ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 166
10.3 Built-in Vector Types . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 167
10.3.1 char, short, int, long, longlong, oat, double . . . . . . . . . . . . . . . . . . . . . . . . . . 167
10.3.2 dim3 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 169
10.4 Built-in Variables . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 169
10.4.1 gridDim . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 169
10.4.2 blockIdx . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 169
10.4.3 blockDim . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 169
10.4.4 threadIdx . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 169
10.4.5 warpSize . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 169
10.5 Memory Fence Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 170
10.6 Synchronization Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 173
10.7 Mathematical Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 174
10.8 Texture Functions . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 174
10.8.1 Texture Object API . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 174
10.8.1.1 tex1Dfetch() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 174
10.8.1.2 tex1D() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 174
10.8.1.3 tex1DLod() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 174
10.8.1.4 tex1DGrad() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 175
10.8.1.5 tex2D() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 175
10.8.1.6 tex2D() for sparse CUDA arrays . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 175
10.8.1.7 tex2Dgather() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 175
10.8.1.8 tex2Dgather() for sparse CUDA arrays . . . . . . . . . . . . . . . . . . . . . . . . . . 175
10.8.1.9 tex2DGrad() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 176
10.8.1.10 tex2DGrad() for sparse CUDA arrays . . . . . . . . . . . . . . . . . . . . . . . . . . . 176
10.8.1.11 tex2DLod() . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 176
iii