JetsonでCUDAやるなら統合メモリが幸せかと思ったらそれは幻想だったのかもしれない

この記事はGROOVE X Advent Calendar 2025の21日目の記事です。

こんにちは、「あず」こと斎藤@aznhe21です。肩掛けスピーカーのSRS-NB10が内部で断線したのでBravia Theatre Uに乗り換えたんですが、低音が激しくて新しい体験でした。首の部分をぐにゃぐにゃ曲げても断線する心配がなさそうなのも良きです。

さて、LOVOT 3.0ではJetson Orinを採用しており、内部ではCUDAも積極的に使用しています。ここでCUDAの統合メモリが便利だったのでご紹介したいと思います（罠と共に）。

JetsonでCUDAやるなら統合メモリが幸せかと思ったらそれは幻想だったのかもしれない

統合メモリとは

CUDAはGPU用のプログラミングモデルであり、基本的にはCPUとGPUでメモリを別途管理する必要があります。というのも一般的なdGPU環境ではCPUとGPUでは物理的なメモリ領域が分かれており、明示的に管理する方が自然なのです。

そこでCUDAの統合メモリ（Unified Memory）を使うとCPUとGPUで同一のポインタを使ってメモリにアクセスできてとても便利です。ただ実際の動作としてはドライバが暗黙的にメモリをコピーする必要があり、明示的に管理するよりもオーバーヘッドが発生します。

ここで、JetsonはCPUとGPUが物理的にメモリを共有しています。そのため、Jetsonでは統合メモリを使うことで物理的なコピーを手軽に省けます。ただしこれがパフォーマンス上有利かと言うとそれはまた別のお話・・・。

ちなみにJetsonのようにCPUとGPUで物理メモリを共有しているアーキテクチャのことをUMA（Unified Memory Architecture）と言います。これはCUDAの統合メモリ（Unified Memory）とは別の概念なので気を付けましょう。 CUDAの統合メモリはソフトウェアの、UMAはハードウェアの話と言えます。

実装

統合メモリが何であるか分かったところで実装を見てみましょう。

シンプルな例

通常のCUDAではGPUで使うメモリをcudaMalloc等で確保しますが、統合メモリではcudaMallocManagedで確保します。

um.cu

#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>

const size_t size = 16;

__global__ void kernel() {}

#define CHECK_CUDA(call)                                                       \
  do {                                                                         \
    cudaError_t err = call;                                                    \
    if (err != cudaSuccess) {                                                  \
      fprintf(stderr, "CUDA error in %s (%s:%d): %s\n", #call, __FILE__,       \
              __LINE__, cudaGetErrorString(err));                              \
      exit(EXIT_FAILURE);                                                      \
    }                                                                          \
  } while (0)

int main() {
  CHECK_CUDA(cudaSetDevice(0));

  cudaStream_t stream;
  CHECK_CUDA(cudaStreamCreate(&stream));

  float *ptr = nullptr;
  CHECK_CUDA(cudaMallocManaged((void **)&ptr, size * sizeof(float)));

  const dim3 threads_per_block(1);
  const dim3 blocks(1, 1);

  kernel<<<blocks, threads_per_block, 0, stream>>>();
  cudaError_t err = cudaGetLastError();
  if (err != cudaSuccess) {
    fprintf(stderr, "Kernel launch error: %s\n", cudaGetErrorString(err));
    exit(EXIT_FAILURE);
  }

  // このcudaStreamSynchronizeをコメントアウトすると・・・？
  CHECK_CUDA(cudaStreamSynchronize(stream));

  for (size_t i = 0; i < size; i++) {
    // volatileで最適化を無効化
    volatile float _val = ptr[i];
  }

  CHECK_CUDA(cudaFree(ptr));
  CHECK_CUDA(cudaStreamDestroy(stream));
}

この例ではCPUとGPUで同じポインタを通じてメモリにアクセスしていますが、統合メモリには罠があります。

CUDA実行中にCPUからメモリにアクセスできません（Jetsonでは）。

・・・と書くと強すぎで、より厳密には書くならこうでしょう。

cudaDevAttrConcurrentManagedAccessが0の環境（主にJetsonとWSL環境）において、ある統合メモリに任意のCUDA Streamがアクセスできる状態でCUDAカーネルが実行されている場合、その統合メモリに対するCPUアクセスでセグフォ（Segmentation fault）が発生します。

上記サンプルのコメントにある通り、cudaStreamSynchronizeをコメントアウトしてJetsonで実行するとセグフォが発生します。カーネルではメモリにアクセスしていないにも関わらず、です。メモリがGPUから使われているかどうかは関係なく、Streamからメモリにアクセスし得ることが問題なのです。

このように単純な例ならばcudaStreamSynchronizeを忘れないようにすれば問題ないのですが、実際のプログラムのように複数スレッド・複数StreamでCUDAを同時に実行すると問題が発生します。

複数スレッド・複数Streamでは問題が起きる

um-mt-issue.cu

#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>
#include <thread>
#include <vector>

__global__ void kernel() {}

#define CHECK_CUDA(call)                                                       \
  do {                                                                         \
    cudaError_t err = call;                                                    \
    if (err != cudaSuccess) {                                                  \
      fprintf(stderr, "CUDA error in %s (%s:%d): %s\n", #call, __FILE__,       \
              __LINE__, cudaGetErrorString(err));                              \
      exit(EXIT_FAILURE);                                                      \
    }                                                                          \
  } while (0)

