jacobtolar

  shr.s32   %r6, %r5, 1;           ; r6 = bc /2
  mov.u32   %r7, %ctaid.y;         
  mul.lo.u32  %r8, %r7, 2;         ; r8 = blkIdx.y * 2
  add.u32   %r9, %r6, %r8;         ; mbh = r6 + r8
  and.b32   %r10, %r5, 1;          ; r10 = bc & 1
  mov.u32   %r11, %ctaid.x;
  mul.lo.u32  %r12, %r11, 2;       ; r12 = blkIdx.x * 2
  add.u32   %r13, %r10, %r12;      ; mbw = r10 + r12

jacobtolar

  mbw = (blockIdx.x<<1) + (bcount & 0x0001); 
  mbh = (blockIdx.y<<1) + (bcount>>1);

jacobtolar

  mov.u32   %r6, 0;
  setp.eq.s32   %p1, %r5, %r6;         ; bc == 0?
  @%p1 bra  $Lt_0_258;                  
  mov.u32   %r7, 1;                    
  setp.eq.s32   %p2, %r5, %r7;         ; bc == 1?
  @%p2 bra  $Lt_0_770;
  mov.u32   %r8, 2;                    
  setp.eq.s32   %p3, %r5, %r8;         ; bc == 2?
  @%p3 bra  $Lt_0_1026;                
  mov.u32   %r9, 3;

jacobtolar

// set bcount somewhere above; assume 0 >= bcount <= 3:
// 
switch (bcount) {
 case 0: mbh = blockIdx.y*2;
         mbw=blockIdx.x*2;
         break;
 case 1: mbh = blockIdx.y*2;
         mbw=blockIdx.x*2+1;
         break;
 case 2: mbh = blockIdx.y*2+1;