Urgent help with Bitonic Sort Please help us with Bitonic Sort

What is the problem with this code? This code works for an array of 512 elements. However, I was trying to perform local bitonic sorts on every 512 elements. This means that every 512 element is sorted in ascending order. (from 0-511, 512-1023, and so on).

main.cpp:

#include <stdio.h>
#include <stdlib.h>

#define NUM 1024

extern “C” void bitonic(float * values, int nSize);

int main(int argc, char *argv)
{
float values[NUM];

for(int i = 0; i < NUM; i++)
{
    values[i] = rand();
}

bitonic (values, NUM);

bool passed = true;
for(int i = 1; i < NUM; i++)
{
    if (values[i-1] > values[i])
    {
        passed = false;
    }
}

for(int i = 0; i < NUM; i++)
{
    if (i % 512 == 0)
      printf ("\n\n\n");
    printf ("%f\t", values[i]);
}

printf( "Test %s\n", passed ? "PASSED" : "FAILED");

}

bitonic_kernel.cu:

#include <stdio.h>
#include <stdlib.h>

device inline void swap(float & a, float & B)
{
// Alternative swap doesn’t use a temporary register:
// a ^= b;
// b ^= a;
// a ^= b;

float tmp = a;
a = b;
b = tmp;

}

global static void bitonicSort(float * values, int nSize)
{
extern shared float shared;

unsigned int tid = threadIdx.x;
unsigned int i   = blockIdx.x * blockDim.x + threadIdx.x;

shared[tid] = values[i];
__syncthreads();

// Parallel bitonic sort.
for (int k = 2; k <= nSize; k *= 2)
{
    // Bitonic merge:
    for (int j = k / 2; j>0; j /= 2)
    {
        int ixj = tid ^ j;
        
        if (ixj > tid)
        {
            if ((tid & k) == 0)
            {
                if (shared[tid] > shared[ixj])
                {
                    swap(shared[tid], shared[ixj]);
                }
            }
            else
            {
                if (shared[tid] < shared[ixj])
                {
                    swap(shared[tid], shared[ixj]);
                }
            }
        }
        
        __syncthreads();
    }
}

// Write result.
values[i] = shared[tid];

}

extern “C” void bitonic(float * values, int nSize)
{
float * valuesd;

cudaMalloc((void**)&valuesd, sizeof(float) * nSize);
cudaMemcpy(valuesd, values, sizeof(float) * nSize, cudaMemcpyHostToDevice);

if (nSize <= 512)
bitonicSort<<<1, nSize, sizeof(float) * 512>>>(valuesd, nSize);
else
bitonicSort<<<ceil((float)nSize/512.0f), 512, sizeof(float) * 512>>>(valuesd, nSize);

cudaMemcpy(values, valuesd, sizeof(float) * nSize, cudaMemcpyDeviceToHost);

cudaFree(valuesd);
}

Output:

0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 523.000000 532.000000 523.000000 532.000000 931.000000 949.000000 931.000000 949.000000

0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.0000000.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 940.000000 949.000000 940.000000 949.000000 940.000000 949.000000 940.000000 949.000000 Test FAILED

but if I use 512, this is the output.

