cudaArray & texture binding... Need help

Hi,

I’ve associated an array to a texture, but I can’t navigate correctly in this.

i don’t understand why.

If someone can help me, thanks in advance.

Here is my code :

typedef struct __align__(16) _paquetInfoGPU

{

	unsigned long int posStart;		// paquet begins at this offset 'posStart' from the beginning of the file

	unsigned long int length;		// Length of the paquet

	unsigned long int selector;		// paquet type

} paquetInfoGPU;

////////////////////////////////////////////////////// VARIABLES //////////////////////////////////////////////////////

// Variables on GLOBAL MEMORY

// --------------------------

void* d_BufferIn = NULL;												// Buffer where are stored paquets to analyze (on Device)

void* h_BufferIn = NULL;												// Buffer where are stored paquets to analyze (on Host)

// Variables on TEXTURE MEMORY

// ---------------------------

texture<ulong1, 2, cudaReadModeElementType> texture_listpaquetinfo;		// Texture binded to paquets information => d_list_paquetinfo

cudaArray* d_list_paquetinfo = NULL;								// Informations about the paquets to extract : paquetInfoGPU[] : start position, length, seletor

////////////////////////////////////////////////////// CONSTANTS //////////////////////////////////////////////////////

#define NBR_PAQUET_EXTRACT	10000

#define DECPLAN_MAX_SIZE	52

#define SHARED_MEM_SIZE	296

//////////////////////////////////////////////////////////////////////////////

///

///	Copy paquets information (start, length, type) to the texture memory

///

///	Parameters :

///		@param[in] nbPaquet - Number of paquets in the list

///		@param[in] pListPaquetInfoGPU - Paquets List

///

///	Return :

///		void

///

///	Source : Host

//////////////////////////////////////////////////////////////////////////////

void

copybindPaquetInfo(const unsigned long int nbPaquet, paquetInfoGPU* pListPaquetInfoGPU)

{

	cudaChannelFormatDesc chDescPaquet = {0};				// Descriptor

	if (d_list_paquetinfo == NULL)

	{

		// PaquetInfo : Create on the device the array

		// --------------------------------------------

		chDescPaquet = cudaCreateChannelDesc<ulong1>();

		cutilSafeCall(cudaMallocArray(&d_list_paquetinfo, &chDescPaquet, sizeof(paquetInfoGPU), nbPaquet));

	}

	// PaquetInfo : Copy data to the array

	// -----------------------------------

	cutilSafeCall(cudaMemcpyToArray(d_list_paquetinfo, 0, 0, (void*)pListPaquetInfoGPU, nbPaquet*sizeof(paquetInfoGPU), cudaMemcpyHostToDevice));

	texture_listpaquetinfo.addressMode[0] = cudaAddressModeClamp;

	texture_listpaquetinfo.filterMode = cudaFilterModePoint;

	texture_listpaquetinfo.normalized = false;

	// PaquetInfo : Bind data with texture

	// ---------------------------------

	cutilSafeCall(cudaBindTextureToArray(texture_listpaquetinfo, d_list_paquetinfo, chDescPaquet));

}

__global__ void

kernel_copyIntoSharedMem(const void* pBufferIn)

{

	extern __shared__ unsigned char shareBufferIn[];	// paquet content

	// Retrieve paquet information

	// ---------------------------

	unsigned long posStart	= ulong1(	tex2D(texture_listpaquetinfo, 0, blockIdx.x)	).x;							// paquet begins at this offset 'posStart' from the beginning

	unsigned long length		= ulong1(	tex2D(texture_listpaquetinfo, 1, blockIdx.x)	).x / sizeof(unsigned int);		// paquet length

	// Copy paquet content into the block shared memory

	// ------------------------------------------------

	if (threadIdx.x < length)

	{

		*((unsigned int*)((unsigned int)&shareBufferIn + (threadIdx.x * sizeof(unsigned int)))) = *((unsigned int*)((unsigned int)pBufferIn + posStart + (threadIdx.x * sizeof(unsigned int))));

	}

}

