__device__ code to be executed on the device (i.e. GPU) only and only called by GPU code (i.e. __global__ or __device__ code)
Host
__host__ code to be executed on the host (i.e. CPU) only
Global
__global__ declares a kernel entry point that executes on the GPU. It is typically launched from host code with a grid and block configuration; with dynamic parallelism, it can also be launched from GPU code.
Worth noting, __host__ __device__ foo() is not the same as __global__ foo(). The former causes the compiler to generate separate host and device versions, while the latter generates a GPU kernel entry point.
Warp
A group of 32 threads that an SM schedules together for execution
SM
Streaming Multiprocessor
My understanding is that it's a hub for some shared register files (to be allocated dynamically to threads), a shared "L1" cache, and a bunch (~128 ish? Architecture dependent) of ALUs which support executing threads in parallel
Worth noting, multiple warps can occur at the same time on an SM
Notes
GPUs execute simple programs in parallel across massive numbers of threads. The tradeoffs include:
No out-of-order execution in a single thread (AFAIK), i.e. decreased smartness
Worth noting, warps can still be parked while other warps execute
Increased latency
More programmatic decomposition (i.e. via hand or compiler) needed
Before Volta, all 32 threads in a warp shared a single program counter, while an active mask tracked which threads participated in the current instruction. Volta introduced per-thread execution state through Independent Thread Scheduling.
Historically, GPUs became popular because of video games. Turns out they are also very relevant to ML workloads.
Reducing latency can be much more expensive than increasing throughput in terms of power and chip area