__device__ and __host__ Functions
Where each kind of function may run.
__device__ and __host__ Functions is a free CUDA Academy lesson on CoddyKit — lesson 2 of 4. You can read the complete lesson below for free — then practise it hands-on in the browser with a built-in code editor and a 24/7 AI tutor. It is part of the CUDA Academy learning path, one of 4 lessons in the course, and your progress syncs across the web and the CoddyKit app.
Beyond the Kernel
Kernels are entry points, but most code lives in helpers. CUDA gives two more qualifiers to say where a plain function may run. Let us meet them.
The __device__ Qualifier
A __device__ function runs on the GPU and is callable only from other GPU code, like a kernel or another device function.
__device__ int square(int x) {
return x * x;
}Helpers for Kernels
Think of __device__ functions as your GPU toolbox. They keep kernels tidy by holding logic the device threads reuse.
The __host__ Qualifier
The __host__ qualifier means ordinary CPU code. It is the default, so plain C++ functions are already host functions.
__host__ void setup() {
// runs on the CPU
}Default Is Host
Leave off all qualifiers and your function is host-only by default. That is why normal C++ still compiles inside a .cu file.
Both at Once
Combine the two and a function compiles for both sides. nvcc builds one version for the CPU and one for the GPU.
__host__ __device__ int dbl(int x) {
return x + x;
}Why Dual-Compile?
Marking a helper __host__ __device__ lets the same math run in CPU reference code and inside kernels, with zero duplication. ✨
Calling Rules
Host code may not call a __device__ function, and device code may not call a host-only one. The compiler enforces the boundary.
No __global__ from Device
A __device__ function is not launchable. Only __global__ kernels accept the triple-angle-bracket launch from the host.
Inlining for Speed
The compiler often inlines small __device__ functions, so wrapping logic in helpers usually costs nothing at runtime.
Pick the Right Tag
Choose by where it runs: __device__ for GPU helpers, __host__ for CPU code, and both when the logic is shared.
Quick Check
Time to test the device and host qualifiers.
Recap
You learned __device__ runs on the GPU, __host__ runs on the CPU, the default is host, and combining both shares one function across worlds. 🎯
Frequently asked questions
Is the “__device__ and __host__ Functions” lesson free?
Yes — the full text of “__device__ and __host__ Functions” is free to read here on the web, and the CUDA Academy course includes 4 lessons in total. To practise it interactively (a built-in code editor and a 24/7 AI tutor) and unlock the rest of the CUDA Academy course, upgrade to CoddyKit PRO.
What will I learn in “__device__ and __host__ Functions”?
Where each kind of function may run. You practise CUDA Academy with hands-on code you run directly in the browser, and a 24/7 AI tutor answers your questions as you work through the lesson.
Do I need any experience to start CUDA Academy?
No prior experience is required. CUDA Academy on CoddyKit is structured for beginners through advanced learners; this is — lesson 2 of 4, so you can start here or from the beginning and move at your own pace.
How long does the “__device__ and __host__ Functions” lesson take?
Most CoddyKit lessons take about 5–10 minutes. Each one is bite-sized and interactive, so you make steady progress and pick up exactly where you left off across the web and the app.
Can I write and run code in this CUDA Academy lesson?
Yes. Every CUDA Academy lesson includes a built-in code editor, so you write and run real code right in your browser and get instant AI feedback — no local setup required.
All lessons in this course
- The __global__ Function Qualifier
- __device__ and __host__ Functions
- Separate Address Spaces
- The Life of a CUDA Program