2416949.000000 6939507.000000 7684930.000000 8936987.000000 11614769.00000012260289.000000 12895151.000000 31308902.000000 35005212.000000 42999168.00000070636432.000000 76065816.000000 78012496.000000 84353896.000000 87755424.000000105575576.000000 110613200.000000 111537760.000000 112805728.000000114723504.000000 116087760.000000 135497280.000000 136495344.000000 137806864.000000 149798320.000000 150122848.000000150517568.000000 155324912.000000 155789216.000000 159259472.000000 160051520.000000 165344816.000000 168002240.000000184803520.000000 188213264.000000 195740080.000000 200747792.000000 201305632.000000 201690608.000000 212251744.000000213801968.000000 213975408.000000 221558432.000000 231602416.000000 233665120.000000 235649152.000000 238962592.000000243268144.000000 246247248.000000 255789536.000000 260152960.000000 260401248.000000 269441504.000000 269455296.000000270744736.000000 278722848.000000 282828192.000000 289700736.000000 292218016.000000 294702560.000000 296864832.000000298625216.000000 304089184.000000 317097472.000000 324763904.000000 327254592.000000 328298272.000000 332266752.000000336465792.000000 337453824.000000 337739296.000000 338888224.000000 349517440.000000 350322240.000000 352118592.000000352406208.000000 356426816.000000 364228448.000000 364319520.000000 364686240.000000 378409504.000000 382697728.000000387346496.000000 387451648.000000 389040736.000000 392035584.000000 396473728.000000 400000576.000000 402724288.000000402903168.000000 404158656.000000 407487136.000000 410409120.000000 412776096.000000 424238336.000000 425245984.000000430253408.000000 434248640.000000 437116480.000000 438792352.000000 439493440.000000 446340704.000000 452867616.000000463480576.000000 468703136.000000 474613984.000000 476667360.000000 480298496.000000 484238048.000000 485560288.000000491705408.000000 492067904.000000 496060032.000000 496987744.000000 498777856.000000 500619008.000000 501772896.000000502278624.000000 511702304.000000 515530016.000000 521595360.000000 524688224.000000 524872352.000000 532670688.000000552473408.000000 552910272.000000 555996672.000000 559301056.000000 559412928.000000 561718016.000000 570073856.000000572660352.000000 588219776.000000 593209472.000000 596516672.000000 601385664.000000 603570496.000000 608413760.000000610515456.000000 619054080.000000 619290048.000000 620145536.000000 624549824.000000 625032192.000000 628175040.000000630668864.000000 631704576.000000 635723072.000000 638422080.000000 648031296.000000 653448064.000000 653468864.000000654887360.000000 655495360.000000 655858688.000000 660260736.000000 661761152.000000 669908544.000000 672655360.000000677870464.000000 680467008.000000 697517696.000000 700108608.000000 705178752.000000 706043328.000000 707900992.000000709393600.000000 711645632.000000 711845888.000000 719346240.000000 719885376.000000 722308544.000000 726371136.000000733053120.000000 739273280.000000 740759360.000000 745425664.000000 749241856.000000 752392768.000000 756898560.000000760313728.000000 767066240.000000 771151424.000000 774044608.000000 776532032.000000 777210496.000000 777720512.000000780821376.000000 783368704.000000 791698944.000000 805750848.000000 816504768.000000 820697728.000000 822262784.000000822890688.000000 824272832.000000 828388032.000000 841148352.000000 846811136.000000 846930880.000000 846942592.000000847228032.000000 855636224.000000 859484416.000000 860516096.000000 861021504.000000 861543936.000000 882160384.000000889023296.000000 894429696.000000 915256192.000000 917679296.000000 927612928.000000 932026304.000000 933110208.000000937370176.000000 939819584.000000 943947712.000000 945117248.000000 950390848.000000 959997312.000000 964445888.000000968338112.000000 971899200.000000 982907008.000000 982936768.000000 990892928.000000 995706880.000000 1012502976.000000 1017679552.000000 1023457728.000000 1025202368.000000 1034514496.000000 1034949312.000000 1036140800.000000 1037127808.000000 1046741248.000000 1057418432.000000 1057467584.000000 1059961408.000000 1063958016.000000 1065103360.000000 1065311680.000000 1067854528.000000 1069755904.000000 1070575296.000000 1081174272.000000 1096689792.000000 1100661376.000000 1101513984.000000 1102520064.000000 1111783936.000000 1113502208.000000 1117142656.000000 1120048768.000000 1125898112.000000 1129033344.000000 1129566464.000000 1131176192.000000 1137949952.000000 1139901440.000000 1141616128.000000 1143408256.000000 1144278016.000000 1144522496.000000 1159126528.000000 1172755584.000000 1175526272.000000 1176911360.000000 1186090368.000000 1186452608.000000 1189641472.000000 1192707584.000000 1194953856.000000 1197352320.000000 1219408000.000000 1231192320.000000 1237379072.000000 1238433408.000000 1239036032.000000 1242608896.000000 1244316416.000000 1253207680.000000 1255179520.000000 1264095104.000000 1266235136.000000 1272469760.000000 1273911936.000000 1275373696.000000 1277849984.000000 1280321664.000000 1285228800.000000 1295166336.000000 1303455744.000000 1307565952.000000 1308044928.000000 1309383296.000000 1315209216.000000 1315634048.000000 1328104320.000000 1329132160.000000 1329202944.000000 1330573312.000000 1335354368.000000 1335939840.000000 1343606016.000000 1344247680.000000 1346811264.000000 1350489984.000000 1350573824.000000 1351797376.000000 1359512192.000000 1365180544.000000 1369133056.000000 1369321856.000000 1371499392.000000 1373226368.000000 1374344064.000000 1376710144.000000 1380171648.000000 1387036160.000000 1389867264.000000 1395132032.000000 1395235072.000000 1398295552.000000 1402586752.000000 1402961664.000000 1407392256.000000 1409959680.000000 1411154304.000000 1411549696.000000 1414647680.000000 1414829184.000000 1415505408.000000 1424268928.000000 1431419392.000000 1432114560.000000 1433102848.000000 1433925888.000000 1447267584.000000 1448703744.000000 1450573568.000000 1450956032.000000 1455590912.000000 1469261952.000000 1469348096.000000 1469834496.000000 1470503424.000000 1472576384.000000 1472713728.000000 1473144448.000000 1473442048.000000 1474612352.000000 1476153216.000000 1477171072.000000 1486222848.000000 1494613760.000000 1498617600.000000 1501252992.000000 1503885184.000000 1504569856.000000 1516266752.000000 1520223232.000000 1529195776.000000 1540383488.000000 1543755648.000000 1544617472.000000 1545032448.000000 1548233344.000000 1548348160.000000 1555319296.000000 1562402304.000000 1566288768.000000 1569229312.000000 1572276992.000000 1573363328.000000 1583571072.000000 1585990400.000000 1586903168.000000 1590079488.000000 1597322368.000000 1600028672.000000 1605894400.000000 1605908224.000000 1610120704.000000 1622597504.000000 1626276096.000000 1631518208.000000 1632621696.000000 1633108096.000000 1635905408.000000 1642548864.000000 1649760512.000000 1653377408.000000 1656478080.000000 1662739712.000000 1662981760.000000 1665947520.000000 1671294848.000000 1676902016.000000 1679895424.000000 1681692800.000000 1682085248.000000 1687926656.000000 1703964672.000000 1704365056.000000 1713258240.000000 1714636928.000000 1722060032.000000 1726956416.000000 1734575232.000000 1737518976.000000 1739000704.000000 1749698560.000000 1750003072.000000 1756915712.000000 1760243584.000000 1760281984.000000 1775473792.000000 1776808960.000000 1780172288.000000 1780695808.000000 1781999744.000000 1782280576.000000 1782436864.000000 1784639488.000000 1789366144.000000 1789376384.000000 1795519104.000000 1797073920.000000 1799878144.000000 1801979776.000000 1804289408.000000 1812282112.000000 1820388480.000000 1823089408.000000 1827336320.000000 1841585792.000000 1843993344.000000 1850952960.000000 1856669184.000000 1857962496.000000 1858721920.000000 1866000128.000000 1869470080.000000 1875335936.000000 1876855552.000000 1884167680.000000 1884661248.000000 1886087040.000000 1887658368.000000 1889947136.000000 1892066560.000000 1896306688.000000 1900553600.000000 1908518784.000000 1909002880.000000 1911165184.000000 1911760000.000000 1914544896.000000 1918502656.000000 1927495936.000000 1937477120.000000 1939964416.000000 1941690368.000000 1943327744.000000 1947346560.000000 1950955904.000000 1953443328.000000 1956297600.000000 1957747840.000000 1960709888.000000 1967513984.000000 1967681152.000000 1973388032.000000 1973594368.000000 1974806400.000000 1975960320.000000 1977648512.000000 1982275840.000000 1983690368.000000 1984210048.000000 1987230976.000000 1989806336.000000 1998898816.000000 2001100544.000000 2001229952.000000 2002495488.000000 2006812032.000000 2007905792.000000 2025187200.000000 2027907712.000000 2038664320.000000 2040332928.000000 2040651392.000000 2044897792.000000 2053999872.000000 2058657152.000000 2077486720.000000 2084420864.000000 2086206720.000000 2089018496.000000 2103318784.000000 2112255744.000000 2113903872.000000 2114738048.000000 2114937728.000000 2118422016.000000 2130794368.000000 2142756992.000000 2145174016.000000 2147469824.000000 Test PASSED

Thank you very much. I think I missed something with shared memory.
:wacko: :( <img src=‘http://hqnveipbwb20/public/style_emoticons/<#EMO_DIR#>/crying.gif’ class=‘bbc_emoticon’ alt=’:’(’ />

I tried my best and it is sort of working now.

Shared memory is shared only per block, so you have to bear in mind maximum threads in block = 512

// Parallel bitonic sort.
for (int k = 2; k <= 512; k *= 2)

I recommend using k<<=1) // instead of k *= 2)
and j>>=1) // instead of j /= 2)

for NUM=1024 you finally get two ascending sequences.

best regards
.m.