デバイスメモリ

CUDAのプログラムモデルは各々のメモリを持つホストとデバイスからなるシステムを仮定している。 カーネルはデバイスメモリの外から操作するので、ランタイムはホストメモリとデバイスメモリ間のデータ転送を行う関数だけでなく、デバイスメモリの確保、解放、デバイスメモリ内のコピーを行う関数も提供している。

デバイスメモリは線形メモリ(linear memory)かCUDA配列(CUDA array)のどちらかとして確保される。

CUDA配列はテクスチャフェッチのために最適化された不透明なメモリレイアウトである。 詳しくはテクスチャとサーフェスメモリーにて。

線形メモリは一つの統一されたメモリ空間にアロケートされる。 つまり、例えば2分木やリンクリストで、別々にアロケーションされたものをポインターを通してお互いに参照することが出来る。 アドレス空間の大きさはホストシステム(CPU)と、使っているGPUのcompute capabilityに依存する。

compute capability	x86_64(AMD64)	ARM64
>= 6.0 (Pascal)	up to 47bit	up to 48bit
<= 5.3 (Maxwell)	40bit	40bit

線型メモリのアドレス空間

線形メモリは一般的にcudaMalloc()を使って確保され、cudaFree()で解放される。 ホストメモリとデバイスメモリ間のデータ転送はcudaMemcpy()を用いて行われることが多い。 ベクトル和のコードサンプルでは、ベクトルはホストメモリからデバイスメモリへコピーされなければならない。

#include <cassert>

#include <cuda_runtime.h>

// Device code
__global__ void VecAdd(float* A, float* B, float* C, int N) {
  int i = blockDim.x * blockIdx.x + threadIdx.x;
  if (i < N) {
    C[i] = A[i] + B[i];
  }
}

int main() {
  constexpr int    N    = 1e6;
  constexpr size_t size = N * sizeof(float);

  // Allocate input vectors h_A and h_B in host memory
  float* h_A = (float*)std::malloc(size);
  float* h_B = (float*)std::malloc(size);
  float* h_C = (float*)std::malloc(size);

  // Initialize input vectors
  for (int i = 0; i < N; ++i) {
    h_A[i] = i;
  }

  for (int i = 0; i < N; ++i) {
    h_B[i] = 2 * i;
  }

  // Allocate vectors in device memory
  float* d_A;
  cudaMalloc(&d_A, size);
  float* d_B;
  cudaMalloc(&d_B, size);
  float* d_C;
  cudaMalloc(&d_C, size);

  // Copy vectors from host memory to device memory
  cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
  cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

  // Invoke kernel
  int threadsPerBlock = 256;
  int blocksPerGrid   = (N + threadsPerBlock - 1) / threadsPerBlock;
  VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

  // Copy result from device memory to host memory
  // h_C contains the result in host memory
  cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

  // Check result
  for (int i = 0; i < N; ++i) {
    assert(h_C[i] == 3 * i);
  }

  // Free device memory
  cudaFree(d_A);
  cudaFree(d_B);
  cudaFree(d_C);

  // Free host memory
  std::free(h_A);
  std::free(h_B);
  std::free(h_C);
}

線形メモリはcudaMallocPitchとcudaMalloc3Dを通しても確保出来る。 これらの関数は2次元または3次元配列のメモリ確保がアライメントの要件に合うように適切にパディングされるようにしたいときにおすすめ。 詳しくはデバイスメモリへのアクセスにて。 返されたピッチ(pitch)(またはストライド(stride))は配列要素にアクセスするように使わなければならない。 以下のコードサンプルでは、column * rowの2次元配列をアロケートし、デバイスコード内で配列上をループする方法を示す。

#include <cuda_runtime.h>

__device__ float& at(float* M, std::size_t pitch, int row, int column) {
  return *((float*)((char*)M + row * pitch) + column);
}

__global__ void MatrixAdd(
    float*      A,
    std::size_t pitch_A,
    float*      B,
    std::size_t pitch_B,
    float*      C,
    std::size_t pitch_C,
    int         row,
    int         column
) {
  int i = blockDim.x * blockIdx.x + threadIdx.x;
  int j = blockDim.y * blockIdx.y + threadIdx.y;
  if (i < row and j < column) {
    at(C, pitch_C, i, j) = at(A, pitch_A, i, j) + at(B, pitch_B, i, j);
  }
}

template <class T>
void AllocateMatrix(T** dev_ptr, std::size_t* pitch, int row, int column) {
  auto e = cudaMallocPitch(dev_ptr, pitch, sizeof(T) * column, row);
  assert(e == cudaSuccess);
}

void FreeMatrix(void* dev_ptr) {
  auto e = cudaFree(dev_ptr);
  assert(e == cudaSuccess);
}

以下のコードは、要素数がwidth * height * depthの3次元配列を確保するコードである。

#include <cuda_runtime.h>

__device__ float& at(cudaPitchedPtr devPitchedPtr, int height, int x, int y, int z) {
  char*  devPtr     = (char*)(devPitchedPtr.ptr);
  size_t pitch      = devPitchedPtr.pitch;
  size_t slicePitch = pitch * height;
  float* row        = (float*)(devPtr + z * slicePitch + y * pitch);
  return row[x];
}

__global__ void TensorAdd(
    cudaPitchedPtr A, cudaPitchedPtr B, cudaPitchedPtr C, int width, int height, int depth
) {
  int i = blockDim.x * blockIdx.x + threadIdx.x;
  int j = blockDim.y * blockIdx.y + threadIdx.y;
  int k = blockDim.z * blockIdx.z + threadIdx.z;
  if (i < width and j < height and k < depth) {
    at(C, height, i, j, k) = at(A, height, i, j, k) + at(B, height, i, j, k);
  }
}

void AllocateTensor(cudaPitchedPtr* devPitchedPtr, int width, int height, int depth) {
  auto e = cudaMalloc3D(devPitchedPtr, make_cudaExtent(width * sizeof(float), height, depth));
  assert(e == cudaSuccess);
}

void FreeTensor(void* dev_ptr) {
  auto e = cudaFree(dev_ptr);
  assert(e == cudaSuccess);
}

void CopyTensor(cudaPitchedPtr dst, cudaPitchedPtr src, int width, int height, int depth) {
  cudaMemcpy3DParms myParms = {0};
  myParms.srcPtr            = src;
  myParms.dstPtr            = dst;
  myParms.extent            = make_cudaExtent(width * sizeof(float), height, depth);
  myParms.kind              = cudaMemcpyDefault;

リファレンスマニュアルには、cudaMalloc()で確保された線形メモリやcudaMallocPitch()やcudaMalloc3D()で確保された線形メモリ、CUDA配列、グローバルまたは定数メモリ空間で宣言された変数に対して確保されたメモリ間のコピーに使われる関数が色々ある。

以下のコードは、ランタイムAPIを使った、グローバル変数へアクセスする様々な方法を示す。

cudaGetSymbolAddress()はグローバルメモリ空間で宣言された変数に対して確保されたメモリを指すアドレスを取り出すために使われる。 例はUnderstanding of CUDA's cudaGetSymbolAddressにて。 確保されたメモリのサイズはcudaGetSymbolSize()で得られる。