OSX Multibeam OpenCL question

[Raistmer]

last call is oclFFT one.

[Joe Fox]

I will be back to working on this today and should be able to get the parameters for the oclFFT call that is causing the error. Again to reiterate though, the OPEN_CL_INTEL version works correctly with accurate test results when run on an nvidia as well as AMD based Mac machines.

[Raistmer]

Intel version probably runs slower than would be ATi or NV builds. Cause code paths are different for all these 3 builds.
Also, intel uses most precise but slowest trigonometry. Other vendors have enough precision in native trig functions to allow simpler and faster math.

EDIT: in other words, if you limited in time you could build all types with same "intel" path (take care for platform selection code though!) but if you can better to use optimized path for each of vendors.

//"intel" path currently almost identical to ATi non-HD5 path. And it's not right cause iGPU has real local memory. I'm in process of changing that so soon USE_OPENCL_INTEL could give enother results for your type of testing.

[Joe Fox]

Ah, thanks for the clarification! That explains what I'm observing.

[Joe Fox]

I'm not sure if this information will be useful or not, but here is the debug message I entered along with the code that was causing the error:

Code: [Select]

        fprintf(stderr,"INFO: oclFFT done no strip. plan_ptr: %p, FtfNum=%d, fftlen=%d, NumBlockFfts=%d, chirplen=%d\n", plan, FftNum, fftlen, NumBlockFfts, chirplen);
        err |= clFFT_ExecuteInterleaved_mb(cq, plan[FftNum],NumBlockFfts, clFFT_Inverse, gpu_ChirpedData,gpu_WorkData, FFTbuf, 0, NULL, NULL);

And the output of that line is:
INFO: oclFFT done no strip. plan_ptr: 0x7fff4fecf060, FtfNum=0, fftlen=8, NumBlockFfts=131072, chirplen=1048576

Error in mb oclFFT_2: -49
ERROR: OpenCL kernel/call 'non-strip fft' call failed (-49) in file /Users/joe/projects/seti/sah_v7_opt/AKv8/client/analyzeFuncs.cpp near line 3823.


This is running the USE_OPENCL_HD5XXX define on an AMD Mac Pro.

I'm not sure what exactly I should be looking for, but plan looks to be non-null, and the index (FftNum) is 0, so I would expect it to be in bounds. Any insight?

[Raistmer]

It's too high-level call again.
As I proposed earlier in this thread try to go inside that oclFFT API call, to the place where params of particular _kernel_ call are set. And add debug code there.

I'll post particular place when will be at my dev PC closer to evening today.

[Raistmer]

Code: [Select]

      //fprintf(stderr,"After getKernelWorkDimensions:\nbatchSize s =%d, gWorkItems =%d, lWorkItems =%d, dir =%d\n", s, gWorkItems, lWorkItems, dir);
err |= clSetKernelArg(kernelInfo->kernel, 0, sizeof(cl_mem), &memObj[currRead]);
err |= clSetKernelArg(kernelInfo->kernel, 1, sizeof(cl_mem), &memObj[currWrite]);
err |= clSetKernelArg(kernelInfo->kernel, 2, sizeof(cl_int), &dir);
err |= clSetKernelArg(kernelInfo->kernel, 3, sizeof(cl_int), &s);
#if USE_OPENCL_INTEL //R: only iGPU uses Taylor trig approx and LuT. Old Ati drivers can't accept NULL as valid buf
err |= clSetKernelArg(kernelInfo->kernel, 4, sizeof(cl_mem), &(plan->cossin_LUT_d1));
            err |= clSetKernelArg(kernelInfo->kernel, 5, sizeof(cl_mem), &(plan->cossin_LUT_d2));
#endif
err |= clEnqueueNDRangeKernel(queue,  kernelInfo->kernel, 1, NULL, &gWorkItems, &lWorkItems, 0, NULL,NULL/*&fft_event*/);
if(err){
fprintf(stderr,"Error in mb oclFFT_1: %d\n",err);
return err;
}
#if OCL_VERBOSE
fprintf(stderr,"INFO: in mb oclFFT_1 ok\n");
#endif
and

Code: [Select]

        //fprintf(stderr,"After getKernelWorkDimensions:\nbatchSize s =%d, gWorkItems =%d, lWorkItems =%d, dir =%d\n", s, gWorkItems, lWorkItems, dir);
    err |= clSetKernelArg(kernelInfo->kernel, 0, sizeof(cl_mem), &memObj[currRead]);
    err |= clSetKernelArg(kernelInfo->kernel, 1, sizeof(cl_mem), &memObj[currWrite]);
    err |= clSetKernelArg(kernelInfo->kernel, 2, sizeof(cl_int), &dir);
    err |= clSetKernelArg(kernelInfo->kernel, 3, sizeof(cl_int), &s);
#if USE_OPENCL_INTEL //R: only iGPU uses Taylor trig approx and LuT. Old Ati drivers can't accept NULL as valid buf
err |= clSetKernelArg(kernelInfo->kernel, 4, sizeof(cl_mem), &(plan->cossin_LUT_d1));
            err |= clSetKernelArg(kernelInfo->kernel, 5, sizeof(cl_mem), &(plan->cossin_LUT_d2));
#endif
    err |= clEnqueueNDRangeKernel(queue,  kernelInfo->kernel, 1, NULL, &gWorkItems, &lWorkItems, 0, NULL,NULL/*&fft_event*/);
