     if (threadIdx.x == 0)
     {
         d_out[TILE_SIZE] = prefix_op.prefix;
-    }
-    else
-    {
-        // In a separate conditional as a workaround for SM10 Win32 bug where we write the wrong prefix
         *d_elapsed = (start > stop) ? start - stop : stop - start;
     }
 }

        Test<128, 2, PREFIX_AGGREGATE, BLOCK_SCAN_RAKING>(SEQ_INC, Sum(), NullType(), prefix, CUB_TYPE_STRING(Sum<TestFoo>));
    }
    else
    {
        // Repeat test sequence
        for (int i = 0; i <= g_repeat; ++i)
        {
            // Run tests for different threadblock sizes
            Test<17>();
            Test<32>();
            Test<62>();
            Test<65>();
//            Test<96>();             // TODO: file bug for UNREACHABLE error for Test<96, 9, BASIC, BLOCK_SCAN_RAKING>(UNIFORM, Sum(), NullType(), make_ulonglong2(17, 21), CUB_TYPE_STRING(Sum<ulonglong2>));
            Test<128>();
        }
    }


    int ITEMS_PER_THREAD>
void Test(GenMode gen_mode)
{
/*
    // primitive
    Test<BLOCK_THREADS, ITEMS_PER_THREAD>(gen_mode, Sum(), (unsigned char) 0, (unsigned char) 99, CUB_TYPE_STRING(Sum<unsigned char>));
    Test<BLOCK_THREADS, ITEMS_PER_THREAD>(gen_mode, Sum(), (unsigned short) 0, (unsigned short) 99, CUB_TYPE_STRING(Sum<unsigned short>));
