CUDALucas (a.k.a. MaclucasFFTW/CUDA 2.3/CUFFTW)

[msft]

Hi, frmky

Originally Posted by frmky

The non-TESRA version works fine, but the TESRA version seems numerically unstable. One compile would work on some GPUs but not others. I'd simply recompile and that would shuffle the GPUs it would work on. When it failed, it did so because the err returned was larger than kErrLimit, so it would keep increasing n until the grid size (n/512) of the rftfsub_kernel call was too large.

GTX2xx's max block number is 65536(Sorry this number from Japanease book), 4194304/512 < 65536.
But, I change next version.

Originally Posted by frmky

I modified the source so that when TESRA is defined, it uses the normalize_kernel from version k rather than breaking it into the three sequential calls. This works correctly on all of the GPUs in the S1070. I'm not sure why breaking it into three kernel calls was causing a problem. Attached is the source as modified.

Version "k" on My GTX260 return correct result, all time, absolutely.

But, a book "CUDA Programming Guide" say.

4.4.2 Synchronization Function
void __syncthreads();
synchronizes all threads in a block. Once all threads have reached this point,
execution resumes normally.

threads in a block is only block not all threads.

D.2.1 cudaThreadSynchronize()
cudaError_t cudaThreadSynchronize(void);
blocks until the device has completed all preceding requested tasks.
cudaThreadSynchronize() returns an error if one of the preceding tasks failed.

What on your mind ?

[frmky]

Originally Posted by msft

Hi, frmky

GTX2xx's max block number is 65536(Sorry this number from Japanease book), 4194304/512 < 65536.
But, I change next version.

That wasn't the real problem. It's just where it failed.

Originally Posted by msft

What on your mind ?

I really don't know. cudaThreadSynchronize() is the right thing to do there. And of course the non-TESRA version works perfectly with the three part normalization, with a couple transposes put in. But at least it's easy to avoid. I've started four double-check runs around 23 million with the modified version v I posted earlier. We'll have the results in a few days.

[msft]

Originally Posted by frmky

I really don't know. cudaThreadSynchronize() is the right thing to do there. And of course the non-TESRA version works perfectly with the three part normalization, with a couple transposes put in. But at least it's easy to avoid. I've started four double-check runs around 23 million with the modified version v I posted earlier. We'll have the results in a few days.

No problem, All of your decision.
Thank you,

[msft]

I test frmky's version source code with -DTESTRA and not -DTESTRA,
M131101 to M1548619 1000 iterations check sum compare to Glucas,it is correct.
Thank you,frmky

"WD1" is a "Public key cryptography" somethig?
It is important?
I use Glucas and Mlucas's Res64 result.

[msft]

Hi,

New version MaclucasFFTW.

checkpoint bug fix(MaclucasFFTW original bug).

Thank you,

[frmky]

Three of the four double-check runs matched the previous residue. This gives me reasonable confidence that the fourth is correct as well:

M23102129
M23102213
M23102267
M23102341

I'll switch to version y.

[msft]

Hi, frmky

I run one double-check runs around 23 million.

Thank you,

[Cruelty]

OT: can't wait to see some FFT implementation using OpenCL or DirectCompute 11, finally I would utilize my HD5870
BTW: AMD plans to expand developer tools for both

[msft]

Hi, Cruelty

Originally Posted by Cruelty

OT: can't wait to see some FFT implementation using OpenCL or DirectCompute 11, finally I would utilize my HD5870
BTW: AMD plans to expand developer tools for both

I hope they release ACML-GPU-FFT.

[msft]

Originally Posted by msft

I run one double-check runs around 23 million.

My GTX260 result.

M23103809

Primenet supoort MacLucasFFTW result, thank you

[nucleon]

How long are the 20M exponent double checks taking to complete?

-- Craig