int main(int argc, char *argv[])

{

		paquetInfoGPU*	pListPaquetInfo = NULL;										// List of paquets information

		// open the card

		// analyse a binary file, and create  an array of "paquetInfoGPU"

		// copy these info into the GPU

		copybindPaquetInfo(NBR_PAQUET_EXTRACT, pListPaquetInfo);

		// copy data to GPU

		copyDataToGPU(streamTM, pListPaquetInfo, NBR_PAQUET_EXTRACT, &size);

		// Execute kernel : copy data into shared mem

		kernel_copyIntoSharedMem <<< NBR_PAQUET_EXTRACT, DECPLAN_MAX_SIZE, SHARED_MEM_SIZE >>>(d_BufferIn);

	// Free resources

	// --------------

	gpumemory_free();

	free(pListPaquetInfo);

	fclose(streamTM);

}

Here is an example of data that are in pListPaquetInfo (array of paquetInfoGPU) :

[   0,  196, 81, ... ]

[196, 196, 81, ... ]

[392, 196, 81, ... ]

[588, 196, 81, ... ]

[784, 196, 81, ... ]

(where … is the dummy data generated by the structure alignment

So, in the function “kernel_copyIntoSharedMem”, with the code :

posStart = ulong1( tex2D(texture_listpaquetinfo, 0, blockIdx.x) ).x; // paquet begins at this offset ‘posStart’ from the beginning

length = ulong1( tex2D(texture_listpaquetinfo, 1, blockIdx.x) ).x / sizeof(unsigned int); // paquet length

I obtain the value posStart = 784, for blockIdx.x=1

why?

I just want to obtain posStart = 196, for blockIdx.x = 1

If i’ve well understand, it works as is : tex2D(texture_listpaquetinfo, x, y), where 0<x<3 in my case and 0<y<NBR_PAQUET_EXTRACT-1

If someone can help me with my problem, because I’m lost.

Thanks in advance for your help

Ok, I’ve found the solution : the texture was not well written, as the copybindPaquetInfo function.

In this one, I’ve fixed the following lines : cudaMallocArray(dest, src, number of element X, number of element Y)…

Here is the working code.

texture<unsigned long int, 2, cudaReadModeElementType> texture_listpaquetinfo;		// Texture binded to paquets information => d_list_paquetinfo

void

copybindPaquetInfo(const unsigned long int nbPaquet, paquetInfoGPU* pListPaquetInfoGPU)

{

	cudaChannelFormatDesc chDescPaquet = {0};				// Content of the array

	if (d_list_paquetinfo == NULL)

	{

		// PaquetInfo : Create on the device the array

		// --------------------------------------------

		chDescPaquet = cudaCreateChannelDesc(32,0,0,0, cudaChannelFormatKindUnsigned);

		cutilSafeCall(cudaMallocArray(&d_list_paquetinfo, &chDescPaquet, sizeof(paquetInfoGPU)/sizeof(unsigned long int), nbPaquet));

	}

	// PaquetInfo : Copy data to the array

	// -----------------------------------

	cutilSafeCall(cudaMemcpyToArray(d_list_paquetinfo, 0, 0, (void*)pListPaquetInfoGPU, nbPaquet*(sizeof(paquetInfoGPU)/sizeof(unsigned long int))*sizeof(unsigned long int), cudaMemcpyHostToDevice));

	// PaquetInfo : Bind data with texture

	// ---------------------------------

	texture_listpaquetinfo.addressMode[0] = cudaAddressModeClamp;

	texture_listpaquetinfo.addressMode[1] = cudaAddressModeClamp;

	texture_listpaquetinfo.filterMode = cudaFilterModePoint;

	texture_listpaquetinfo.normalized = 0;

	cutilSafeCall(cudaBindTextureToArray(texture_listpaquetinfo, d_list_paquetinfo, chDescPaquet));

}