Как успешно читать с 2D текстуры - PullRequest
6 голосов
/ 01 октября 2010

Как я могу:

  1. Привязка плавающей памяти cudaMallocPitch к эталону 2D текстуры
  2. Скопируйте некоторые данные хоста в 2D-массив на устройстве
  3. Добавьте единицу к эталону текстуры и запишите в a.) 2D-массив Pitch ИЛИ b.) Запишите в линейный массив памяти
  4. Прочитайте ответ и отобразите его.

Ниже приведен код, который должен выполнить это. Обратите внимание, что для размеров массива NxN мой код работает. Для NxM, где N! = M, мой код кусает пыль (не правильный результат). Если вы сможете решить эту проблему, я награду вас 1 интернетом (поставка ограничена). Может быть, я сумасшедший, но согласно документации это должно работать (и это работает для квадратных массивов!). Прикрепленный код должен работать с 'nvcc whateveryoucallit.cu -o runit'.

Помощь приветствуется!

#include<stdio.h>
#include<cuda.h>
#include<iostream>
#define height 16
#define width 11
#define BLOCKSIZE 16

using namespace std;

// Device Kernels

//Texture reference Declaration
texture<float,2> texRefEx;


__global__ void kernel_w_textures(float* devMPPtr, float * devMPtr, int pitch)
{
 // Thread indexes
        unsigned int idx = blockIdx.x*blockDim.x + threadIdx.x;
        unsigned int idy = blockIdx.y*blockDim.y + threadIdx.y;

 // Texutre Coordinates
 float u=(idx)/float(width);
 float v=(idy)/float(height);
 devMPtr[idy*width+idx]=devMPPtr[idy*pitch/sizeof(float)+idx];
 // Write Texture Contents to malloc array +1
 devMPtr[idy*width+idx]= tex2D(texRefEx,u,v);//+1.0f;
}
int main()
{
 // memory size
 size_t memsize=height*width;
 size_t offset;
 float * data,  // input from host
  *h_out,  // host space for output
  *devMPPtr, // malloc Pitch ptr
  *devMPtr; // malloc ptr

 size_t pitch;

 // Allocate space on the host
 data=(float *)malloc(sizeof(float)*memsize);
 h_out=(float *)malloc(sizeof(float)*memsize);


// Define data
for (int i = 0; i <  height; i++)
 for (int j=0; j < width; j++)
  data[i*width+j]=float(j);

// Define the grid
dim3 grid((int)(width/BLOCKSIZE)+1,(int)(height/BLOCKSIZE)+1), threads(BLOCKSIZE,BLOCKSIZE);

// allocate Malloc Pitch
cudaMallocPitch((void**)&devMPPtr,&pitch, width * sizeof(float), height);

// Print the pitch
printf("The pitch is %d \n",pitch/sizeof(float));

// Texture Channel Description
//cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32,0,0,0,cudaChannelFormatKindFloat);

// Bind texture to pitch mem:
cudaBindTexture2D(&offset,&texRefEx,devMPPtr,&channelDesc,width,height,pitch);
cout << "My Description x is " << channelDesc.x << endl;
cout << "My Description y is " << channelDesc.y << endl;
cout << "My Description z is " << channelDesc.z << endl;
cout << "My Description w is " << channelDesc.w << endl;
cout << "My Description kind is " << channelDesc.f << endl;
cout << "Offset is " << offset << endl;

// Set mutable properties:
texRefEx.normalized=true;
texRefEx.addressMode[0]=cudaAddressModeWrap;
texRefEx.addressMode[1]=cudaAddressModeWrap;
texRefEx.filterMode= cudaFilterModePoint;

// Allocate cudaMalloc memory
cudaMalloc((void**)&devMPtr,memsize*sizeof(float));

// Read data from host to device
cudaMemcpy2D((void*)devMPPtr,pitch,(void*)data,sizeof(float)*width,
  sizeof(float)*width,height,cudaMemcpyHostToDevice);

//Read back and check this memory
cudaMemcpy2D((void*)h_out,width*sizeof(float),(void*)devMPPtr,pitch,
  sizeof(float)*width,height,cudaMemcpyDeviceToHost);

// Print the memory
 for (int i=0; i<height; i++){
  for (int j=0; j<width; j++){
   printf("%2.2f ",h_out[i*width+j]);
  }
 cout << endl;
 }

 cout << "Done" << endl;
// Memory is fine... 

kernel_w_textures<<<grid,threads>>>(devMPPtr, devMPtr, pitch);

// Copy back data to host
cudaMemcpy((void*)h_out,(void*)devMPtr,width*height*sizeof(float),cudaMemcpyDeviceToHost);


// Print the Result
 cout << endl;
 for (int i=0; i<height; i++){
  for (int j=0; j<width; j++){
   printf("%2.2f ",h_out[i*width+j]);
  }
 cout << endl;
 }
 cout << "Done" << endl;

return(0);
}

