3

CUDA で 3D float 配列を使用したいのですが、コードは次のとおりです。

#define  SIZE_X 128 //numbers in elements
#define  SIZE_Y 128
#define  SIZE_Z 128
typedef float  VolumeType;
cudaExtent volumeSize = make_cudaExtent(SIZE_X, SIZE_Y, SIZE_Z); //The first argument should be SIZE_X*sizeof(VolumeType)??

float *d_volumeMem;
cutilSafeCall(cudaMalloc((void**)&d_volumeMem, SIZE_X*SIZE_Y*SIZE_Z*sizeof(float)));

.....//assign value to d_volumeMem in GPU

cudaArray *d_volumeArray = 0;
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<VolumeType>();
cutilSafeCall( cudaMalloc3DArray(&d_volumeArray, &channelDesc, volumeSize) ); 
cudaMemcpy3DParms copyParams = {0};
copyParams.srcPtr = make_cudaPitchedPtr((void*)d_volumeMem, SIZE_X*sizeof(VolumeType), SIZE_X, SIZE_Y); //
copyParams.dstArray = d_volumeArray;
copyParams.extent = volumeSize;
copyParams.kin = cudaMemcpyDeviceToDevice;
cutilSafeCall( cudaMemcpy3D(&copyParams) );

実際、私のプログラムはうまく動作します。しかし、結果が正しいかどうかはわかりません。ここに私の問題があります.CUDA liberayでは、make_cudaExtentの最初のパラメータは「バイト単位の幅」であり、他の2つは要素の高さと深さです。上記のコードでは、5行目は

cudaExtent volumeSize = make_cudaExtent(SIZE_X*sizeof(VolumeType), SIZE_Y, SIZE_Z);

しかし、この方法では、cutilSafeCall( cudaMemcpy3D(©Params) ); で「引数が無効です」というエラーが発生します。なんで？

もう 1 つのパズルは、CUDA ライブラリが述べているように、strcut cudaExtent です。そのコンポーネントの幅は、「配列メモリを参照する場合は要素の幅、線形メモリを参照する場合はバイト単位」を表します。したがって、コードで volumeSize.width を参照するときは、要素の数値である必要があると思います。ただし、使用する場合

 cudaExtent volumeSize = make_cudaExtent(SIZE_X*sizeof(VolumeType), SIZE_Y, SIZE_Z);

volumeSize.width は SIZE_X*sizeof(VolumeType)(128*4) になります。これは、要素数ではなくバイト数です。

多くの CUDA SDK では、VolumeType として char を使用するため、make_cudaExtent の最初の引数として SIZE_X を使用するだけです。しかし、私のものはフロートなので、これを使用して 3D 配列を作成する必要がある場合、cudaExtent を作成する正しい方法はどれか教えてもらえますか?? どうもありがとう！

4

2 に答える 2

3

ドキュメントの内容を確認しましょうcudaMemcpy3D。

extent フィールドは、転送された領域の寸法を要素で定義します。CUDA 配列がコピーに参加している場合、範囲はその配列の要素に関して定義されます。CUDA 配列がコピーに参加していない場合、エクステントは unsigned char の要素で定義されます。

同様に、cudaMalloc3DArrayノートのドキュメント:

すべての値は要素で指定されます

したがって、両方の呼び出しで形成する必要があるエクステントは、要素の最初の次元を持つ必要があります (の割り当ての 1 つがcudaMemcpy3D配列であるため)。

d_volumeMemただし、を使用して線形メモリ ソースを割り当てているため、コードに別の問題が発生する可能性がありますcudaMalloc。cudaMemcpy3Dは、リニア ソース メモリが互換性のあるピッチで割り当てられていることを期待しています。あなたのコードは、サイズの線形割り当てを使用しているだけです

SIZE_X*SIZE_Y*SIZE_Z*sizeof(float)

選択した寸法が、使用しているハードウェアと互換性のあるピッチを生成する可能性がありますが、そうであるとは限りません。cudaMalloc3D線形ソース メモリの割り当てにも使用することをお勧めします。小さなコード スニペットを中心に構築されたこの拡張デモは、次のようになります。

