How to call _popc

Dec 12, 2012 at 10:48 PM

How does one call _popc from CUDAfy?

Dec 13, 2012 at 2:09 PM
Edited Dec 13, 2012 at 2:18 PM

Wil, if you're willing to download and compile the latest version of cudafy, you'll have a solution for your problem.

GThread has a new static method, called GThread.InsertCode. You can use it to emit a __popc instruction. Example:

int a = 10;

int b = 0; // need to initialize it with something.

GThread.InsertCode("{0} = __popc({1})", b, a); // will generate the CUDA instruction "b = __popc(a);"

 

cheers

 

P.S. Nick, this wasn't me, I swear!

 

Coordinator
Dec 13, 2012 at 4:37 PM

Ha, ye olde trolls are lurking around ;-)

Actually funny how these things work out. I'll be reviewing the code you sent and most like will add shortly. There are indeed a bunch of useful functions that should be in there.

Dec 14, 2012 at 1:14 AM

Thanks a lot Pedritolo1!  I'll give it a shot.  I had worked around it by using [CudafyDummy] but this is much better.

Thanks,

Wil

Dec 14, 2012 at 2:02 AM

It worked!  Thanks again!

I had a bit of a problem because of the name mangling so I had to do the entire method calling InsertCode.  Don't know if there's an easier way.

        [Cudafy]
        public static int PopulationCount(ulong value)
        {
            GThread.InsertCode("int* LValue = (int*)&value;");
            GThread.InsertCode("return __popc(LValue[0]) + __popc(LValue[1]);");
            return 0;
        }
Result:
__device__ int PopulationCount(unsigned long long value)
{
 int* LValue = (int*)&value;
 return __popc(LValue[0]) + __popc(LValue[1]);
 return 0;
}
Dec 14, 2012 at 11:36 AM

Wil, why not use __popcll instead?

GThread.InsertCode("return __popcll({0});", value);

cheers

Dec 14, 2012 at 11:40 AM

Thanks pedritolo1!  I didn't know about __popcll!

The params version of InsertCode won't work because it can't take a pointer :(

Thanks,

Wil

Dec 14, 2012 at 11:43 AM

Never mind, with __popcll I don't need the extra variable so the params version works then.

Thanks again