Массив cudaArray для текстурного кода с несколькими графическими процессорами - PullRequest
Новая
       5

Массив cudaArray для текстурного кода с несколькими графическими процессорами

0 голосов
/ 28 ноября 2018

У меня есть код, который я пытаюсь заставить работать для общего случая с несколькими графическими процессорами, для n количества равных устройств, где n неизвестно во время компиляции.

Для этого кода мне нужно привязать к текстурной памяти некоторый массив, и мне нужны точно такие же данные для привязки к разным графическим процессорам.

Мой единственный код памяти графического процессора для привязки 3D-текстурвыглядит так: 

cudaArray *d_imagedata = 0;
const cudaExtent extent = make_cudaExtent(geo.nVoxelX, geo.nVoxelY, geo.nVoxelZ);
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
cudaMalloc3DArray(&d_imagedata, &channelDesc, extent);
cudaCheckErrors("cudaMalloc3D error 3D tex");

cudaMemcpy3DParms copyParams = { 0 };
copyParams.srcPtr = make_cudaPitchedPtr((void*)img, extent.width*sizeof(float), extent.width, extent.height);
copyParams.dstArray = d_imagedata;
copyParams.extent = extent;
copyParams.kind = cudaMemcpyHostToDevice;
cudaMemcpy3D(&copyParams);

cudaCheckErrors("cudaMemcpy3D fail");

// Configure texture options
tex.normalized = false;
tex.filterMode = cudaFilterModePoint; 
tex.addressMode[0] = cudaAddressModeBorder;
tex.addressMode[1] = cudaAddressModeBorder;
tex.addressMode[2] = cudaAddressModeBorder;

cudaBindTextureToArray(tex, d_imagedata, channelDesc);

Это стандартная копия cudaArray, а затем процесс привязки и настройки, ничего нового здесь.

Чтобы преобразовать этот код в мульти GPU, я знаю, что мне не нужно изменять глобальную ссылку на текстуру tex, поскольку CUDA будет знать, что разные GPU имеют разные tex, однако мне нужно n cudaArray *d_imagedata экземпляров, по одному на каждый графический процессор.

Как создать (и выделить) массив из cudaArray с?

Если бы это были глобальные указатели памяти, это было бы проще, просто процессор malloc с двойным указателем и затем cudaMalloc на каждом из них работал бы, но cudaArray не является стандартным типомЯ не понял, как создать гибкий массив из него.

1 Ответ

0 голосов
/ 28 ноября 2018

Я бы порекомендовал использовать объекты текстуры, а не ссылки на текстуры.

Используя объекты текстуры , тривиальная модификация представленного кода здесь , кажется, работает правильно дляme: 

$ cat t341.cu
#include <helper_cuda.h>
#include <curand.h>
#define NUM_TEX 4

const int SizeNoiseTest = 32;
const int cubeSizeNoiseTest = SizeNoiseTest*SizeNoiseTest*SizeNoiseTest;
static cudaTextureObject_t texNoise[NUM_TEX];

__global__ void AccesTexture(cudaTextureObject_t my_tex)
{
        float test = tex3D<float>(my_tex,(float)threadIdx.x,(float)threadIdx.y,(float)threadIdx.z);//by using this the error occurs
        printf("thread: %d,%d,%d, value: %f\n", threadIdx.x, threadIdx.y, threadIdx.z, test);
}

