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