Vectorized Loads with float4
Wider transactions for more bandwidth.
Vectorized Loads with float4 is a free CUDA Academy lesson on CoddyKit — lesson 3 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.
Move More Per Instruction
A normal load fetches one value at a time. A vectorized load grabs several adjacent values in a single wider instruction, doing more work per request.
Meet float4
CUDA ships built-in vector types. A float4 packs four floats into one 16-byte bundle you can load or store together.
float4 v = make_float4(1, 2, 3, 4);
float first = v.x; // also .y .z .wOne Load, Four Floats
Reading a float4 from memory pulls four floats in a single transaction. That cuts the number of memory instructions a thread must issue by four.
float4 a = reinterpret_cast<float4*>(in)[i];Reinterpret the Pointer
To use vector loads on a plain float array, you cast the pointer. reinterpret_cast lets you read the same bytes as float4 elements.
float4* in4 = reinterpret_cast<float4*>(in);
float4 chunk = in4[idx];Alignment Is Required
A float4 must sit on a 16-byte boundary. If the data is not properly aligned, the load is illegal and your kernel can crash or read garbage.
cudaMalloc Aligns for You
Good news: buffers from cudaMalloc are aligned to at least 256 bytes. So freshly allocated device arrays are safe to read as float4 from the start.
Fewer Requests, More Bandwidth
Wider loads mean fewer outstanding requests for the same data. That can raise achieved bandwidth when a kernel is limited by memory, not math.
Index by Vectors, Not Floats
With float4 each thread now covers four elements. Your global index steps through float4 slots, so the total thread count drops by four.
int i = blockIdx.x * blockDim.x + threadIdx.x;
float4 d = in4[i]; // covers 4 floatsOther Vector Widths
float4 is not the only choice. Types like float2 and int4 give you 8- or 16-byte loads, so you can pick a width that fits your data.
Watch the Leftovers
If your array length is not a multiple of four, handle the extra remainder elements with a normal scalar loop after the vectorized part.
It Costs Registers
A float4 holds four values, so it uses more registers per thread. As always, confirm the bandwidth gain outweighs any drop in occupancy.
Quick Check
You reinterpret a float array as float4 but the kernel crashes. Why?
Recap: Wider Is Faster
You learned that float4 loads four floats at once, cutting instructions and lifting bandwidth. Keep data aligned and handle leftover elements. 📦
Frequently asked questions
Is the “Vectorized Loads with float4” lesson free?
Yes — the full text of “Vectorized Loads with float4” 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 “Vectorized Loads with float4”?
Wider transactions for more bandwidth. 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 3 of 4, so you can start here or from the beginning and move at your own pace.
How long does the “Vectorized Loads with float4” 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
- Instruction-Level Parallelism
- Loop Unrolling with #pragma unroll
- Vectorized Loads with float4
- Register Pressure and Spills