3

データの配列を保持する構造を実装しようとしていますが、次のような動的配列を実装したいと思います。

struct myStruct {
  float3 *data0, *data1;
};

__global__ void kernel(myStruct input) {
  unsigned int N = 2;
  while(someStatements) {
    data0 = new float3[N];
    // do somethings
    N *= 2;
  }
}

CUDAカーネルでこのようなことをするにはどうすればよいですか?

4

1 に答える 1

1

このコードをコンピューティング機能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;
}

注意すべきいくつかのポイント:

  1. このコードは、FermiまたはKeplerGPU上のCUDA4.xまたは5.0でのみコンパイルおよび実行されます
  2. GPUをコンパイルするには、GPUの正しいアーキテクチャをnvccに渡す必要があります(たとえばnvcc -arch=sm_30 -Xptxas="-v" -o dynstruct dynstruct.cu、LinuxでGTX 670用にコンパイルするために使用しました)
  3. サンプルコードは、「gather」カーネルを使用して、ランタイムヒープ内の構造から、ホストAPIがアクセスできる割り当てにデータをコピーして、結果を印刷できるようにします。cudaMemcpyこれは、実行時ヒープメモリ内のアドレスから直接コピーできないことに関して前述した制限の回避策です。これがCUDA5.0で修正されることを期待していましたが、最新のリリース候補にはまだこの制限があります。
于 2012-08-31T16:38:47.990 に答える