search

ccecka / fmmtl

FMM Template Library
45 stars 14 forks source link

Gpu implementation has high relative error for provided examples #4

Open antequ opened 3 years ago

antequ commented 3 years ago

Hi, I think the GPU implementation has computation correctness errors.

I built everything with CUDA enabled, and the results I get from running the examples in the examples directory have high relative error, meaning the output is unusable. To compare, I also built everything for the cpu, with CUDA disabled, and their results have relative errors close to 0. Can you fix the CUDA computation errors?

Examples: Here are the CUDA-enabled GPU versions:

fmmtl/examples$ ./error_laplace
FMM in 0.216471 secs
FMM in 0.0459567 secs
FMM in 0.0447383 secs
Computing direct matvec...
Direct in 0.922225 secs
Vector  relative error: 0.623028
Average relative error: 0.444972
Maximum relative error: 1.20022
fmmtl/examples$ ./error_biot
FMM in 0.252116 secs
FMM in 0.0779735 secs
FMM in 0.0770103 secs
Computing direct matvec...
Direct in 1.33767 secs
Vector  relative error: 0.684924
Average relative error: 0.489844
Maximum relative error: 7.83254
fmmtl/examples$ ./error_barycentric
FMM in 0.152405 secs
FMM in 0.00364061 secs
FMM in 0.00367243 secs
Computing direct matvec...
Direct in 0.163656 secs
Vector  relative error: 0.999659
Average relative error: 3.72861
Maximum relative error: 13593
fmmtl/examples$ ./error_img
Initializing source and N = 1048576 targets...
Building the kernel matrix...
Performing the kernel matrix-vector mult...
Computing direct kernel matrix-vector mult...
Computing the errors...
Min log error: -16
Max log error: 0

On the other hand, kNN appears to be correct, implying there's just some kind of scale error:

fmmtl/examples$ ./kNN
Construct: 0.000110769
Traverse: 0.00758732
Computing direct...
Direct: 0.0059448
Wrong counts: 0 of 1000
((0.00202416, 907), (0.00437583, 298), (0.00447187, 507), (0.00450509, 178), (0.00470155, 833))

In contrast, here are the errors from the CPU versions (I used make clean && make -j34 error_biot error_laplace error_img NO_CUDA=1 to build this, after I added -fPIC to CXXFLAGS. I couldn't compile error_barycentric, but I don't need it.):

fmmtl/examples$ ./error_laplace
FMM in 0.0921008 secs
FMM in 0.0676527 secs
FMM in 0.067525 secs
Computing direct matvec...
Direct in 0.895511 secs
Vector  relative error: 3.36163e-05
Average relative error: 2.86249e-05
Maximum relative error: 0.000357396
fmmtl/examples$ ./error_biot
FMM in 0.135726 secs
FMM in 0.0971907 secs
FMM in 0.104078 secs
Computing direct matvec...
Direct in 1.3523 secs
Vector  relative error: 4.42274e-05
Average relative error: 5.30808e-05
Maximum relative error: 0.00247941
fmmtl/examples$ ./error_img
Initializing source and N = 1048576 targets...
Building the kernel matrix...
Performing the kernel matrix-vector mult...
Computing direct kernel matrix-vector mult...
Computing the errors...
Min log error: -16
Max log error: -1.90081

I'll include more about the unit tests in the next reply.

Here is some info about my OS and hardware: OS: Ubuntu 18.04 Compiled with Boost 1.65, g++-4.8, and nvcc V10.0.130. GPU: Nvidia Geforce GTX 1080, Driver 440.82, Cuda Driver 10.2 / Runtime 10.0, capability 6.1.

I also tried replacing all doubles with floats in the codebase (and changing some of the epsilon tolerances accordingly), and tested error_biot. The CPU version matched the correct results above, and the GPU version produced the same erroneous results.

If you're interested, here's the result from my deviceQuery:

Device 0: "GeForce GTX 1080"
  CUDA Driver Version / Runtime Version          10.2 / 10.0
  CUDA Capability Major/Minor version number:    6.1
  Total amount of global memory:                 8118 MBytes (8512602112 bytes)
  (20) Multiprocessors, (128) CUDA Cores/MP:     2560 CUDA Cores
  GPU Max Clock rate:                            1734 MHz (1.73 GHz)
  Memory Clock rate:                             5005 Mhz
  Memory Bus Width:                              256-bit
  L2 Cache Size:                                 2097152 bytes
  Maximum Texture Dimension Size (x,y,z)         1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384)
  Maximum Layered 1D Texture Size, (num) layers  1D=(32768), 2048 layers
  Maximum Layered 2D Texture Size, (num) layers  2D=(32768, 32768), 2048 layers
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       49152 bytes
  Total number of registers available per block: 65536
  Warp size:                                     32
  Maximum number of threads per multiprocessor:  2048
  Maximum number of threads per block:           1024
  Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
  Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
  Maximum memory pitch:                          2147483647 bytes
  Texture alignment:                             512 bytes
  Concurrent copy and kernel execution:          Yes with 2 copy engine(s)
  Run time limit on kernels:                     No
  Integrated GPU sharing Host Memory:            No
  Support host page-locked memory mapping:       Yes
  Alignment requirement for Surfaces:            Yes
  Device has ECC support:                        Disabled
  Device supports Unified Addressing (UVA):      Yes
  Device supports Compute Preemption:            Yes
  Supports Cooperative Kernel Launch:            Yes
  Supports MultiDevice Co-op Kernel Launch:      Yes
  Device PCI Domain ID / Bus ID / location ID:   0 / 79 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