#include <cstdio>

typedef float  VolumeType;

const size_t SIZE_X = 8;
const size_t SIZE_Y = 8;
const size_t SIZE_Z = 8;
const size_t width = sizeof(VolumeType) * SIZE_X;

texture<VolumeType, cudaTextureType3D, cudaReadModeElementType> tex; 

__global__ void testKernel(VolumeType * output, int dimx, int dimy, int dimz)
{
    int tidx = threadIdx.x + blockIdx.x * blockDim.x;
    int tidy = threadIdx.y + blockIdx.y * blockDim.y;
    int tidz = threadIdx.z + blockIdx.z * blockDim.z;

    float x = float(tidx)+0.5f;
    float y = float(tidy)+0.5f;
    float z = float(tidz)+0.5f;

    size_t oidx = tidx + tidy*dimx + tidz*dimx*dimy;
    output[oidx] = tex3D(tex, x, y, z);
}

inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }

template<typename T>
void init(char * devPtr, size_t pitch, int width, int height, int depth)
{
    size_t slicePitch = pitch * height;
    int v = 0;
    for (int z = 0; z < depth; ++z) {
        char * slice = devPtr + z * slicePitch;
        for (int y = 0; y < height; ++y) {
            T * row = (T *)(slice + y * pitch);
            for (int x = 0; x < width; ++x) {
                row[x] = T(v++);
            }
        }
    }
}

int main(void)
{
    VolumeType *h_volumeMem, *d_output, *h_output;

    cudaExtent volumeSizeBytes = make_cudaExtent(width, SIZE_Y, SIZE_Z);
    cudaPitchedPtr d_volumeMem; 
    gpuErrchk(cudaMalloc3D(&d_volumeMem, volumeSizeBytes));

    size_t size = d_volumeMem.pitch * SIZE_Y * SIZE_Z;
    h_volumeMem = (VolumeType *)malloc(size);
    init<VolumeType>((char *)h_volumeMem, d_volumeMem.pitch, SIZE_X, SIZE_Y, SIZE_Z);
    gpuErrchk(cudaMemcpy(d_volumeMem.ptr, h_volumeMem, size, cudaMemcpyHostToDevice));

    cudaArray * d_volumeArray;
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<VolumeType>();
    cudaExtent volumeSize = make_cudaExtent(SIZE_X, SIZE_Y, SIZE_Z);
    gpuErrchk( cudaMalloc3DArray(&d_volumeArray, &channelDesc, volumeSize) ); 

    cudaMemcpy3DParms copyParams = {0};
    copyParams.srcPtr = d_volumeMem;
    copyParams.dstArray = d_volumeArray;
    copyParams.extent = volumeSize;
    copyParams.kind = cudaMemcpyDeviceToDevice;
    gpuErrchk( cudaMemcpy3D(&copyParams) ); 

    tex.normalized = false;                      
    tex.filterMode = cudaFilterModeLinear;      
    tex.addressMode[0] = cudaAddressModeWrap;   
    tex.addressMode[1] = cudaAddressModeWrap;
    tex.addressMode[2] = cudaAddressModeWrap;
    gpuErrchk(cudaBindTextureToArray(tex, d_volumeArray, channelDesc));

    size_t osize = 64 * sizeof(VolumeType);
    gpuErrchk(cudaMalloc((void**)&d_output, osize));

    testKernel<<<1,dim3(4,4,4)>>>(d_output,4,4,4);
    gpuErrchk(cudaPeekAtLastError());

    h_output = (VolumeType *)malloc(osize);
    gpuErrchk(cudaMemcpy(h_output, d_output, osize, cudaMemcpyDeviceToHost));

    for(int i=0; i<64; i++)
        fprintf(stdout, "%d %f\n", i, h_output[i]);

    return 0;
}

テクスチャ読み取りの出力がホスト上の元のソース メモリと一致することを自分で確認できます。

于 2012-05-19T08:30:12.763 に答える