Cuda — копирование из глобальной памяти устройства в текстурную память

Я пытаюсь выполнить две задачи (разделенные на 2 ядра) на GPU, используя Cuda и C ++. В качестве входных данных я беру матрицу NxM (хранится в памяти на хосте как массив с плавающей запятой). Затем я буду использовать ядро, которое выполняет некоторые операции с этой матрицей, чтобы сделать ее матрицей NxMxD. Затем у меня есть второе ядро, которое выполняет некоторые операции с этой трехмерной матрицей (я просто читаю значения, мне не нужно записывать значения в нее).

Работа с текстурной памятью, кажется, намного быстрее для моей задачи, поэтому мой вопрос, возможно ли скопировать мои данные из глобальной памяти на устройстве после ядра 1 и перенести их непосредственно в текстурную память для ядра 2, не возвращая их на хост? ?

ОБНОВИТЬ

Я добавил код, чтобы лучше проиллюстрировать мою проблему.

Вот два ядра. Первый пока просто заполнитель и копирует 2D матрицу в 3D.

__global__ void computeFeatureVector(float* imData3D_dev, int imX, int imY, int imZ) {

//calculate each thread global index
int xindex=blockIdx.x*blockDim.x+threadIdx.x;
int yindex=blockIdx.y*blockDim.y+threadIdx.y;

#pragma unroll
for (int z=0; z<imZ; z++) {
imData3D_dev[xindex+yindex*imX + z*imX*imY] = tex2D(texImIp,xindex,yindex);
}
}

Второй возьмет эту трехмерную матрицу, теперь представленную как текстуру, и выполнит над ней некоторые операции. Бланк пока.

__global__ void kernel2(float* resData_dev, int imX) {
//calculate each thread global index
int xindex=blockIdx.x*blockDim.x+threadIdx.x;
int yindex=blockIdx.y*blockDim.y+threadIdx.y;

resData_dev[xindex+yindex*imX] = tex3D(texImIp3D,xindex,yindex, 0);

return;
}

Тогда основная часть кода выглядит следующим образом:

// declare textures
texture<float,2,cudaReadModeElementType> texImIp;
texture<float,3,cudaReadModeElementType> texImIp3D;

void main_fun() {

// constants
int imX = 1024;
int imY = 768;
int imZ = 16;

// input data
float* imData2D  = new float[sizeof(float)*imX*imY];
for(int x=0; x<imX*imY; x++)
imData2D[x] = (float) rand()/RAND_MAX;

//create channel to describe data type
cudaArray* carrayImIp;
cudaChannelFormatDesc channel;
channel=cudaCreateChannelDesc<float>();

//allocate device memory for cuda array
cudaMallocArray(&carrayImIp,&channel,imX,imY);

//copy matrix from host to device memory
cudaMemcpyToArray(carrayImIp,0,0,imData2D,sizeof(float)*imX*imY,cudaMemcpyHostToDevice);

// Set texture properties
texImIp.filterMode=cudaFilterModePoint;
texImIp.addressMode[0]=cudaAddressModeClamp;
texImIp.addressMode[1]=cudaAddressModeClamp;

// bind texture reference with cuda array
cudaBindTextureToArray(texImIp,carrayImIp);

// kernel params
dim3 blocknum;
dim3 blocksize;
blocksize.x=16; blocksize.y=16; blocksize.z=1;
blocknum.x=(int)ceil((float)imX/16);
blocknum.y=(int)ceil((float)imY/16);

// store output here
float* imData3D_dev;
cudaMalloc((void**)&imData3D_dev,sizeof(float)*imX*imY*imZ);

// execute kernel
computeFeatureVector<<<blocknum,blocksize>>>(imData3D_dev, imX, imY, imZ);

//unbind texture reference to free resource
cudaUnbindTexture(texImIp);

// check copied ok
float* imData3D  = new float[sizeof(float)*imX*imY*imZ];
cudaMemcpy(imData3D,imData3D_dev,sizeof(float)*imX*imY*imZ,cudaMemcpyDeviceToHost);
cout << " kernel 1" << endl;
for (int x=0; x<10;x++)
cout << imData3D[x] << " ";
cout << endl;
delete [] imData3D;//
// kernel 2
//// copy data on device to 3d array
cudaArray* carrayImIp3D;
cudaExtent volumesize;
volumesize = make_cudaExtent(imX, imY, imZ);
cudaMalloc3DArray(&carrayImIp3D,&channel,volumesize);
cudaMemcpyToArray(carrayImIp3D,0,0,imData3D_dev,sizeof(float)*imX*imY*imZ,cudaMemcpyDeviceToDevice);

// texture params and bind
texImIp3D.filterMode=cudaFilterModePoint;
texImIp3D.addressMode[0]=cudaAddressModeClamp;
texImIp3D.addressMode[1]=cudaAddressModeClamp;
texImIp3D.addressMode[2]=cudaAddressModeClamp;
cudaBindTextureToArray(texImIp3D,carrayImIp3D,channel);

// store output here
float* resData_dev;
cudaMalloc((void**)&resData_dev,sizeof(float)*imX*imY);

// kernel 2
kernel2<<<blocknum,blocksize>>>(resData_dev, imX);
cudaUnbindTexture(texImIp3D);

//copy result matrix from device to host memory
float* resData  = new float[sizeof(float)*imX*imY];
cudaMemcpy(resData,resData_dev,sizeof(float)*imX*imY,cudaMemcpyDeviceToHost);

// check copied ok
cout << " kernel 2" << endl;
for (int x=0; x<10;x++)
cout << resData[x] << " ";
cout << endl;delete [] imData2D;
delete [] resData;
cudaFree(imData3D_dev);
cudaFree(resData_dev);
cudaFreeArray(carrayImIp);
cudaFreeArray(carrayImIp3D);

}

Я рад, что первое ядро ​​работает правильно, но 3D-матрица imData3D_dev, похоже, неправильно привязана к текстуре texImIp3D.

ОТВЕТ

Я решил свою проблему с помощью cudaMemcpy3D. Вот пересмотренный код для второй части основной функции. imData3D_dev содержит трехмерную матрицу в глобальной памяти из первого ядра.

    cudaArray* carrayImIp3D;
cudaExtent volumesize;
volumesize = make_cudaExtent(imX, imY, imZ);
cudaMalloc3DArray(&carrayImIp3D,&channel,volumesize);
cudaMemcpy3DParms copyparms={0};

copyparms.extent = volumesize;
copyparms.dstArray = carrayImIp3D;
copyparms.kind = cudaMemcpyDeviceToDevice;
copyparms.srcPtr = make_cudaPitchedPtr((void*)imData3D_dev, sizeof(float)*imX,imX,imY);
cudaMemcpy3D(&copyparms);

// texture params and bind
texImIp3D.filterMode=cudaFilterModePoint;
texImIp3D.addressMode[0]=cudaAddressModeClamp;
texImIp3D.addressMode[1]=cudaAddressModeClamp;
texImIp3D.addressMode[2]=cudaAddressModeClamp;

cudaBindTextureToArray(texImIp3D,carrayImIp3D,channel);

// store output here
float* resData_dev;
cudaMalloc((void**)&resData_dev,sizeof(float)*imX*imY);

kernel2<<<blocknum,blocksize>>>(resData_dev, imX);

// ... clean up