float4でベクトル化ロードを行う
より幅広いトランザクションで帯域幅を拡大
「float4でベクトル化ロードを行う」はCoddyKit上の無料CUDA Academyレッスンです。 これはレッスン3/4です。 下記で完全なレッスンを無料で読むことができます。その後、ブラウザ内の組み込みコードエディタと24時間対応のAIチューターでハンズオン演習できます。 これはCUDA Academy学習パスの一部であり、ウェブとCoddyKitアプリ全体で進捗が同期されます。 CUDA Academyコースには全4レッスンが含まれています。
このレッスンの一部はまだ翻訳されておらず、英語で表示されています。
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. 📦
よくある質問
「float4でベクトル化ロードを行う」レッスンは無料ですか?
はい。「float4でベクトル化ロードを行う」の完全なテキストはこのウェブで無料で読めます。インタラクティブに演習し(組み込みコードエディタと24時間対応のAIチューター)、CUDA Academyコースの残りをアンロックするには、CoddyKit PROにアップグレードしてください。 CUDA Academyコースには全4レッスンが含まれています。
「float4でベクトル化ロードを行う」で何を学びますか?
より幅広いトランザクションで帯域幅を拡大 ブラウザで直接実行するハンズオンコードでCUDA Academyを演習し、24時間対応のAIチューターがレッスンを進める中での質問に答えます。
CUDA Academyを始めるのに経験は必要ですか?
事前経験は必要ありません。CoddyKitのCUDA Academyは初級者から上級者向けに構成されているため、ここから始めるか最初から始めて、自分のペースで進むことができます。 これはレッスン3/4です。
「float4でベクトル化ロードを行う」レッスンにはどのくらい時間がかかりますか?
ほとんどのCoddyKitレッスンは約5~10分かかります。各レッスンはコンパクトでインタラクティブなので、着実に進歩し、ウェブとアプリ全体で正確に前回の場所から再開できます。
このCUDA Academyレッスンでコードを書いて実行できますか?
はい。すべてのCUDA Academyレッスンに組み込みコードエディタが含まれているため、ブラウザでリアルコードを書いて実行し、即座のAIフィードバックを取得できます。ローカル設定は不要です。
このコースのすべてのレッスン
- 命令レベルの並列性
- #pragma unrollでループを展開する
- float4でベクトル化ロードを行う
- レジスタ圧力とスピル