CUDA Academy · レッスン

shared配列を宣言する

ブロックごとに使える高速なスクラッチパッドメモリを学びます。

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

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

A Scratchpad Per Block

Every block gets a tiny, fast pool of on-chip memory called shared memory. Think of it as a scratchpad the whole block can write to and read from together.

Why It Is So Fast

Shared memory sits right on the streaming multiprocessor, so it is roughly 100x faster than global memory. Use it to avoid hammering slow off-chip DRAM. 🚀

The shared Keyword

You declare it with the __shared__ qualifier inside a kernel. This one array is then shared by every thread in the block.

__global__ void kern() {
    __shared__ float tile[256];
}

Fixed Size at Compile Time

When you give a size in brackets, that is static shared memory. The compiler must know the size, so it has to be a constant, not a runtime value.

__shared__ int counts[128];

One Copy, Not One Per Thread

This is the key idea: a __shared__ array is created once per block, not once per thread. All threads see the exact same array.

Threads Cooperate Through It

Because every thread sees the same data, shared memory lets threads cooperate. One thread can stash a value and a neighbor can pick it up.

Block Scope, Not Beyond

Its lifetime matches the block. The array is born when the block starts and gone when it ends, so it has block scope only. 🧱

Each Thread Owns a Slot

A common pattern is one slot per thread, indexed by threadIdx.x. Each thread loads its element into shared memory in parallel.

__shared__ float s[256];
s[threadIdx.x] = input[i];

A Tiny but Precious Resource

Shared memory is small, often just 48 to 100 KB per SM. Asking for too much per block reduces how many blocks can run at once.

The Classic Use: Staging Tiles

The most common job is staging a tile of global data so the block can reuse it many times without going back to slow DRAM.

Not Visible to Other Blocks

Remember the boundary: shared memory is private to its block. Two different blocks each get their own separate copy and cannot peek at each other.

Quick Check

Let us check how shared memory is scoped.

Recap

You learned that __shared__ gives each block a fast on-chip scratchpad, created once per block and ideal for staging reusable data. Next: keeping threads in step. 🎯

よくある質問