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

2

Я хочу иметь массив с плавающей точкой в 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) ); 

На самом деле, моя программа работает хорошо. Но я не уверен, что результат правильный. Вот моя проблема, в Libra CUDA, он сказал, что первый параметр make_cudaExtent это "Ширина в байтах" и два других - высота и глубина в элементах. Так что я думаю, что в моем коде выше, пятая строка должна быть

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

Но в этом случае будет ошибка "неверный аргумент" в cutilSafeCall (cudaMemcpy3D (& amp; 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), то есть число в байтах вместо числа в элементах.

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

Мне интересно, как вы можете написать "На самом деле, моя программа работает хорошо. Но я не уверен, что результат правильный. Разве это не противоречие? И как можно не быть уверенным, что результат «правильный»? Конечно, вы знаете, какой должен быть правильный результат? talonmies
Я имею в виду, что программа может работать и выводить результат, но я не уверен, что результат правильный. На самом деле мне нужна 3d текстура для рисования картинки, но я не знаю, какой должна быть картинка. И дело в том, как я копирую данные в трехмерный массив, который я описал выше. Если это не правильно, я думаю, что результат может иметь некоторые ошибки. TonyLic

Ваш Ответ

2   ответа
2

cudaMemcpy3D говорит:

The extent field defines the dimensions of the transferred area in elements. If a CUDA array is participating in the copy, the extent is defined in terms of that array's elements. If no CUDA array is participating in the copy then the extents are defined in elements of unsigned char.

и аналогично документация дляcudaMalloc3DArray заметки:

All values are specified in elements

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

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

-1

так как задействован cudaArray. ChannelDesc, предоставленный массиву, содержит информацию о размере с плавающей точкой (4 байта). Ваша спецификация степени с & quot; * sizeof (VolumeType) & quot; было бы правильно копировать между двумя указателями памяти (с srcPtr, используется dstPtr). Кроме того, srcPos и dstPos затем должны быть заданы в байтах, то есть первый параметр «* sizeof (VolumeType)».

Проблемы с высотой звука по-прежнему могут возникать при 3d операциях в зависимости от графического процессора / драйвера Я видел это, но редко (2 ^ n размеры должны быть в порядке). Вы также можете разбить его, используя cudaMemCpy2DToArray в одном цикле for, так как он должен быть более терпимым. CudaMalloc2D не существует, поэтому SDK выдает любые корректные шаги для операций 2d.

Похожие вопросы