Kernels are launched as grids of blocks of threads. Threads can further be divided into 32-thread warps, and each thread in a warp is called a lane.

Round half to even
A tie-breaking rule that is less biased is round half to even, namely: If the fraction of y is 0.5, then q is the even integer nearest to y.
Thus, for example, +23.5 becomes +24, as does +24.5; while −23.5 becomes −24, as does −24.5.
This method treats positive and negative values symmetrically, and is therefore free of sign bias. More importantly, for reasonable distributions of y values, the expected (average) value of the rounded numbers is the same as that of the original numbers. However, this rule will introduce a towards-zero bias when y − 0.5 is even, and a towards-infinity bias for when it is odd.
This variant of the round-to-nearest method is also called unbiased rounding, convergent rounding, statistician's rounding, Dutch rounding, Gaussian rounding, odd–even rounding or bankers' rounding.
This is the default rounding mode used in IEEE 754 computing functions and operators.

A CUDA program is a unified source code encompassing both host and device code. The NVIDIA C compiler (nvcc) separates the two during the compilation process. The host code is straight ANSI C code; it is further compiled with the host’s standard C compilers and runs as an ordinary CPU process. The device code is written using ANSI C extended with keywords for labeling data-parallel functions, called kernels, and their associated data structures. The device code is typically further compiled by the nvcc and executed on a GPU device.

It is worth noting that CUDA threads are of much lighter weight than the CPU threads. CUDA programmers can assume that these threads take very few cycles to generate and schedule due to efficient hardware support. This is in contrast with the CPU threads that typically require thousands of clock cycles to generate and schedule.

The execution starts with host (CPU) execution. When a kernel function is invoked, or launched, the execution is moved to a device (GPU), where a large number of threads are generated to take advantage of abundant data parallelism.
All the threads that are generated by a kernel during an invocation are collectively called a grid.
When all threads of a kernel complete their execution, the corresponding grid terminates, and the execution continues on the host until another kernel is invoked.

In CUDA, the host and devices have separate memory spaces. This reflects the reality that devices are typically hardware cards that come with their own dynamic random access memory (DRAM).
In order to execute a kernel on a device, the programmer needs to allocate memory on the device and transfer pertinent data from the host memory to the allocated device memory.
After device execution, the programmer needs to transfer result data from the device memory back to the host memory and free up the device memory that is no longer needed.

The first parameter of the cudaMalloc() function is the address of a pointer variable that must point to the allocated object after allocation. The address of the pointer variable should be cast to (void **) because the function expects a generic pointer value; the memory allocation function is a generic function that is not restricted to any particular type of objects. This address allows the cudaMalloc() function to write the address of the allocated object into the pointer variable.
The second parameter of the cudaMalloc() function gives the size of the object to be allocated, in terms of bytes.

Because all of these threads execute the same code, CUDA programming is an instance of the well-known single-program, multiple-data (SPMD) parallel programming style.

The __global__ keyword indicates that the function being declared is a CUDA kernel function. The function will be executed on the device and can only be called from the host to generate a grid of threads on a device.
The __device__ keyword indicates that the function being declared is a CUDA device function. A device function executes on a CUDA device and can only be called from a kernel function or another device function. Device functions can have neither recursive function calls nor indirect function calls through pointers in them.
The __host__ keyword indicates that the function being declared is a CUDA host function. A host function is simply a traditional C function that executes on the host and can only be called from another host function. By default, all functions in a CUDA program are host functions if they do not have any of the CUDA keywords in their declaration. This makes sense, as many CUDA applications are ported from CPU-only execution environments. The programmer would add kernel functions and device functions during the porting process. The original functions remain as host functions. Having all functions default into host functions spares the programmer the tedious work of changing all original function declarations.

When a kernel is invoked, or launched, it is executed as grid of parallel threads. Threads in a grid are organized into a two-level hierarchy. In reality, a grid will typically consist of many more threads.
At the top level, each grid consists of one or more thread blocks. All blocks in a grid have the same number of threads.
Each thread block is, in turn, organized as a three-dimensional array of threads with a total size of up to 512 threads. The coordinates of threads in a block are uniquely defined by three thread indices: threadIdx.x,
threadIdx.y, and threadIdx.z. Not all applications will use all three dimensions of a thread block.

The warp is the unit of thread scheduling in SM.

With enough warps around, the hardware will likely find a warp to execute at any point in time, thus making full use of the execution hardware in spite of these long-latency operations. The selection of ready warps for execution does not introduce any idle time into the execution timeline, which is referred to as zero-overhead thread scheduling(零开销的线程调度). With warp scheduling, the long waiting time of warp instructions is hidden by executing instructions from other warps.

Once a block is assigned to an SM, it is further partitioned into warps. At any time, the SM executes only a subset of its resident warps for execution. This allows the other warps to wait for long-latency operations without slowing down the overall execution throughput of the massive number of execution units.

不明觉厉

In CUDA each thread block is actually executed as 32 thread warps. So we are goning to divide each thread block into these 32 thread warps. And execute each warp asd a SIMD unit. So 32 of these CUDA threads will actually be executed in the SIMD execution model, where only one instruction will be fetched and all the 32 of those threads will be execution the same instruction but based on their own data.