if(err){
fprintf(stderr,"Error in mb oclFFT_2: %d\n",err);
return err;
}
#if OCL_VERBOSE
fprintf(stderr,"INFO: in mb oclFFT_2 ok\n");
#endif


[Joe Fox]

ah, ok. Sorry for the misunderstanding.

[Joe Fox]

Finally got some time to trace down to the KernelArg that is throwing the error.

Here is the relevant code:
fft_execute.cpp line 583-603

Code: [Select]

    getKernelWorkDimensions(plan, kernelInfo, &s, &gWorkItems, &lWorkItems);
        fprintf(stderr,"After getKernelWorkDimensions:\nbatchSize s =%d, gWorkItems =%lu, lWorkItems =%lu, dir =%d\n", s, gWorkItems, lWorkItems, dir);
    err |= clSetKernelArg(kernelInfo->kernel, 0, sizeof(cl_mem), &memObj[currRead]);
            fprintf(stderr, "after param 0: err=%d\n", err);
    err |= clSetKernelArg(kernelInfo->kernel, 1, sizeof(cl_mem), &memObj[currWrite]);
            fprintf(stderr, "after param 1: err=%d\n", err);
    err |= clSetKernelArg(kernelInfo->kernel, 2, sizeof(cl_int), &dir);
            fprintf(stderr, "after param 2: err=%d\n", err);
    err |= clSetKernelArg(kernelInfo->kernel, 3, sizeof(cl_int), &s);
            fprintf(stderr, "after param 3: err=%d\n", err);
            err |= clSetKernelArg(kernelInfo->kernel, 4, sizeof(cl_mem), &(plan->cossin_LUT_d1));
            fprintf(stderr, "after param 4: err=%d\n", err);
            err |= clSetKernelArg(kernelInfo->kernel, 5, sizeof(cl_mem), &(plan->cossin_LUT_d2));
            fprintf(stderr, "after param 5: err=%d\n", err);

    err |= clEnqueueNDRangeKernel(queue,  kernelInfo->kernel, 1, NULL, &gWorkItems, &lWorkItems, 0, NULL,NULL/*&fft_event*/);
            fprintf(stderr, "after enqueue: err=%d\n", err);
if(err){
fprintf(stderr,"Error in mb oclFFT_2: %d\n",err);
return err;
}

and the resulting output:

Code: [Select]

INFO: oclFFT done no strip. plan_ptr: 0x7fff500f0060, FtfNum=0, fftlen=8, NumBlockFfts=131072, chirplen=1048576
After getKernelWorkDimensions:
batchSize s =131072, gWorkItems =131072, lWorkItems =64, dir =1
after param 0: err=0
after param 1: err=0
after param 2: err=0
after param 3: err=0
after param 4: err=-49
after param 5: err=-49
after enqueue: err=-49
Error in mb oclFFT_2: -49
ERROR: OpenCL kernel/call 'non-strip fft' call failed (-49) in file /Users/joe/projects/seti/sah_v7_opt/AKv8/client/analyzeFuncs.cpp near line 3823.

So it looks like line 593

Code: [Select]

            err |= clSetKernelArg(kernelInfo->kernel, 4, sizeof(cl_mem), &(plan->cossin_LUT_d1));
is the culprit.

This is all on svn revision 2760 which appears to be the latest.

And advice on next steps?

[Raistmer]

The advise is quite simple

that's how really current code looks:

          err |= clSetKernelArg(kernelInfo->kernel, 3, sizeof(cl_int), &s);
#if USE_OPENCL_INTEL //R: only iGPU uses Taylor trig approx and LuT. Old Ati drivers can't accept NULL as valid buf
         err |= clSetKernelArg(kernelInfo->kernel, 4, sizeof(cl_mem), &(plan->cossin_LUT_d1));
            err |= clSetKernelArg(kernelInfo->kernel, 5, sizeof(cl_mem), &(plan->cossin_LUT_d2));
#endif
          err |= clEnqueueNDRangeKernel(queue,  kernelInfo->kernel, 1, NULL, &gWorkItems, &lWorkItems, 0, NULL,NULL/*&fft_event*/);

So, do update to the latest revision and don't forget this time that SETI_opt source tree includes not only AKv8 directory but also src one that common between both MB and AP. Seems you use outdated src now.

[Joe Fox]

Ok, I see now. There are two copies of fft_execute. The shared one in src/OpenCL_FFT and another one in AKv8/mac_build/OpenCL_FFT/ which is out of date. I will update my project to point to the shared one. Sorry for the trouble and thanks for the help.

[Joe Fox]

And to clarify, the duplication was due to a mistake that I made. I have cleaned up my project and removed the duplicate copy of OpenCL_FFT, and now I am getting a successful run using the ati version on the mac pro. I will also check the nvidia version later today.

Thanks again for the assistance.

[Raistmer]

Great! Looking for OSX supporting OpenCL MB fully 

[Joe Fox]

All three builds are working correctly based on my tests. I have submitted the binaries to Eric Korpela. If you are curious, and want to play with the builds before he posts them, you can find them here: https://www.dropbox.com/s/hlh9m3z1cl6nlbl/seti%40home.zip?dl=0

[Raistmer]

Fine!
I have no OS X systems, maybe Urs would try them.
Now we only have to wait when SETI servers will be fixed ultimately...