¿Cómo puede I:Cómo leer con éxito de una textura 2D
- Enlazar cudaMallocPitch memoria flotador a una referencia de textura 2D
- copiar algunos datos de host a la matriz 2D en el dispositivo
- Añadir uno a la referencia de textura y escribir en a.) la matriz de tono 2D O b.) escribir en una matriz de memoria lineal
- Lea la respuesta y visualícela.
A continuación se muestra un código que debe lograr esto. Tenga en cuenta que para los tamaños de matriz NxN, mi código funciona. Para NxM donde N! = M, mi código muerde el polvo (no el resultado correcto). Si puede resolver este problema, le otorgaré 1 internets (suministro limitado). Tal vez estoy loco, pero de acuerdo con la documentación esto debería funcionar (¡y funciona para matrices cuadradas!). El código adjunto debería ejecutarse con 'nvcc whateveryoucallit.cu -o runit'.
¡Ayuda se aprecia!
#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 October 17: Así que todavía no he encontrado una solución a este problema. Nvidia guarda silencio sobre esto, parece que el mundo también lo es. Encontré una solución utilizando memorios compartidos, pero si alguien tiene una solución de textura, sería muy, por favor.
Editar Octoboer 26: Todavía no hay solución, pero sigue interesado en uno si alguien sabe.
Editar 26 de julio: Wow, han pasado 9 meses, y había pasado por alto la respuesta correcta todo el tiempo. El truco fue:
if (idx < width && idy < height){//.... code }
Como se ha señalado anteriormente. ¡Gracias a todos los que contribuyeron!
He hecho algunas pruebas y parece que puedo ejecutar con éxito el código para dimensiones de matriz de (16 * M por 32 * N) donde M = 1,2,3,4 ... y N = 1,2 , 4,8 ... etc. ¡Si este es el caso, entonces Nvidia debería poner esto en su guía de programación! – Marm0t
¿Cuál es su GPU? – karlphillip
He probado en dos (GTX 285, Fermi uno de los más nuevos) – Marm0t