Unrecognized Token in kernel

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,

The following kernel is given in the document Scan using CUDA on page number 10 to 12.

When I compile the code I get the following errors:

[root@kiran Download]# nvcc Scan13.cu -deviceemu -o Scan

Scan13.cu(59): error: unrecognized token

Scan13.cu(59): error: expected a “]”

Scan13.cu(59): error: unrecognized token

Scan13.cu(59): error: unrecognized token

4 errors detected in the compilation of “/tmp/tmpxft_00001307_00000000-4_Scan13.cpp1.ii”.

I have folllwoing questions here:

1- In the line : #define CONFLICT_FREE_OFFSET(n) \ ((n) >> NUM_BANKS + (n) >> (2 * LOG_NUM_BANKS))
what is the meaning ""? should it not be “/” ?

2- I am getting all the four errors in the following line of the code:

if (thid==0) { temp[n - 1 + CONFLICT_FREE_OFFSET(n - 1)] = 0; } //// clear the last element

3 of these 4 errors say : “unrecognized token”. What is the meaning of token here with respect to line 59
where I am clearing last element. And why the same error is coming up
three times when I think the instruction is executed only once during compilation.

3- I am unable to understand “/tmp/tmpxft_00001307_00000000-4_Scan13.cpp1.ii”. What is tmp, tmpxft_00001307_00000000-4_Scan13.cpp1.ii
What is .ii file etc?

4-why "expected a “]” error is coming in the line 59 i.e:
if (thid==0) { temp[n - 1 + CONFLICT_FREE_OFFSET(n - 1)] = 0; } //// clear the last element

Thanks in advance!!

==============

The kernel along with the Macro is given below (please notice the bold lines):

#define NUM_BANKS 16
#define LOG_NUM_BANKS 4
#ifdef ZERO_BANK_CONFLICTS

#define CONFLICT_FREE_OFFSET(n) \ ((n) >> NUM_BANKS + (n) >> (2 * LOG_NUM_BANKS))#else

#define CONFLICT_FREE_OFFSET(n) ((n) >> LOG_NUM_BANKS)
#endif

global void prescan(float *g_odata, float *g_idata, int n)
{
extern shared float temp;// allocated on invocation
int thid = threadIdx.x;
int offset = 1;

int ai = thid;
int bi = thid + (n/2);
int bankOffsetA = CONFLICT_FREE_OFFSET(ai);
int bankOffsetB = CONFLICT_FREE_OFFSET(ai);
temp[ai + bankOffsetA] = g_idata[ai];
temp[bi + bankOffsetB] = g_idata[bi];

for (int d = n>>1; d > 0; d >>= 1) // build sum in place up the tree
{
__syncthreads();
if (thid < d)
{

int ai = offset*(2thid+1)-1;
int bi = offset(2*thid+2)-1;
ai += CONFLICT_FREE_OFFSET(ai);
bi += CONFLICT_FREE_OFFSET(bi);

temp[bi] += temp[ai];
}
offset *= 2;
}

if (thid==0) { temp[n - 1 + CONFLICT_FREE_OFFSET(n - 1)] = 0; } //// clear the last element

for (int d = 1; d < n; d *= 2) // traverse down tree & build scan
{
offset >>= 1;
__syncthreads();
if (thid < d)
{

int ai = offset*(2thid+1)-1;
int bi = offset(2*thid+2)-1;
ai += CONFLICT_FREE_OFFSET(ai);
bi += CONFLICT_FREE_OFFSET(bi);

float t = temp[ai];

temp[ai] = temp[bi];
temp[bi] += t;

}
}

__syncthreads();

g_odata[ai] = temp[ai + bankOffsetA];
g_odata[bi] = temp[bi + bankOffsetB];

}

2

You should fix this

<b>#define CONFLICT_FREE_OFFSET(n) \ ((n) >> NUM_BANKS + (n) >> (2 * LOG_NUM_BANKS))</b>#else

#define CONFLICT_FREE_OFFSET(n) ((n) >> LOG_NUM_BANKS)

#endif

to this:

#define CONFLICT_FREE_OFFSET(n) \ 

  ((n) >> NUM_BANKS + (n) >> (2 * LOG_NUM_BANKS))

#else

  #define CONFLICT_FREE_OFFSET(n) ((n) >> LOG_NUM_BANKS)

#endif

plain macro typo… you’re missing a line-feed there

eyal

3

Thanks a lot eyal!!

Can you tell me why I am getting “Unrecognised Token” error three times and why it is expecting “]” in line 59 wher I am clearing the last element?

4

Its a compiler complain - nothing to do with CUDA :)

you’ll get an error for every line that has the faulty macro.

eyal

5

Ok. Got it!!!

By the way do you think the way temp has been declared has nothing to do with the error and there is no error in the decerlarion:

extern shared float temp;// allocated on invocation

6

What’s the problem with this line? its not a macro if thats what you mean…

eyal