ここでの最初の問題はx_dev
、デバイス シンボルではないことです。デバイス メモリ内のアドレスが含まれている可能性がありますが、そのアドレスを に渡すことはできませんcudaMemcpyToSymbol
。呼び出しは次のようにする必要があります。
cudaMemcpyToSymbol(xdev, ......);
これにより、2番目の問題が発生します。これを行う:
cudaMemcpyToSymbol(xdev, x, sizeof(collapsed)*size);
違法だろう。xdev
はポインタであるため、コピーできる唯一の有効な値xdev
はデバイス アドレスです。x
がデバイス メモリ内の のアドレスである 場合struct collapsed
、このメモリ転送操作の唯一の有効なバージョンは
cudaMemcpyToSymbol(xdev, &x, sizeof(collapsed *));
すなわち。x
デバイスに割り当てられたメモリのアドレスに以前に設定されている必要があります。
collapsed *x;
cudaMalloc((void **)&x, sizeof(collapsed)*size);
cudaMemcpy(x, host_src, sizeof(collapsed)*size, cudaMemcpyHostToDevice);
約束どおり、ここに完全な動作例があります。最初のコード:
#include <cstdlib>
#include <iostream>
#include <cuda_runtime.h>
struct collapsed {
char **seq;
int num;
};
__device__ collapsed xdev;
__global__
void kernel(const size_t item_sz)
{
if (threadIdx.x < xdev.num) {
char *p = xdev.seq[threadIdx.x];
char val = 0x30 + threadIdx.x;
for(size_t i=0; i<item_sz; i++) {
p[i] = val;
}
}
}
#define gpuQ(ans) { gpu_assert((ans), __FILE__, __LINE__); }
void gpu_assert(cudaError_t code, const char *file, const int line)
{
if (code != cudaSuccess)
{
std::cerr << "gpu_assert: " << cudaGetErrorString(code) << " "
<< file << " " << line << std::endl;
exit(code);
}
}
int main(void)
{
const int nitems = 32;
const size_t item_sz = 16;
const size_t buf_sz = size_t(nitems) * item_sz;
// Gpu memory for sequences
char *_buf;
gpuQ( cudaMalloc((void **)&_buf, buf_sz) );
gpuQ( cudaMemset(_buf, 0x7a, buf_sz) );
// Host array for holding sequence device pointers
char **seq = new char*[nitems];
size_t offset = 0;
for(int i=0; i<nitems; i++, offset += item_sz) {
seq[i] = _buf + offset;
}
// Device array holding sequence pointers
char **_seq;
size_t seq_sz = sizeof(char*) * size_t(nitems);
gpuQ( cudaMalloc((void **)&_seq, seq_sz) );
gpuQ( cudaMemcpy(_seq, seq, seq_sz, cudaMemcpyHostToDevice) );
// Host copy of the xdev structure to copy to the device
collapsed xdev_host;
xdev_host.num = nitems;
xdev_host.seq = _seq;
// Copy to device symbol
gpuQ( cudaMemcpyToSymbol(xdev, &xdev_host, sizeof(collapsed)) );
// Run Kernel
kernel<<<1,nitems>>>(item_sz);
// Copy back buffer
char *buf = new char[buf_sz];
gpuQ( cudaMemcpy(buf, _buf, buf_sz, cudaMemcpyDeviceToHost) );
// Print out seq values
// Each string should be ASCII starting from ´0´ (0x30)
char *seq_vals = buf;
for(int i=0; i<nitems; i++, seq_vals += item_sz) {
std::string s;
s.append(seq_vals, item_sz);
std::cout << s << std::endl;
}
return 0;
}
そして、ここでコンパイルして実行します:
$ /usr/local/cuda/bin/nvcc -arch=sm_12 -Xptxas=-v -g -G -o erogol erogol.cu
./erogol.cu(19): Warning: Cannot tell what pointer points to, assuming global memory space
ptxas info : 8 bytes gmem, 4 bytes cmem[14]
ptxas info : Compiling entry function '_Z6kernelm' for 'sm_12'
ptxas info : Used 5 registers, 20 bytes smem, 4 bytes cmem[1]
$ /usr/local/cuda/bin/cuda-memcheck ./erogol
========= CUDA-MEMCHECK
0000000000000000
1111111111111111
2222222222222222
3333333333333333
4444444444444444
5555555555555555
6666666666666666
7777777777777777
8888888888888888
9999999999999999
::::::::::::::::
;;;;;;;;;;;;;;;;
<<<<<<<<<<<<<<<<
================
>>>>>>>>>>>>>>>>
????????????????
@@@@@@@@@@@@@@@@
AAAAAAAAAAAAAAAA
BBBBBBBBBBBBBBBB
CCCCCCCCCCCCCCCC
DDDDDDDDDDDDDDDD
EEEEEEEEEEEEEEEE
FFFFFFFFFFFFFFFF
GGGGGGGGGGGGGGGG
HHHHHHHHHHHHHHHH
IIIIIIIIIIIIIIII
JJJJJJJJJJJJJJJJ
KKKKKKKKKKKKKKKK
LLLLLLLLLLLLLLLL
MMMMMMMMMMMMMMMM
NNNNNNNNNNNNNNNN
OOOOOOOOOOOOOOOO
========= ERROR SUMMARY: 0 errors
いくつかのメモ:
- 少し単純化するために、
_buf
すべての文字列データを保持するために 1 つのメモリ割り当てのみを使用しました。の各値はseq
、 内の異なるアドレスに設定されます_buf
。cudaMalloc
これは、ポインタごとに個別の呼び出しを実行するのと機能的に同等ですが、はるかに高速です。
- 重要な概念は、アクセスしたい構造のコピーをホスト メモリ内のデバイスにアセンブルし、それをデバイスにコピーすることです。my のすべてのポインター
xdev_host
はデバイスポインターです。CUDA API には、ディープ コピーや自動ポインター変換機能がまったくないため、これが正しいことを確認するのはプログラマーの責任です。
- カーネル内の各スレッドは、そのシーケンスを異なる ASCII 文字で埋めるだけです。xdev を構造体へのポインターではなく構造体として宣言し、
__device__
シンボルへの参照ではなく値をコピーしたことに注意してください (これも少し単純化するためです)。しかし、それ以外の場合は、設計パターンを機能させるために一連の操作が必要になります。
- コンピューティング 1.x デバイスにしかアクセスできないため、コンパイラは警告を発行します。One Compute 2.x と 3.x では、これらのデバイスのメモリ モデルが改善されているため、これは発生しません。警告は正常であり、安全に無視できます。
- 各シーケンスは の異なる部分に書き込まれるだけなので
_buf
、1 回の cudaMemcpy 呼び出しですべてのシーケンスをホストに戻すことができます。