このコードをコンピューティング機能2.xまたは3、xデバイスのいずれかで、最新バージョンのCUDAを使用して実行する場合、カーネルコードはほぼ正確です。C ++new
演算子は、FermiおよびKeplerハードウェアのCUDA4.xおよび5.0でサポートされています。デバイスのランタイムヒープを使用して割り当てられた、new
または割り当てられたメモリに注意してください。malloc
作成されたコンテキストの存続期間がありますが、現在、CUDAホストAPIから(cudaMemcpy
または同様の方法で)直接アクセスすることはできません。
私はあなたの構造とカーネルを簡単なサンプルコードに変えました。それがどのように機能するかを自分で試してみることができます。
#include <cstdio>
struct myStruct {
float *data;
};
__device__
void fill(float * x, unsigned int n)
{
for(int i=0; i<n; i++) x[i] = (float)i;
}
__global__
void kernel(myStruct *input, const unsigned int imax)
{
for(unsigned int i=0,N=1; i<imax; i++, N*=2) {
float * p = new float[N];
fill(p, N);
input[i].data = p;
}
}
__global__
void kernel2(myStruct *input, float *output, const unsigned int imax)
{
for(unsigned int i=0,N=1; i<imax; i++, N*=2) {
output[i] = input[i].data[N-1];
}
}
inline void gpuAssert(cudaError_t code, char * file, int line, bool Abort=true)
{
if (code != 0) {
fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code),file,line);
if (Abort) exit(code);
}
}
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
int main(void)
{
const unsigned int nvals = 16;
struct myStruct * _s;
float * _f, * f;
gpuErrchk( cudaMalloc((void **)&_s, sizeof(struct myStruct) * size_t(nvals)) );
size_t sz = sizeof(float) * size_t(nvals);
gpuErrchk( cudaMalloc((void **)&_f, sz) );
f = new float[nvals];
kernel<<<1,1>>>(_s, nvals);
gpuErrchk( cudaPeekAtLastError() );
kernel2<<<1,1>>>(_s, _f, nvals);
gpuErrchk( cudaPeekAtLastError() );
gpuErrchk( cudaMemcpy(f, _f, sz, cudaMemcpyDeviceToHost) );
gpuErrchk( cudaDeviceReset() );
for(int i=0; i<nvals; i++) {
fprintf(stdout, "%d %f\n", i, f[i]);
}
return 0;
}
注意すべきいくつかのポイント:
- このコードは、FermiまたはKeplerGPU上のCUDA4.xまたは5.0でのみコンパイルおよび実行されます
- GPUをコンパイルするには、GPUの正しいアーキテクチャをnvccに渡す必要があります(たとえば
nvcc -arch=sm_30 -Xptxas="-v" -o dynstruct dynstruct.cu
、LinuxでGTX 670用にコンパイルするために使用しました)
- サンプルコードは、「gather」カーネルを使用して、ランタイムヒープ内の構造から、ホストAPIがアクセスできる割り当てにデータをコピーして、結果を印刷できるようにします。
cudaMemcpy
これは、実行時ヒープメモリ内のアドレスから直接コピーできないことに関して前述した制限の回避策です。これがCUDA5.0で修正されることを期待していましたが、最新のリリース候補にはまだこの制限があります。