Как использовать make_cudaExtent для правильного определения cudaExtent?

Question

Я хочу иметь 3d буев в CUDA, вот мой код:

#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 является «Ширина в байтах», а две другие - высота и глубина элементов. Так что я думаю, что в моем коде выше, пятая строка должна быть

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

Но таким образом, было бы ошибкой «неверный аргумент» в cutilSafeCall (cudaMemcpy3D (& copyParams)); Зачем?

И еще одна головоломка - это strcut cudaExtent, как указано в CUDA-библиотеке, ее ширина компонента означает «Ширина элементов при обращении к памяти массива, в байтах при обращении к линейной памяти». Поэтому я думаю, что в моем коде, когда я ссылаюсь на volumeSize.width, он должен быть числом в элементах. Однако, если я использую

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

volumeSize.width будет SIZE_X * SizeOf (VolumeType) (128 * 4), то есть число в байтах, а не число в элементах.

Во многих SDK CUDA они используют char как VolumeType, поэтому они просто используют SIZE_X в качестве первого аргумента в make_cudaExtent. Но мой плавает, так что любой может сказать мне, что это правильный способ создать cudaExtent, если мне нужно использовать его для создания 3D-массива? Большое спасибо!

Answer 1

Давайте рассмотрим, что документация для cudaMemcpy3D говорит:

Поле степени определяет размеры передаваемой площади в элементов. Если в копии участвует массив CUDA, то размер равен , определенному в терминах элементов этого массива. Если ни один массив CUDA не является , участвующим в копии, то экстенты определяются в элементах unsigned char.

и точно так же документация по cudaMalloc3DArray примечаниями:

Все значения указаны в элементах

Так насколько вам необходимо сформировать для обоих вызовов необходимо иметь первое измерение в элементах (поскольку одно из распределений в 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; 
} 

Вы можете подтвердить для себя, что выход текстур читает соответствует первоисточнику память на хосте.