Come posso:Come leggere con successo da una texture 2D
- Bind cudaMallocPitch memoria galleggiante ad un riferimento di texture 2D
- copiare alcuni dati host a matrice 2D sul dispositivo
- Aggiungere uno a il riferimento di trama e scrivere su a.) l'array Pitch 2D OR b.) scrivere su un array di memoria lineare
- Leggere la risposta indietro e visualizzarla.
Di seguito è riportato un codice che dovrebbe eseguire ciò. Si noti che per le dimensioni degli array NxN, il mio codice funziona. Per NxM dove N! = M, il mio codice morde la polvere (non il risultato corretto). Se puoi risolvere questo problema ti assegnerò 1 internet (fornitura limitata). Forse sono pazzo, ma secondo la documentazione questo dovrebbe funzionare (e funziona con array quadrati!). Il codice allegato dovrebbe essere eseguito con 'nvcc whateveryoucallit.cu -o runit'.
L'aiuto è apprezzato!
#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);
}
Modifica 17 ottobre: quindi non ho ancora trovato una soluzione a questo problema. Nvidia è abbastanza silenzioso su questo sembra che sia anche il mondo. Ho trovato una soluzione alternativa utilizzando mem condivisa, ma se qualcuno ha una soluzione di texture sarei molto felice.
Modifica Octoboer 26: Ancora nessuna soluzione, ma ancora interessata a uno se qualcuno lo sa.
Modifica 26 luglio: Wow sono trascorsi 9 mesi e non ho tenuto conto della risposta corretta per tutto il tempo. Il trucco era:
if (idx < width && idy < height){//.... code }
Come era stato sottolineato prima. Grazie a tutti quelli che hanno contribuito!
Ho eseguito un piccolo test e sembra che possa eseguire correttamente il codice per le dimensioni dell'array di (16 * M per 32 * N) dove M = 1,2,3,4 ... e N = 1,2 , 4,8 ... ecc. Se questo è il caso, Nvidia dovrebbe metterlo in questa guida di programmazione! – Marm0t
Qual è la tua GPU? – karlphillip
Ho provato su due (GTX 285, Fermi uno dei più nuovi) – Marm0t