void CreateTexture()
{

    for (int i = 0; i < NUM_TEX; i++){
        cudaSetDevice(i);
        float *d_NoiseTest;//Device Array with random floats
        cudaMalloc((void **)&d_NoiseTest, cubeSizeNoiseTest*sizeof(float));//Allocation of device Array
        //curand Random Generator (needs compiler link -lcurand)
        curandGenerator_t gen;
        curandCreateGenerator(&gen,CURAND_RNG_PSEUDO_DEFAULT);
        curandSetPseudoRandomGeneratorSeed(gen,1235ULL+i);
        curandGenerateUniform(gen, d_NoiseTest, cubeSizeNoiseTest);//writing data to d_NoiseTest
        curandDestroyGenerator(gen);

        //cudaArray Descriptor
        cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
        //cuda Array
        cudaArray *d_cuArr;
        checkCudaErrors(cudaMalloc3DArray(&d_cuArr, &channelDesc, make_cudaExtent(SizeNoiseTest*sizeof(float),SizeNoiseTest,SizeNoiseTest), 0));
        cudaMemcpy3DParms copyParams = {0};


        //Array creation
        copyParams.srcPtr   = make_cudaPitchedPtr(d_NoiseTest, SizeNoiseTest*sizeof(float), SizeNoiseTest, SizeNoiseTest);
        copyParams.dstArray = d_cuArr;
        copyParams.extent   = make_cudaExtent(SizeNoiseTest,SizeNoiseTest,SizeNoiseTest);
        copyParams.kind     = cudaMemcpyDeviceToDevice;
        checkCudaErrors(cudaMemcpy3D(&copyParams));
        //Array creation End

        cudaResourceDesc    texRes;
        memset(&texRes, 0, sizeof(cudaResourceDesc));
        texRes.resType = cudaResourceTypeArray;
        texRes.res.array.array  = d_cuArr;
        cudaTextureDesc     texDescr;
        memset(&texDescr, 0, sizeof(cudaTextureDesc));
        texDescr.normalizedCoords = false;
        texDescr.filterMode = cudaFilterModeLinear;
        texDescr.addressMode[0] = cudaAddressModeClamp;   // clamp
        texDescr.addressMode[1] = cudaAddressModeClamp;
        texDescr.addressMode[2] = cudaAddressModeClamp;
        texDescr.readMode = cudaReadModeElementType;
        checkCudaErrors(cudaCreateTextureObject(&texNoise[i], &texRes, &texDescr, NULL));}
}

int main(int argc, char **argv)
{
        CreateTexture();
        cudaSetDevice(0);
        AccesTexture<<<1,dim3(2,2,2)>>>(texNoise[0]);
        cudaSetDevice(1);
        AccesTexture<<<1,dim3(2,2,2)>>>(texNoise[1]);
        cudaSetDevice(2);
        AccesTexture<<<1,dim3(2,2,2)>>>(texNoise[2]);
        checkCudaErrors(cudaPeekAtLastError());
        cudaSetDevice(0);
        checkCudaErrors(cudaDeviceSynchronize());
        cudaSetDevice(1);
        checkCudaErrors(cudaDeviceSynchronize());
        cudaSetDevice(2);
        checkCudaErrors(cudaDeviceSynchronize());
        return 0;
}
$ nvcc -arch=sm_30 -I/usr/local/cuda/samples/common/inc -lcurand -o t341 t341.cu
$ cuda-memcheck ./t341
========= CUDA-MEMCHECK
thread: 0,0,0, value: 0.310691
thread: 1,0,0, value: 0.627906
thread: 0,1,0, value: 0.638900
thread: 1,1,0, value: 0.665186
thread: 0,0,1, value: 0.167465
thread: 1,0,1, value: 0.565227
thread: 0,1,1, value: 0.397606
thread: 1,1,1, value: 0.503013
thread: 0,0,0, value: 0.809163
thread: 1,0,0, value: 0.795669
thread: 0,1,0, value: 0.808565
thread: 1,1,0, value: 0.847564
thread: 0,0,1, value: 0.853998
thread: 1,0,1, value: 0.688446
thread: 0,1,1, value: 0.733255
thread: 1,1,1, value: 0.649379
thread: 0,0,0, value: 0.040824
thread: 1,0,0, value: 0.087417
thread: 0,1,0, value: 0.301392
thread: 1,1,0, value: 0.298669
thread: 0,0,1, value: 0.161962
thread: 1,0,1, value: 0.316443
thread: 0,1,1, value: 0.452077
thread: 1,1,1, value: 0.477722
========= ERROR SUMMARY: 0 errors
$