Please look into this and correct the GPU implementation!

Thanks, Ante

antequ commented 3 years ago

The CPU and GPU unit tests have very similar results, with the exception of correctness and dual_correctness: the cpu version is wrong on 0 of 10000, while the GPU version is wrong on 10000 of 10000.

Unit tests (GPU version): compiles: success correctness: wrong on 10000 of 10000, most errors range from 0.1 to 0.4. dual_correctness: same results as correctness (wrong 10000 of 10000), with similar errors multi_level:

fmmtl/unit_tests$ ./multi_level
WARNING: Expansion does not have a correct M2T!
WARNING: Expansion does not have a correct M2T!
rexact = 0.589133 -0.200385 -0.200385 -0.200385
rm2t1 = 0 0 0 0
    [-0.589133 0.200385 0.200385 0.200385]
rm2t2 = 0 0 0 0
    [-0.589133 0.200385 0.200385 0.200385]
rfmm = 0.589147 -0.200343 -0.200343 -0.200343
    [1.44709e-05 4.28269e-05 4.28269e-05 4.28269e-05]

single_level:

fmmtl/unit_tests$ ./single_level
WARNING: Expansion does not have a correct M2T!
DIST: (0.8, 0.8, 0.8) : 1.38564
rexact = 0.589133 -0.200385 -0.200385 -0.200385
rm2t = 0 0 0 0
    [-0.589133 0.200385 0.200385 0.200385]
rfmm = 0.589056 -0.201209 -0.201209 -0.201209
    [-7.64378e-05 -0.00082331 -0.00082331 -0.00082331]

test_balltree: looks fine test_bbfmm:

fmmtl/unit_tests$ ./test_bbfmm
has_eval_op: 1
has_transpose: 1
has_vector_S2T_symm: 0
has_vector_S2T_asymm: 0
has_init_multipole: 1
has_init_local: 1
has_S2M: 1
  has_scalar_S2M: 0
  has_vector_S2M: 1
has_S2L: 0
  has_scalar_S2L: 0
  has_vector_S2L: 0
has_M2M: 1
has_M2L: 1
has_L2L: 1
has_M2T: 0
  has_scalar_M2T: 0
  has_vector_M2T: 0
has_L2T: 1
  has_scalar_L2T: 0
  has_vector_L2T: 1
has_dynamic_MAC: 0
FMM in 0.158716 secs
FMM in 0.0063701 secs
FMM in 0.0063367 secs
Computing direct matvec...
Direct in 0.226268 secs
Vector  relative error: 5.69574e-01
Average relative error: 3.63698e+00
Maximum relative error: 4.98108e+03

test_direct: fine test_expansion: Lots of issues. Pasting them for 16LaplaceSpherical, for example:

16LaplaceSpherical:
has_eval_op: 1
has_transpose: 1
has_vector_S2T_symm: 0
has_vector_S2T_asymm: 0
has_init_multipole: 1
has_init_local: 1
has_S2M: 1
  has_scalar_S2M: 1
  has_vector_S2M: 0
has_S2L: 0
  has_scalar_S2L: 0
  has_vector_S2L: 0
has_M2M: 1
has_M2L: 1
has_L2L: 1
has_M2T: 0
  has_scalar_M2T: 0
  has_vector_M2T: 0
has_L2T: 1
  has_scalar_L2T: 1
  has_vector_L2T: 0
has_dynamic_MAC: 0

test_gpu:

fmmtl/unit_tests$ ./test_gpu
terminate called after throwing an instance of 'thrust::system::system_error'
  what():  parallel_for failed: no kernel image is available for execution on the device
Aborted (core dumped)

test_kdtree: looks reasonable test_kernel: Just pasted a few:

10BiotSavart:
0 0 0
has_eval_op: 1
has_transpose: 1
has_vector_S2T_symm: 0
has_vector_S2T_asymm: 0

14RosenheadMoore:
0 0 0
has_eval_op: 1
has_transpose: 1
has_vector_S2T_symm: 0
has_vector_S2T_asymm: 0

test_ndtree: looks reasonable test_s2t:

fmmtl/unit_tests$ ./test_s2t
CPU-GPU:
  Vector  relative error: 1
  Average relative error: 1
  Maximum relative error: 1

CPU-GPU Blocked:
  Vector  relative error: 1
  Average relative error: 1
  Maximum relative error: 1

test_vec:

fmmtl/unit_tests$ ./test_vec
Is POD: 0
Is trivial: 0
Is standard layout: 1
0 0 0
0 8 20
21.5407
8.24
4.38972
3.14
1 2.1 3.14 2

version:

fmmtl/unit_tests$ ./version
Using Thrust v1.9