errors wen executing bitonic sort for a structure array I have modified the bitonic sort to sort an

Accelerated Computing CUDA CUDA Programming and Performance
1

Error 1 internal error: assertion failed: gen_expr: bad expr node kind (D:/Bld/rel/gpgpu/toolkit/r2.1/compiler/edg/EDG_3.9/src/cp_gen_be.c, line 9510) C:\Program Files\Microsoft Visual Studio 8\VC\INCLUDE\xutility 3048

Error 2 fatal error C1075: end of file found before the left parenthesis ‘(’ at ‘C:\Program Files\Microsoft Visual Studio 8\VC\INCLUDE\xutility(3015)’ was matched C:\Program Files\Microsoft Visual Studio 8\VC\INCLUDE\xutility 3017

My structure defination is as follows in bitonic.h header file which i have included in bitonic_kernel.cu
struct Frequency_Table{
char ch;
int frequency;
};

and my main code is as follows (the remaining code is as the original bitonic code)

Frequency_Table hTfreqTable[256];
for(int i = 0; i < 256; i++)
{
hTfreqTable[i].ch = char(i);
hTfreqTable[i].frequency=rand();
}

Frequency_Table *dTfreqTable;
cutilSafeCall(cudaMalloc((void**)&dTfreqTable, sizeof(Frequency_Table) * 256));
cutilSafeCall(cudaMemcpy(dTfreqTable, hTfreqTable, sizeof(Frequency_Table) * 256, cudaMemcpyHostToDevice));

FT_bitonicSort<<<1, 256, sizeof(Frequency_Table) * 256>>>(dTfreqTable);

cutilSafeCall(cudaMemcpy(hTfreqTable, dTfreqTable, sizeof(Frequency_Table) * 256, cudaMemcpyDeviceToHost));

AND MY KERNEL FUNCTIONS ARE AS FOLLOWS
device inline void FT_swap(Frequency_Table & a, Frequency_Table & B)
{
// Alternative swap doesn’t use a temporary register:
// a ^= b;
// b ^= a;
// a ^= b;

Frequency_Table tmp;
tmp.frequency= a.frequency;
a.frequency= b.frequency;
b.frequency = tmp.frequency;

tmp.ch= a.ch;
a.ch= b.ch;
b.ch = tmp.ch;

}

global static void FT_bitonicSort(Frequency_Table * values)
{
extern shared Frequency_Table FT_shared;

const unsigned int tid = threadIdx.x;

// Copy input to shared mem.
FT_shared[tid].frequency = values[tid].frequency;
FT_shared[tid].ch = values[tid].ch;

__syncthreads();

// Parallel bitonic sort.
for (unsigned int k = 2; k <= NUM; k *= 2)
{
    // Bitonic merge:
    for (unsigned int j = k / 2; j>0; j /= 2)
    {
        unsigned int ixj = tid ^ j;
        
        if (ixj > tid)
        {
            if ((tid & k) == 0)
            {
                if (FT_shared[tid].frequency > FT_shared[ixj].frequency)
                {
                    FT_swap(FT_shared[tid], FT_shared[ixj]);
                }
            }
            else
            {
                if (FT_shared[tid].frequency < FT_shared[ixj].frequency)
                {
                    FT_swap(FT_shared[tid], FT_shared[ixj]);
                }
            }
        }
        
        __syncthreads();
    }
}

// Write result.
values[tid].frequency = FT_shared[tid].frequency;
values[tid].ch = FT_shared[tid].ch;

}

please somebody help me in debuging this