Обратите внимание, что для простоты представления эта функция CreateTexture() перезаписывает ранее назначенные указатели устройства, такие как d_NoiseTest и d_cuArr, во время обработки цикла.Это не является незаконной или функциональной проблемой, но повышает вероятность утечек памяти.(Но см. Ниже пример того, как этого избежать.)

РЕДАКТИРОВАТЬ: Исходя из вопроса в комментариях, ни один из них не должен зависеть от времени компиляции.Вот модификация приведенного выше кода, демонстрирующая это: 

$ cat t342.cu
#include <helper_cuda.h>
#include <curand.h>

const int SizeNoiseTest = 32;
const int cubeSizeNoiseTest = SizeNoiseTest*SizeNoiseTest*SizeNoiseTest;

__global__ void AccesTexture(cudaTextureObject_t my_tex)
{
        float test = tex3D<float>(my_tex,(float)threadIdx.x,(float)threadIdx.y,(float)threadIdx.z);//by using this the error occurs
        printf("thread: %d,%d,%d, value: %f\n", threadIdx.x, threadIdx.y, threadIdx.z, test);
}

void CreateTexture(int num, cudaTextureObject_t *texNoise, cudaArray **d_cuArr, float **d_NoiseTest)
{

    for (int i = 0; i < num; i++){
        cudaSetDevice(i);
        cudaMalloc((void **)&d_NoiseTest[i], cubeSizeNoiseTest*sizeof(float));//Allocation of device Array
        //curand Random Generator (needs compiler link -lcurand)
        curandGenerator_t gen;
        curandCreateGenerator(&gen,CURAND_RNG_PSEUDO_DEFAULT);
        curandSetPseudoRandomGeneratorSeed(gen,1235ULL+i);
        curandGenerateUniform(gen, d_NoiseTest[i], cubeSizeNoiseTest);//writing data to d_NoiseTest
        curandDestroyGenerator(gen);

        //cudaArray Descriptor
        cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
        //cuda Array
        checkCudaErrors(cudaMalloc3DArray(&d_cuArr[i], &channelDesc, make_cudaExtent(SizeNoiseTest*sizeof(float),SizeNoiseTest,SizeNoiseTest), 0));
        cudaMemcpy3DParms copyParams = {0};


        //Array creation
        copyParams.srcPtr   = make_cudaPitchedPtr(d_NoiseTest[i], SizeNoiseTest*sizeof(float), SizeNoiseTest, SizeNoiseTest);
        copyParams.dstArray = d_cuArr[i];
        copyParams.extent   = make_cudaExtent(SizeNoiseTest,SizeNoiseTest,SizeNoiseTest);
        copyParams.kind     = cudaMemcpyDeviceToDevice;
        checkCudaErrors(cudaMemcpy3D(&copyParams));
        //Array creation End

        cudaResourceDesc    texRes;
        memset(&texRes, 0, sizeof(cudaResourceDesc));
        texRes.resType = cudaResourceTypeArray;
        texRes.res.array.array  = d_cuArr[i];
        cudaTextureDesc     texDescr;
        memset(&texDescr, 0, sizeof(cudaTextureDesc));
        texDescr.normalizedCoords = false;
        texDescr.filterMode = cudaFilterModeLinear;
        texDescr.addressMode[0] = cudaAddressModeClamp;   // clamp
        texDescr.addressMode[1] = cudaAddressModeClamp;
        texDescr.addressMode[2] = cudaAddressModeClamp;
        texDescr.readMode = cudaReadModeElementType;
        checkCudaErrors(cudaCreateTextureObject(&texNoise[i], &texRes, &texDescr, NULL));}
}
void FreeTexture(int num, cudaTextureObject_t *texNoise, cudaArray **d_cuArr, float **d_NoiseTest)
{
   for (int i = 0; i < num; i++){
     cudaFree(d_NoiseTest[i]);
     cudaDestroyTextureObject(texNoise[i]);
     cudaFreeArray(d_cuArr[i]);}
}