void proc() {
  const size_t size = 100;

  CHECK_CUDA(cudaSetDevice(0));

  cudaStream_t stream;
  CHECK_CUDA(cudaStreamCreate(&stream));

  float *ptr = nullptr;
  CHECK_CUDA(cudaMallocManaged((void **)&ptr, size * sizeof(float)));

  for (int i = 0; i < 100; i++) {
    const dim3 threads_per_block(1);
    const dim3 blocks(1, 1);

    kernel<<<blocks, threads_per_block, 0, stream>>>();
    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess) {
      fprintf(stderr, "Kernel launch error: %s\n", cudaGetErrorString(err));
      exit(EXIT_FAILURE);
    }

    CHECK_CUDA(cudaStreamSynchronize(stream));

    for (size_t j = 0; j < size; j++) {
      // volatileで最適化を無効化
      volatile float _val = ptr[j];
    }
  }

  CHECK_CUDA(cudaFree(ptr));
  CHECK_CUDA(cudaStreamDestroy(stream));
}

int main() {
  std::vector<std::thread> threads;

  for (int i = 0; i < 10; i++) {
    threads.emplace_back(proc);
  }

  for (auto &t : threads) {
    t.join();
  }
}

この例では10個のスレッドで各々のStreamを生成し、カーネルを起動・同期を取った上で統合メモリにアクセスしています。一見バグは無いように見えますが、Jetsonで実行してみるとセグフォが発生します。

これはcudaMallocManagedで確保したメモリはどのStreamからもアクセスできることから、他スレッドでのカーネル起動により統合メモリがGPUから使われていると見なされ、 CPUからのアクセス時にセグフォが発生しているのです。

すなわちスレッド1でカーネルを起動・同期を取っても、スレッド2でカーネルを実行中の場合、最初のサンプル同様同期が取れていないのと同じ意味になるのです。

そこで、統合メモリをStreamと紐付けることで、別のStreamが実行中であってもCPUからアクセスできるようにします。

統合メモリとStreamを紐付ける

cudaStreamAttachMemAsyncにより統合メモリとStreamを紐付けることができ、そのStreamと同期が取れていれば他のStreamの状態に関わらずCPUからアクセスができるようになります。

意図をより明確化するため、cudaMallocManagedにcudaMemAttachHostを渡しておくと良いでしょう。

um-mt-fixed.cu

#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>
#include <thread>
#include <vector>

__global__ void kernel() {}

#define CHECK_CUDA(call)                                                       \
  do {                                                                         \
    cudaError_t err = call;                                                    \
    if (err != cudaSuccess) {                                                  \
      fprintf(stderr, "CUDA error in %s (%s:%d): %s\n", #call, __FILE__,       \
              __LINE__, cudaGetErrorString(err));                              \
      exit(EXIT_FAILURE);                                                      \
    }                                                                          \
  } while (0)

void proc() {
  const size_t size = 100;

  CHECK_CUDA(cudaSetDevice(0));

  cudaStream_t stream;
  CHECK_CUDA(cudaStreamCreate(&stream));

  float *ptr = nullptr;
  // cudaMemAttachHostによりどのStreamにも見せない
  CHECK_CUDA(cudaMallocManaged((void **)&ptr, size * sizeof(float), cudaMemAttachHost));
  // このstreamにのみ紐付け
  CHECK_CUDA(cudaStreamAttachMemAsync(stream, ptr, 0, cudaMemAttachSingle));

  for (int i = 0; i < 100; i++) {
    const dim3 threads_per_block(1);
    const dim3 blocks(1, 1);

    kernel<<<blocks, threads_per_block, 0, stream>>>();
    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess) {
      fprintf(stderr, "Kernel launch error: %s\n", cudaGetErrorString(err));
      exit(EXIT_FAILURE);
    }

    CHECK_CUDA(cudaStreamSynchronize(stream));

    for (size_t j = 0; j < size; j++) {
      // volatileで最適化を無効化
      volatile float _val = ptr[j];
    }
  }

  CHECK_CUDA(cudaFree(ptr));
  CHECK_CUDA(cudaStreamDestroy(stream));
}

int main() {
  std::vector<std::thread> threads;

  for (int i = 0; i < 10; i++) {
    threads.emplace_back(proc);
  }

  for (auto &t : threads) {
    t.join();
  }
}

この記事を書いて気付いた幻想

この調査をした際はベンチマークを書いてパフォーマンスが悪化していないことは確認しており、コードがシンプルになったので満足していました。ただそのベンチマークは概ねCUDAカーネルを起動するだけのものであり、CPUからメモリにアクセスしていないなど、実態に沿っていませんでした。

今回記事を書くに当たって裏取りをしていたところ、どうやら統合メモリではパフォーマンスが悪化することがあるようです。というのも、GPUで書き込んだメモリにCPUからアクセスする場合はキャッシュコヒーレンシが維持できず、 CPUキャッシュが無効化されてしまうようなのです。つまり物理的なコピーの削減によって速度が向上しても、CPUキャッシュが効かないことによって速度が低減し、却って遅くなることがある、ということのようです。とは言えパフォーマンスへの影響がないことも多いようなので、実装・保守コストとのバランス次第でしょう（文献を読んだだけで実動作は未確認なので「ようです」ばかりです）。

というわけで、CPUアクセスも含めると全体のパフォーマンスは悪化していた可能性があります。もちろんメトリクスで大きな問題がないことは確認しているとは言え、今回のことでマイクロベンチマークの幻想に気付かされました。

皆さんは最適化・リファクタリングするならちゃんと実際の動作で計測して比較しましょうね。

さいごに

GROOVE XではCUDAをカリカリにチューニングしたい画像認識エンジニアを募集しています。

Inside of LOVOT

GROOVE X 技術ブログ