CUDA，如何實現在CUDA內核結構的動態數組

我試圖實現保存數據的陣列結構，我想實現動態數組，這樣的：CUDA，如何實現在CUDA內核結構的動態數組

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內核中這樣的東西？

來源

2012-08-31 ratin_sat

如果您要在具有最新版本的CUDA的計算能力2.x或3 x設備上運行此代碼，那麼您的內核代碼幾乎是正確的。 Fermi和Kepler硬件上的CUDA 4.x和5.0支持C++ new運算符。請注意，使用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或Kepler GPU上使用CUDA 4.x或5.0編譯並運行
您必須將GPU的正確架構傳遞給nvcc進行編譯（例如，我使用nvcc -arch=sm_30 -Xptxas="-v" -o dynstruct dynstruct.cu編譯了Linux上的GTX 670）
示例代碼使用「收集」內核將數據從運行時堆中的結構複製到主機API可以訪問的分配，以便可以打印出結果。這是一個解決我之前提到的有關cudaMemcpy不能直接從運行時堆內存中的地址複製的限制。我希望這可能在CUDA 5.0中得到修復，但最新發布的候選版本仍然有此限制。

來源

2012-08-31 16:38:47 talonmies

CUDA，如何實現在CUDA內核結構的動態數組

回答

相關問題