int main(int argc, char **argv)
{
        int num_dev = 1;
        if (argc > 1) num_dev = atoi(argv[1]);
        cudaTextureObject_t *texNoise = new cudaTextureObject_t[num_dev];
        cudaArray **d_cuArr = new cudaArray*[num_dev];
        float **d_NoiseTest = new float*[num_dev];
        CreateTexture(num_dev, texNoise, d_cuArr, d_NoiseTest);
        for (int i = 0; i < num_dev; i++){
          cudaSetDevice(i);
          AccesTexture<<<1,dim3(2,2,2)>>>(texNoise[i]);}
        checkCudaErrors(cudaPeekAtLastError());
        for (int i = 0; i < num_dev; i++){
          cudaSetDevice(i);
          checkCudaErrors(cudaDeviceSynchronize());}
        FreeTexture(num_dev, texNoise, d_cuArr, d_NoiseTest);
        delete[] d_cuArr;
        delete[] d_NoiseTest;
        delete[] texNoise;
        return 0;
}
$ nvcc -I/usr/local/cuda/samples/common/inc -lcurand -o t342 t342.cu
$ cuda-memcheck ./t342
========= CUDA-MEMCHECK
thread: 0,0,0, value: 0.310691
thread: 1,0,0, value: 0.627906
thread: 0,1,0, value: 0.638900
thread: 1,1,0, value: 0.665186
thread: 0,0,1, value: 0.167465
thread: 1,0,1, value: 0.565227
thread: 0,1,1, value: 0.397606
thread: 1,1,1, value: 0.503013
========= ERROR SUMMARY: 0 errors
$ cuda-memcheck ./t342 2
========= CUDA-MEMCHECK
thread: 0,0,0, value: 0.310691
thread: 1,0,0, value: 0.627906
thread: 0,1,0, value: 0.638900
thread: 1,1,0, value: 0.665186
thread: 0,0,1, value: 0.167465
thread: 1,0,1, value: 0.565227
thread: 0,1,1, value: 0.397606
thread: 1,1,1, value: 0.503013
thread: 0,0,0, value: 0.809163
thread: 1,0,0, value: 0.795669
thread: 0,1,0, value: 0.808565
thread: 1,1,0, value: 0.847564
thread: 0,0,1, value: 0.853998
thread: 1,0,1, value: 0.688446
thread: 0,1,1, value: 0.733255
thread: 1,1,1, value: 0.649379
========= ERROR SUMMARY: 0 errors
$ cuda-memcheck ./t342 3
========= CUDA-MEMCHECK
thread: 0,0,0, value: 0.310691
thread: 1,0,0, value: 0.627906
thread: 0,1,0, value: 0.638900
thread: 1,1,0, value: 0.665186
thread: 0,0,1, value: 0.167465
thread: 1,0,1, value: 0.565227
thread: 0,1,1, value: 0.397606
thread: 1,1,1, value: 0.503013
thread: 0,0,0, value: 0.809163
thread: 1,0,0, value: 0.795669
thread: 0,1,0, value: 0.808565
thread: 1,1,0, value: 0.847564
thread: 0,0,1, value: 0.853998
thread: 1,0,1, value: 0.688446
thread: 0,1,1, value: 0.733255
thread: 1,1,1, value: 0.649379
thread: 0,0,0, value: 0.040824
thread: 1,0,0, value: 0.087417
thread: 0,1,0, value: 0.301392
thread: 1,1,0, value: 0.298669
thread: 0,0,1, value: 0.161962
thread: 1,0,1, value: 0.316443
thread: 0,1,1, value: 0.452077
thread: 1,1,1, value: 0.477722
========= ERROR SUMMARY: 0 errors
$

Этот код был запущен в системе, которая имеет (как минимум) 3 графических процессора.Я также обновил приведенный выше пример, чтобы он продемонстрировал, как создать массив указателей типа cudaArray, а также продемонстрировал, как можно избежать утечек памяти.

Добро пожаловать на сайт PullRequest, где вы можете задавать вопросы и получать ответы от других членов сообщества.

Нет похожих вопросов

...