CUDA Academy · レッスン

コンスタントメモリとそのキャッシュ

読み取り専用の値を低コストでブロードキャストします。

「コンスタントメモリとそのキャッシュ」はCoddyKit上の無料CUDA Academyレッスンです。これはレッスン3/4です。下記で完全なレッスンを無料で読むことができます。その後、ブラウザ内の組み込みコードエディタと24時間対応のAIチューターでハンズオン演習できます。これはCUDA Academy学習パスの一部であり、ウェブとCoddyKitアプリ全体で進捗が同期されます。 CUDA Academyコースには全4レッスンが含まれています。

このレッスンの一部はまだ翻訳されておらず、英語で表示されています。

A Space for Read-Only Values

Constant memory is a small region built for data that every thread reads but none of them ever writes during a kernel.

Declaring It

You mark a global-scope array with the __constant__ qualifier. It lives outside any function, visible to all your kernels.

__constant__ float weights[256];

It Is Genuinely Small

Total constant memory is only 64 KB on typical GPUs. It is meant for coefficients and parameters, not for large datasets.

Filling It from the Host

You write to constant memory from the CPU using cudaMemcpyToSymbol, since kernels themselves cannot modify it.

cudaMemcpyToSymbol(weights, h_w,
    256 * sizeof(float));

Backed by a Special Cache

Reads flow through a dedicated constant cache on each multiprocessor. A cached value comes back almost as fast as a register.

The Broadcast Superpower

When a whole warp reads the same address, the hardware does one fetch and broadcasts it to all 32 threads in a single step. 📡

Same Address Is the Key

The win depends on uniform access. If threads in a warp read different constant addresses, those reads serialize and the speedup is lost.

Reading It Looks Normal

Inside a kernel you read constant memory like any array. The compiler quietly routes the access through the fast constant cache.

__global__ void k(float *out, int i) {
    out[i] = weights[0] * 2.0f;
}

A Perfect Fit for Filters

Convolution kernels and lookup tables shine here, because every thread reuses the same small set of coefficients over and over.

Set Once, Reuse Often

You typically upload to constant memory a single time, then launch many kernels that all read those unchanging values cheaply.

When Not to Use It

Skip constant memory if your data is large or if threads read scattered addresses. In those cases plain global memory serves you better.

Quick Check

When does constant memory give its biggest speedup?

Recap: Small, Cached, Broadcast

You learned constant memory is a tiny read-only space whose cache broadcasts uniform reads to a whole warp. Use it for small, shared values. ✨

よくある質問

「コンスタントメモリとそのキャッシュ」レッスンは無料ですか？

はい。「コンスタントメモリとそのキャッシュ」の完全なテキストはこのウェブで無料で読めます。インタラクティブに演習し（組み込みコードエディタと24時間対応のAIチューター）、CUDA Academyコースの残りをアンロックするには、CoddyKit PROにアップグレードしてください。 CUDA Academyコースには全4レッスンが含まれています。

「コンスタントメモリとそのキャッシュ」で何を学びますか？

読み取り専用の値を低コストでブロードキャストします。ブラウザで直接実行するハンズオンコードでCUDA Academyを演習し、24時間対応のAIチューターがレッスンを進める中での質問に答えます。

CUDA Academyを始めるのに経験は必要ですか？

事前経験は必要ありません。CoddyKitのCUDA Academyは初級者から上級者向けに構成されているため、ここから始めるか最初から始めて、自分のペースで進むことができます。これはレッスン3/4です。

「コンスタントメモリとそのキャッシュ」レッスンにはどのくらい時間がかかりますか？

ほとんどのCoddyKitレッスンは約5～10分かかります。各レッスンはコンパクトでインタラクティブなので、着実に進歩し、ウェブとアプリ全体で正確に前回の場所から再開できます。

このCUDA Academyレッスンでコードを書いて実行できますか？

はい。すべてのCUDA Academyレッスンに組み込みコードエディタが含まれているため、ブラウザでリアルコードを書いて実行し、即座のAIフィードバックを取得できます。ローカル設定は不要です。

このコースのすべてのレッスン

← CUDA Academyに戻る