Edit 17 октября: так что я до сих пор не нашел решение этой проблемы. Nvidia довольно молчит об этом, кажется, что мир тоже. Я нашел обходной путь, используя общий мем, но если у кого-нибудь есть текстурное решение, я буду очень рад.

Редактировать Октобур 26: Все еще нет решения, но все еще заинтересован в одном, если кто-нибудь знает.

Редактировать 26 июля: Ух, прошло уже 9 месяцев - и я все время игнорировал правильный ответ. Трюк был:

if ( idx < width  && idy < height){//.... code }

Как было указано ранее. Спасибо всем, кто внес свой вклад!

Ответы [ 5 ]

3 голосов
/ 27 октября 2010

Это может иметь отношение к размеру вашего блока.В этом коде вы пытаетесь записать блок из потоков 16x16 в блок памяти 11x16.Это означает, что некоторые из ваших потоков пишут в нераспределенную память.Это также объясняет, почему ваши тесты (16 * M на 32 * N) работали: не было потоков, записывающих в нераспределенную память, поскольку ваши измерения были кратны 16.

Простой способ решить эту проблему - эточто-то вроде этого:

if ((x < width) && (y < height)) {
   // write output 
  devMPtr[idy*width+idx]= tex2D(texRefEx,u,v); 
}

Вам нужно будет либо передать высоту и ширину функции ядра, либо скопировать константу на карту перед вызовом ядра.

2 голосов
/ 13 сентября 2012

Я думаю:

 float u=(idx)/float(width);
 float v=(idy)/float(height);

должно быть

 float u=(idx+0.5f)/float(width);
 float v=(idy+0.5f)/float(height);

Чтобы получить идентичный ввод / вывод, в противном случае второй столбец вывода равен первому столбцу ввода, а не второму, а второй последний столбец вывода также неверен.

Пожалуйста, поправьте меня, если у вас другое наблюдение.

1 голос
/ 28 октября 2010
 // Texutre Coordinates
 float u=(idx + 0.5)/float(width);
 float v=(idy + 0.5)/float(height);

Вам нужно смещение, чтобы добраться до центра текселя.Я думаю, что может быть некоторая ошибка округления для вашей не кратной 16 текстур.Я попробовал это, и это сработало для меня (оба выхода были идентичны).

0 голосов
/ 20 декабря 2010

Возможно, взгляните на эту ветку: http://forums.nvidia.com/index.php?showtopic=186585

Еще один очень полезный пример кода в настоящее время находится в NVIDIA SDK;как упоминалось в вышеупомянутой теме на форумах NVIDIA, пример simplePitchLinearTexture работает хорошо.

Поскольку мы используем текстурную память, я считаю, что размеры 2D-сетки должны быть степенями 2 на некоторых аппаратных средствах, такжепредлагается в одном из ответов выше.

0 голосов
/ 01 октября 2010

Графические карты обычно ожидают, что текстуры будут иметь размеры, равные степени 2, это особенно верно для карт nVidia. Cuda cudaMallocPitch и cudaMemcpy2D работают с этими высотами и, глядя на ваш код, самое безопасное решение - настроить ширину и высоту самостоятельно, чтобы быть на безопасной стороне. В противном случае Cuda может записать в недопустимую память, потому что он будет ожидать неправильных смещений:

#define height 16
#define width 11

...

size_t roundUpToPowerOf2(size_t v)
{
  // See http://graphics.stanford.edu/~seander/bithacks.html#RoundUpPowerOf2
  --v;
  v |= v >> 1;
  v |= v >> 2;
  v |= v >> 4;
  v |= v >> 8;
  v |= v >> 16;
  ++v;
  return v;
}
...

size_t horizontal_pitch = roundUpToPowerOf2(width);
size_t vertical_pitch = roundUpToPowerOf2(height);
size_t memsize = horizontal_pitch * vertical_pitch;

...

// Read data from host to device
cudaMemcpy2D((void*)devMPPtr,pitch,(void*)data,sizeof(float)*horizontal_pitch,
  sizeof(float)*width,height,cudaMemcpyHostToDevice);

//Read back and check this memory
cudaMemcpy2D((void*)h_out,horizontal_pitch*sizeof(float),(void*)devMPPtr,pitch,
  sizeof(float)*width,height,cudaMemcpyDeviceToHost);

// Print the memory
 for (int i=0; i<height; i++){
  for (int j=0; j<width; j++){
   printf("%2.2f ",h_out[i*horizontal_pitch+j]);
  }
 cout << endl;
 }

...

// Copy back data to host
cudaMemcpy((void*)h_out,(void*)devMPtr,horizontal_pitch*vertical_pitch*sizeof(float),cudaMemcpyDeviceToHost);

// Print the Result
 cout << endl;
 for (int i=0; i<height; i++){
  for (int j=0; j<width; j++){
   printf("%2.2f ",h_out[i*horizontal_pitch+j]);
  }
 cout << endl;
 }
 cout << "Done" << endl;

Надеюсь, я не упустил из виду ни одного места, где вместо горизонтальной ширины / высоты следует использовать Horizontal_pitch / vertical_pitch.

...