Problem in calling nested macro function

Accelerated Computing CUDA CUDA Programming and Performance
1

I have multiple macro functions in my project.

Here is the first macro defined as following

#if !defined(bswap_32)
#define bswap_32(x) ((rotr32((x), 24) & 0x00ff00ff) | (rotr32((x), 8) & 0xff00ff00))
#endif

Here is the second macro function, which calls the first macro function

#if defined(SWAP_BYTES)
#define bsw_32(p,n) \
    { int _i = (n); while(_i--) ((uint32_t*)p)[_i] = bswap_32(((uint32_t*)p)[_i]); }

At compilation level, the nvcc complains the following

error: calling a __host__ function("__bswap_32") from a __device__ function is not allowed

When i substitute the contents of bswap_32 in the second macro , it compiles just fine.

The problem occurs when i call the second macro function from the first.

2

Most likely whatever headers you’re including already define bswap_32 as a macro. So your first macro never gets defined, and the existing macro calls host function __bswap_32.

3

Resolved. thanks

4

In device code, the fastest way to perform a 32-bit bswap operation should be the following, which I would expect to compile into a single PRMT instruction:

__byte_perm (word, 0, 0x0123)