private static uint getTop5Cards(GThread thread, int ranks) { int c = 31 - thread.clz(ranks); int top5cards = c << TOP_CARD_SHIFT; ranks ^= 1 << c; c = 31 - thread.clz(ranks); top5cards += c << SECOND_CARD_SHIFT; ranks ^= 1 << c; c = 31 - thread.clz(ranks); top5cards += c << THIRD_CARD_SHIFT; ranks ^= 1 << c; c = 31 - thread.clz(ranks); top5cards += c << FOURTH_CARD_SHIFT; ranks ^= 1 << c; c = 31 - thread.clz(ranks); top5cards += c << FIFTH_CARD_SHIFT; return((uint)top5cards); }
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; } }