2

Cプログラムに動的に宣言された2D配列があり、その内容をCUDAカーネルに転送してさらに処理したいと考えています。処理したら、Cコードで動的に宣言された2D配列に、CUDAで処理されたデータを入力します。これは静的な2DC配列では実行できますが、動的に宣言されたC配列では実行できません。どんな入力でも大歓迎です！

私は動的配列の動的配列を意味します。私が書いたテストコードは以下の通りです。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <conio.h>
#include <math.h>
#include <stdlib.h>

const int nItt = 10;
const int nP = 5;

__device__ int d_nItt = 10;
__device__ int d_nP = 5;


__global__ void arr_chk(float *d_x_k, float *d_w_k, int row_num)
{

int index = (blockIdx.x * blockDim.x) + threadIdx.x;
int index1 = (row_num * d_nP) + index; 
if ( (index1 >= row_num * d_nP) && (index1 < ((row_num +1)*d_nP)))              //Modifying only one row data pertaining to one particular iteration
{

        d_x_k[index1] = row_num * d_nP;
        d_w_k[index1] = index;
}

}

float **mat_create2(int r, int c)
{
float **dynamicArray;
dynamicArray = (float **) malloc (sizeof (float)*r);
for(int i=0; i<r; i++)
    {
    dynamicArray[i] = (float *) malloc (sizeof (float)*c);
        for(int j= 0; j<c;j++)
        {
            dynamicArray[i][j] = 0;
        }
    }
return dynamicArray;
}

/* Freeing memory - here only number of rows are passed*/
void cleanup2d(float **mat_arr, int x)
{
int i;
for(i=0; i<x; i++)
{
    free(mat_arr[i]);
}
free(mat_arr);
}

int main()
{

//float w_k[nItt][nP]; //Static array declaration - works!
//float x_k[nItt][nP];
// if I uncomment this dynamic declaration and comment the static one, it does not work.....
float **w_k = mat_create2(nItt,nP); 
float **x_k = mat_create2(nItt,nP);
float *d_w_k, *d_x_k;       // Device variables for w_k and x_k
int nblocks, blocksize, nthreads;
for(int i=0;i<nItt;i++)
{
    for(int j=0;j<nP;j++)
    {
        x_k[i][j] = (nP*i);
        w_k[i][j] = j;
    }
}

for(int i=0;i<nItt;i++)
{
    for(int j=0;j<nP;j++)
    {
        printf("x_k[%d][%d] = %f\t",i,j,x_k[i][j]);
        printf("w_k[%d][%d] = %f\n",i,j,w_k[i][j]);
    }
}
int size1 = nItt * nP * sizeof(float);
printf("\nThe array size in memory bytes is: %d\n",size1);
cudaMalloc( (void**)&d_x_k, size1 );
cudaMalloc( (void**)&d_w_k, size1 );

if((nP*nItt)<32)    
{
    blocksize = nP*nItt;
    nblocks = 1;
}
else
{
    blocksize = 32;     // Defines the number of threads running per block. Taken equal to warp size
    nthreads = blocksize;
    nblocks =  ceil(float(nP*nItt) / nthreads);     // Calculated total number of blocks thus required
}

for(int i = 0; i< nItt; i++)
{
    cudaMemcpy( d_x_k, x_k, size1,cudaMemcpyHostToDevice ); //copy of x_k to device
    cudaMemcpy( d_w_k, w_k, size1,cudaMemcpyHostToDevice ); //copy of w_k to device
    arr_chk<<<nblocks, blocksize>>>(d_x_k,d_w_k,i);
    cudaMemcpy( x_k, d_x_k, size1, cudaMemcpyDeviceToHost );
    cudaMemcpy( w_k, d_w_k, size1, cudaMemcpyDeviceToHost );
}
printf("\nVerification after return from gpu\n");
for(int i = 0; i<nItt; i++)
{
    for(int j=0;j<nP;j++)
    {
        printf("x_k[%d][%d] = %f\t",i,j,x_k[i][j]);
        printf("w_k[%d][%d] = %f\n",i,j,w_k[i][j]);
    }
}
cudaFree( d_x_k );
cudaFree( d_w_k );
cleanup2d(x_k,nItt);
cleanup2d(w_k,nItt);
getch();
return 0;
4

1 に答える 1

5

私は動的配列の動的配列を意味します。

まあ、それはまさに問題があるところです。動的配列の動的配列は、配列の各行に1つずつ、互いに素なメモリブロックの束で構成されます（ループのmalloc内側からはっきりとわかります）。したがって、 *を1回呼び出すだけでは、このようなデータ構造をデバイスメモリにコピーすることはできません。代わりに、どちらかを行う必要がありますformat_create2cudaMemcpy

  • また、CUDAで動的配列の動的配列を使用します。これを行うには、基本的に、の代わりにmat_create2を使用して関数を再作成してから、各行を個別にコピーする必要があります。cudaMallocmalloc

  • 今と同じように、CUDAで「タイトな」2Dアレイを使用します（これは、少なくともパフォーマンスの面では良いことです！）。ただし、ホストメモリでdyn-dyn-arraysを使用し続ける場合は、次のように各行を個別にコピーする必要があります。

    for(int i=0; i<r; ++i){
  cudaMemcpy(d_x_k + i*c, x_k[i], c*sizeof(float), cudaMemcpyHostToDevice)
}

「なぜ静的な2D配列で機能したのか」と疑問に思うかもしれません。さて、 Cの静的2D配列は、一度にコピーできる適切でタイトな配列です。実際には完全に異なる動作をするため、これらがdyn-dyn配列（）とまったく同じ構文でインデックス付けされているのは少し混乱します。arr[x][y]

ただし、ホストメモリでタイトな配列を使用することも検討する必要があります。

typedef struct {
  float* data;
  int n_rows, n_cols;
} tight2dFloatArray;

#define INDEX_TIGHT2DARRAY(arr, y, x)\
  (arr).data[(y)*(arr).n_cols + (x)]

もちろん、このようなアプローチは、C++クラスとしてはるかに安全に実装できます。

*memcpyまた、実際のデータではなく、ポインタの配列のみをコピーする1つだけでメインメモリ内にコピーすることはできません。

于 2012-10-13T10:52:05.523 に答える