public static void integerIntrinsicsInt32(GThread thread, int[] input, int[] output)
{
int i = 0;
int x = 0;
output[i++] = thread.popcount((uint)0x55555555); // 16
output[i++] = thread.clz(0x1FFFE00); // 7
output[i++] = (int)thread.mul24((int)0x00000042, (int)0x00000042);
output[i++] = (int)thread.umul24((uint)0x00000042, (uint)0x00000042);
output[i++] = (int)thread.mulhi(0x0AFFEEDD, 0x0DEEFFAA);
output[i++] = (int)thread.umulhi((uint)0x0AFFEEDD, (uint)0x0DEEFFAA);
}
public static void integerIntrinsicsInt32(GThread thread, int[] input, int[] output)
{
int i = 0;
int x = 0;
output[i++] = thread.popcount((uint)0x55555555);// 16
output[i++] = thread.clz(0x1FFFE00); // 7
output[i++] = (int)thread.mul24((int)0x00000042, (int)0x00000042);
output[i++] = (int)thread.umul24((uint)0x00000042, (uint)0x00000042);
output[i++] = (int)thread.mulhi(0x0AFFEEDD, 0x0DEEFFAA);
output[i++] = (int)thread.umulhi((uint)0x0AFFEEDD, (uint)0x0DEEFFAA);
}
private static void evaluate(GThread thread, ulong[] hands, int numCards, int numHands, uint[] oranks)
{
int tid = thread.threadIdx.x + thread.blockIdx.x * thread.blockDim.x +
(thread.threadIdx.y + thread.blockIdx.y * thread.blockDim.y) * thread.gridDim.x * thread.blockDim.x;
uint retval = 0, four_mask, three_mask, two_mask;
if (tid < numHands)
{
ulong cards = hands[tid];
// Seperate out by suit
uint sc = (uint)(cards & 8191ul); //(uint)((cards >> (CLUB_OFFSET)) & 0x1fffUL);
ulong c_ = cards / 8192ul;
uint sd = (uint)(c_ & 8191ul); //(uint)((cards >> (DIAMOND_OFFSET)) & 0x1fffUL);
c_ /= 8192ul;
uint sh = (uint)(c_ & 8191ul); //((cards >> (HEART_OFFSET)) & 0x1fffUL);
c_ /= 8192ul;
uint ss = (uint)(c_ & 8191ul); // ((cards >> (SPADE_OFFSET)) & 0x1fffUL);
uint ranks = sc | sd | sh | ss;
int n_ranks = thread.popcount(ranks);
int n_dups = numCards - n_ranks;
if (thread.popcount(sc) >= 5)
{
if ((sc & S10JQKA) == S10JQKA)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 10;
}
else if ((sc & S910JQK) == S910JQK)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 9;
}
else if ((sc & S8910JQ) == S8910JQ)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 8;
}
else if ((sc & S78910J) == S78910J)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 7;
}
else if ((sc & S678910) == S678910)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 6;
}
else if ((sc & S56789) == S56789)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 5;
}
else if ((sc & S45678) == S45678)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 4;
}
else if ((sc & S34567) == S34567)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 3;
}
else if ((sc & S23456) == S23456)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 2;
}
else if ((sc & SA2345) == SA2345)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 1;
}
else
{
retval = HANDTYPE_VALUE_FLUSH + (uint)getTop5Cards(thread, (int)sc);
}
}
else if (thread.popcount(sd) >= 5)
{
if ((sd & S10JQKA) == S10JQKA)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 10;
}
else if ((sd & S910JQK) == S910JQK)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 9;
}
else if ((sd & S8910JQ) == S8910JQ)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 8;
}
else if ((sd & S78910J) == S78910J)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 7;
}
else if ((sd & S678910) == S678910)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 6;
}
else if ((sd & S56789) == S56789)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 5;
}
else if ((sd & S45678) == S45678)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 4;
}
else if ((sd & S34567) == S34567)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 3;
}
else if ((sd & S23456) == S23456)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 2;
}
else if ((sd & SA2345) == SA2345)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 1;
}
else
{
retval = HANDTYPE_VALUE_FLUSH + (uint)getTop5Cards(thread, (int)sc);
}
}
else if (thread.popcount(sh) >= 5)
{
if ((sh & S10JQKA) == S10JQKA)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 10;
}
else if ((sh & S910JQK) == S910JQK)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 9;
}
else if ((sh & S8910JQ) == S8910JQ)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 8;
}
else if ((sh & S78910J) == S78910J)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 7;
}
else if ((sh & S678910) == S678910)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 6;
}
else if ((sh & S56789) == S56789)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 5;
}
else if ((sh & S45678) == S45678)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 4;
}
else if ((sh & S34567) == S34567)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 3;
}
else if ((sh & S23456) == S23456)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 2;
}
else if ((sh & SA2345) == SA2345)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 1;
}
else
{
retval = HANDTYPE_VALUE_FLUSH + (uint)getTop5Cards(thread, (int)sc);
}
}
else if (thread.popcount(ss) >= 5)
{
if ((ss & S10JQKA) == S10JQKA)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 10;
}
else if ((ss & S910JQK) == S910JQK)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 9;
}
else if ((ss & S8910JQ) == S8910JQ)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 8;
}
else if ((ss & S78910J) == S78910J)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 7;
}
else if ((ss & S678910) == S678910)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 6;
}
else if ((ss & S56789) == S56789)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 5;
}
else if ((ss & S45678) == S45678)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 4;
}
else if ((ss & S34567) == S34567)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 3;
}
else if ((ss & S23456) == S23456)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 2;
}
else if ((ss & SA2345) == SA2345)
{
retval = HANDTYPE_VALUE_STRAIGHTFLUSH + 1;
}
else
{
retval = HANDTYPE_VALUE_FLUSH + (uint)getTop5Cards(thread, (int)sc);
}
}
else
{
if ((ranks & S10JQKA) == S10JQKA)
{
retval = HANDTYPE_VALUE_STRAIGHT + 10;
}
else if ((ranks & S910JQK) == S910JQK)
{
retval = HANDTYPE_VALUE_STRAIGHT + 9;
}
else if ((ranks & S8910JQ) == S8910JQ)
{
retval = HANDTYPE_VALUE_STRAIGHT + 8;
}
else if ((ranks & S78910J) == S78910J)
{
retval = HANDTYPE_VALUE_STRAIGHT + 7;
}
else if ((ranks & S678910) == S678910)
{
retval = HANDTYPE_VALUE_STRAIGHT + 6;
}
else if ((ranks & S56789) == S56789)
{
retval = HANDTYPE_VALUE_STRAIGHT + 5;
}
else if ((ranks & S45678) == S45678)
{
retval = HANDTYPE_VALUE_STRAIGHT + 4;
}
else if ((ranks & S34567) == S34567)
{
retval = HANDTYPE_VALUE_STRAIGHT + 3;
}
else if ((ranks & S23456) == S23456)
{
retval = HANDTYPE_VALUE_STRAIGHT + 2;
}
else if ((ranks & SA2345) == SA2345)
{
retval = HANDTYPE_VALUE_STRAIGHT + 1;
}
}
if (retval == 0)
{
switch (n_dups)
{
case 0:
retval = HANDTYPE_VALUE_HIGHCARD + (uint)getTop5Cards(thread, (int)ranks);
break;
case 1:
{
/* It's a one-pair hand */
uint t, kickers;
two_mask = ranks ^ (sc ^ sd ^ sh ^ ss);
retval = (uint)(HANDTYPE_VALUE_PAIR + ((31 - thread.clz((int)two_mask)) << TOP_CARD_SHIFT));
t = ranks ^ two_mask; /* Only one bit set in two_mask */
/* Get the top five cards in what is left, drop all but the top three
* cards, and shift them by one to get the three desired kickers */
kickers = ((uint)getTop5Cards(thread, (int)t) >> CARD_WIDTH) & ~FIFTH_CARD_MASK;
retval += kickers;
break;
}
case 2:
/* Either two pair or trips */
two_mask = ranks ^ (sc ^ sd ^ sh ^ ss);
if (two_mask != 0)
{
uint t = ranks ^ two_mask; /* Exactly two bits set in two_mask */
retval = (uint)(HANDTYPE_VALUE_TWOPAIR
+ (getTop5Cards(thread, (int)two_mask)
& (TOP_CARD_MASK | SECOND_CARD_MASK))
+ ((31 - thread.clz((int)t)) << THIRD_CARD_SHIFT));
break;
}
else
{
uint t, second;
three_mask = ((sc & sd) | (sh & ss)) & ((sc & sh) | (sd & ss));
retval = (uint)(HANDTYPE_VALUE_TRIPS + ((31 - thread.clz((int)three_mask)) << TOP_CARD_SHIFT));
t = ranks ^ three_mask; /* Only one bit set in three_mask */
second = (uint)(31 - thread.clz((int)t));
retval += (second << SECOND_CARD_SHIFT);
t ^= (1U << (int)second);
retval += (uint)((31 - thread.clz((int)t)) << THIRD_CARD_SHIFT);
break;
}
default:
/* Possible quads, fullhouse, straight or flush, or two pair */
four_mask = sh & sd & sc & ss;
if (four_mask != 0)
{
uint tc = (uint)(31 - thread.clz((int)four_mask));
retval = (uint)(HANDTYPE_VALUE_FOUR_OF_A_KIND
+ (tc << TOP_CARD_SHIFT)
+ ((31 - thread.clz((int)(ranks ^ (1U << (int)tc)))) << SECOND_CARD_SHIFT));
break;
}
;
/* Technically, three_mask as defined below is really the set of
* bits which are set in three or four of the suits, but since
* we've already eliminated quads, this is OK */
/* Similarly, two_mask is really two_or_four_mask, but since we've
* already eliminated quads, we can use this shortcut */
two_mask = ranks ^ (sc ^ sd ^ sh ^ ss);
if (thread.popcount(two_mask) != n_dups)
{
/* Must be some trips then, which really means there is a
* full house since n_dups >= 3 */
uint tc, t;
three_mask = ((sc & sd) | (sh & ss)) & ((sc & sh) | (sd & ss));
retval = HANDTYPE_VALUE_FULLHOUSE;
tc = (uint)(31 - thread.clz((int)three_mask));
retval += (tc << TOP_CARD_SHIFT);
t = (two_mask | three_mask) ^ (1U << (int)tc);
retval += (uint)(31 - thread.clz((int)t) << SECOND_CARD_SHIFT);
break;
}
;
if (retval == 0)
{
/* Must be two pair */
uint top, second;
retval = HANDTYPE_VALUE_TWOPAIR;
top = (uint)(31 - thread.clz((int)two_mask));
retval += (top << TOP_CARD_SHIFT);
second = (uint)(31 - thread.clz((int)(two_mask ^ (1 << (int)top))));
retval += (second << SECOND_CARD_SHIFT);
retval += (uint)((31 - thread.clz((int)(ranks ^ (1U << (int)top) ^ (1 << (int)second)))) << THIRD_CARD_SHIFT);
}
break;
}
}
/*
* By the time we're here, either:
* 1) there's no five-card hand possible (flush or straight), or
* 2) there's a flush or straight, but we know that there are enough
* duplicates to make a full house / quads possible.
*/
oranks[tid] = retval;
}
}
