Lbgpu node vel #2878

fweik · 2019-05-31T18:51:29Z

Description of changes:

Factored out velocity getter, and obey boundary velocity for
both interpolation schemes,
Removed some globals
Leak less

codecov · 2019-05-31T19:12:17Z

Codecov Report

Merging #2878 into python will decrease coverage by <1%.
The diff coverage is n/a.

@@          Coverage Diff           @@
##           python   #2878   +/-   ##
======================================
- Coverage      82%     82%   -1%     
======================================
  Files         525     525           
  Lines       26807   26807           
======================================
- Hits        22015   22013    -2     
- Misses       4792    4794    +2

Impacted Files	Coverage Δ
src/core/electrostatics_magnetostatics/p3m.cpp	`85% <0%> (-1%)`	⬇️

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update 326c261...e416e7e. Read the comment docs.

fweik · 2019-06-03T08:40:16Z

@mkuron do you have seen the rocm error before? (..### HCC STATUS_CHECK Error: HSA_STATUS_ERROR_INVALID_ISA (0x100f) at file:mcwamp_hsa.cpp line:1193, looks like a toolchain issue to me)

mkuron · 2019-06-03T16:07:55Z

No idea. Could be some variation of an out-of-memory error or out-of-registers.

KaiSzuttor · 2019-06-06T09:08:28Z

/**
   * The instruction set architecture is invalid.
   */
  HSA_STATUS_ERROR_INVALID_ISA = 0x100F,

https://github.com/RadeonOpenCompute/ROCR-Runtime/blob/faa3ae51388517353afcdaf9c16621f879ef0a59/src/inc/hsa.h#L196-L199

KaiSzuttor · 2019-06-06T09:14:05Z

actually I don't understand why the compilation is not terminated, there are a number of compile errors:

/builds/espressomd/espresso/src/core/electrostatics_magnetostatics/p3m_gpu_error_cuda.cu:104:37: error:  'Utils::sinc':  no overloaded function has restriction specifiers that are compatible with the ambient context 'p3m_k_space_error_gpu_kernel_ik'
        int_pow<2 * cao>(Utils::sinc(meshi.x * nx) * Utils::sinc(meshi.y * ny) *
                                    ^
/builds/espressomd/espresso/src/core/electrostatics_magnetostatics/p3m_gpu_error_cuda.cu:104:37: error:  'Utils::sinc':  no overloaded function has restriction specifiers that are compatible with the ambient context 'p3m_k_space_error_gpu_kernel_ik'
/builds/espressomd/espresso/src/core/electrostatics_magnetostatics/p3m_gpu_error_cuda.cu:104:65: error:  'Utils::sinc':  no overloaded function has restriction specifiers that are compatible with the ambient context 'p3m_k_space_error_gpu_kernel_ik'
        int_pow<2 * cao>(Utils::sinc(meshi.x * nx) * Utils::sinc(meshi.y * ny) *
                                                                ^
/builds/espressomd/espresso/src/core/electrostatics_magnetostatics/p3m_gpu_error_cuda.cu:104:65: error:  'Utils::sinc':  no overloaded function has restriction specifiers that are compatible with the ambient context 'p3m_k_space_error_gpu_kernel_ik'
/builds/espressomd/espresso/src/core/electrostatics_magnetostatics/p3m_gpu_error_cuda.cu:105:37: error:  'Utils::sinc':  no overloaded function has restriction specifiers that are compatible with the ambient context 'p3m_k_space_error_gpu_kernel_ik'
                         Utils::sinc(meshi.z * nz));
                                    ^
/builds/espressomd/espresso/src/core/electrostatics_magnetostatics/p3m_gpu_error_cuda.cu:105:37: error:  'Utils::sinc':  no overloaded function has restriction specifiers that are compatible with the ambient context 'p3m_k_space_error_gpu_kernel_ik'
/builds/espressomd/espresso/src/core/electrostatics_magnetostatics/p3m_gpu_error_cuda.cu:151:39: error:  'Utils::sinc':  no overloaded function has restriction specifiers that are compatible with the ambient context 'p3m_k_space_error_gpu_kernel_ad'
          U2 = pow((double)Utils::sinc(meshi.x * nmx) *
                                      ^
/builds/espressomd/espresso/src/core/electrostatics_magnetostatics/p3m_gpu_error_cuda.cu:151:39: error:  'Utils::sinc':  no overloaded function has restriction specifiers that are compatible with the ambient context 'p3m_k_space_error_gpu_kernel_ad'
/builds/espressomd/espresso/src/core/electrostatics_magnetostatics/p3m_gpu_error_cuda.cu:152:35: error:  'Utils::sinc':  no overloaded function has restriction specifiers that are compatible with the ambient context 'p3m_k_space_error_gpu_kernel_ad'
                       Utils::sinc(meshi.y * nmy) * Utils::sinc(meshi.z * nmz),
                                  ^
/builds/espressomd/espresso/src/core/electrostatics_magnetostatics/p3m_gpu_error_cuda.cu:152:35: error:  'Utils::sinc':  no overloaded function has restriction specifiers that are compatible with the ambient context 'p3m_k_space_error_gpu_kernel_ad'
/builds/espressomd/espresso/src/core/electrostatics_magnetostatics/p3m_gpu_error_cuda.cu:152:64: error:  'Utils::sinc':  no overloaded function has restriction specifiers that are compatible with the ambient context 'p3m_k_space_error_gpu_kernel_ad'
                       Utils::sinc(meshi.y * nmy) * Utils::sinc(meshi.z * nmz),
                                                               ^
/builds/espressomd/espresso/src/core/electrostatics_magnetostatics/p3m_gpu_error_cuda.cu:152:64: error:  'Utils::sinc':  no overloaded function has restriction specifiers that are compatible with the ambient context 'p3m_k_space_error_gpu_kernel_ad'
/builds/espressomd/espresso/src/core/electrostatics_magnetostatics/p3m_gpu_error_cuda.cu:210:39: error:  'Utils::sinc':  no overloaded function has restriction specifiers that are compatible with the ambient context 'p3m_k_space_error_gpu_kernel_ik_i'
          U2 = pow((double)Utils::sinc(meshi.x * nmx) *
                                      ^
/builds/espressomd/espresso/src/core/electrostatics_magnetostatics/p3m_gpu_error_cuda.cu:210:39: error:  'Utils::sinc':  no overloaded function has restriction specifiers that are compatible with the ambient context 'p3m_k_space_error_gpu_kernel_ik_i'
/builds/espressomd/espresso/src/core/electrostatics_magnetostatics/p3m_gpu_error_cuda.cu:211:35: error:  'Utils::sinc':  no overloaded function has restriction specifiers that are compatible with the ambient context 'p3m_k_space_error_gpu_kernel_ik_i'
                       Utils::sinc(meshi.y * nmy) * Utils::sinc(meshi.z * nmz),
                                  ^
/builds/espressomd/espresso/src/core/electrostatics_magnetostatics/p3m_gpu_error_cuda.cu:211:35: error:  'Utils::sinc':  no overloaded function has restriction specifiers that are compatible with the ambient context 'p3m_k_space_error_gpu_kernel_ik_i'
/builds/espressomd/espresso/src/core/electrostatics_magnetostatics/p3m_gpu_error_cuda.cu:211:64: error:  'Utils::sinc':  no overloaded function has restriction specifiers that are compatible with the ambient context 'p3m_k_space_error_gpu_kernel_ik_i'
                       Utils::sinc(meshi.y * nmy) * Utils::sinc(meshi.z * nmz),
                                                               ^
/builds/espressomd/espresso/src/core/electrostatics_magnetostatics/p3m_gpu_error_cuda.cu:211:64: error:  'Utils::sinc':  no overloaded function has restriction specifiers that are compatible with the ambient context 'p3m_k_space_error_gpu_kernel_ik_i'
/builds/espressomd/espresso/src/core/electrostatics_magnetostatics/p3m_gpu_error_cuda.cu:273:39: error:  'Utils::sinc':  no overloaded function has restriction specifiers that are compatible with the ambient context 'p3m_k_space_error_gpu_kernel_ad_i'
          U2 = pow((double)Utils::sinc(meshi.x * nmx) *
                                      ^
fatal error: too many errors emitted, stopping now [-ferror-limit=]
20 errors generated.

mkuron · 2019-06-06T10:48:14Z

Good find. So this is actually something like an unresolved symbol error.

actually I don't understand why the compilation is not terminated

That‘s a known hipcc bug. It doesn‘t produce nonzero exit codes if compilation failed. Usually that‘s not a problem because the object file would subsequently be missing and CMake terminates anyway, but in the current case, an (incomplete?) object file seems to have already been written. That issue was fixed last month: ROCm/HIP#1117.

jngrad · 2019-06-11T16:36:31Z

Concerning the no overloaded function has restriction specifiers compiler error, the explanation can be found in ROCm/HIP#374 (comment). The error message disappears when using #if defined(__CUDACC__) or defined(__HIPCC__) instead of

espresso/src/utils/include/utils/device_qualifier.hpp

Line 4 in 7eebdb9

#if defined(__CUDACC__)

fweik · 2019-06-12T08:23:30Z

@KaiSzuttor there is a merge regression now that reintroduces the ek_params stuff?

KaiSzuttor · 2019-06-12T08:29:37Z

I'll try again...

fweik · 2019-06-12T08:38:06Z

src/core/grid_based_algorithms/lbgpu_cuda.cu

 */
 __global__ void integrate(LB_nodes_gpu n_a, LB_nodes_gpu n_b, LB_rho_v_gpu *d_v,
                          LB_node_force_density_gpu node_f,
+                          EK_parameters *ek_parameters_gpu,


this was intended by this PR, right?

No, those are not needed and were removed in the master after this branch was forked of

KaiSzuttor · 2019-06-12T08:39:59Z

was there before merge:

espresso/src/core/grid_based_algorithms/lbgpu_cuda.cu

Line 2157 in e016779

EK_parameters *ek_parameters_gpu,

fweik · 2019-06-12T08:40:52Z

Still this reverts #2877 which it should not.

KaiSzuttor · 2019-06-12T08:41:34Z

puh, then maybe you should merge...

…lbgpu_node_vel

fweik · 2019-06-12T08:55:45Z

@KaiSzuttor

fweik · 2019-07-09T11:06:55Z

This is too difficult to debug for me, giving up.

mkuron · 2019-07-09T12:00:48Z

src/core/grid_based_algorithms/lbgpu_cuda.cu

+  float *boundary_velocity = nullptr;
+  int *boundary_node_list = nullptr;
+  int *boundary_index_list = nullptr;
+  size_t size_of_boundindex = 0;


These variables shadow globals, and these globals are not even used anymore. It's possible that that was the cause of the HSA_STATUS_ERROR_INVALID_ISA you were seeing.

HSA_STATUS_ERROR_INVALID_ISA: The instruction set architecture is invalid.

That‘s just a generic error code. It is used in many places for which no specific error code has been defined.

KaiSzuttor · 2019-07-09T12:01:15Z

@fweik it's only the rocm test that fails... maybe we should not throw away this PR just because of that

fweik · 2019-07-09T12:08:41Z

Well you know where to find the code. I'm done with this.

mkuron · 2019-07-09T13:08:44Z

This is where the error comes from: https://github.com/RadeonOpenCompute/ROCR-Runtime/blob/e2d1950bd8fc83dd5b8352a41b829a7c57fc1073/src/core/runtime/amd_aql_queue.cpp#L820. It supposedly means "Out of VGPRs".

On CUDA, an out-of-registers condition would cause an overflow into global memory -- slowing things down, but still working. I guess we don't have infrastructure in place that would detect this kind of performance regression on CUDA though. So it's plausible that a function requiring a very large number of registers would be broken on AMD. Adding -Xptxas=-vto the CUDA flags reveals a few functions that are close to the limit (63 on sm_30 and 255 on later ones) and have probably already overflowed into global memory.

I think AMD can deal with 255 VGPRs too, but it's unclear to me how to measure how many of them a kernel is using. Maybe this pull request just crosses the limit. To reduce register consumption, large kernels would need to be broken up into multiple smaller ones.

mkuron · 2019-07-09T15:55:59Z

It's the call to thrust::transform(..., lb_get_interpolated_velocity_gpu<27>(...)). thrust::transform(..., lb_get_interpolated_velocity_gpu<8>(...)) works fine. #2982 fixes the issue.

@fweik

2982: Reduce excessive loop unrolling in lbgpu velocity interpolation r=KaiSzuttor a=mkuron This caused excessive register usage, especially when combined with thrust. Issue discovered by @fweik in #2878. It turns out that this is a problem for CUDA too, it just exhibits a different behavior. Instead of crashing like on HIP, CUDA just produces a large binary and slower code. In a perfect world, the compiler should display a warning, but I guess neither AMD nor Nvidia operate in a perfect world. Co-authored-by: Michael Kuron <mkuron@users.noreply.github.com>

jngrad · 2019-07-10T13:26:06Z

@KaiSzuttor the webhook didn't trigger CI for 81b9c33

KaiSzuttor · 2019-07-11T09:57:37Z

bors r+

2878: Lbgpu node vel r=KaiSzuttor a=fweik Description of changes: - Factored out velocity getter, and obey boundary velocity for both interpolation schemes, - Removed some globals - Leak less Co-authored-by: Florian Weik <fweik@icp.uni-stuttgart.de> Co-authored-by: RudolfWeeber <weeber@icp.uni-stuttgart.de> Co-authored-by: Kai Szuttor <2150555+kaiszuttor@users.noreply.github.com> Co-authored-by: Kai Szuttor <kai@icp.uni-stuttgart.de>

bors · 2019-07-11T10:55:35Z

Build succeeded

ICP GitLab CI

fweik added 2 commits May 31, 2019 19:56

core: lbgpu: Factor out velocity getter

e9e21c1

core: lbgpu: Removed lb_boundary_velocity global and leak less

2c08ed8

fweik requested a review from KaiSzuttor May 31, 2019 18:51

fweik added 2 commits May 31, 2019 20:55

Formatting

91e2bd6

Doxygen

e016779

fweik changed the title ~~Lbgpu node vel~~ WIP: Lbgpu node vel Jun 12, 2019

fweik commented Jun 12, 2019

View reviewed changes

Merge branch 'python' of https://github.com/espressomd/espresso into …

e97e4d7

…lbgpu_node_vel

fweik changed the title ~~WIP: Lbgpu node vel~~ Lbgpu node vel Jun 12, 2019

RudolfWeeber and others added 3 commits June 15, 2019 17:29

Merge branch 'python' into lbgpu_node_vel

b6d1262

Merge branch 'python' into lbgpu_node_vel

1256369

Merge branch 'python' into lbgpu_node_vel

6c5ace0

KaiSzuttor mentioned this pull request Jul 5, 2019

HIP issue list as discussed in the offline meeting #2973

Closed

fweik closed this Jul 9, 2019

mkuron reviewed Jul 9, 2019

View reviewed changes

mkuron mentioned this pull request Jul 9, 2019

Reduce excessive loop unrolling in lbgpu velocity interpolation #2982

Merged

KaiSzuttor reopened this Jul 10, 2019

Merge branch 'python' into lbgpu_node_vel

81b9c33

core: LB: remove unused globals.

e416e7e

KaiSzuttor approved these changes Jul 11, 2019

View reviewed changes

mkuron approved these changes Jul 11, 2019

View reviewed changes

bors bot merged commit e416e7e into espressomd:python Jul 11, 2019

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Lbgpu node vel #2878

Lbgpu node vel #2878

fweik commented May 31, 2019

codecov bot commented May 31, 2019 •

edited

Loading

fweik commented Jun 3, 2019

mkuron commented Jun 3, 2019

KaiSzuttor commented Jun 6, 2019

KaiSzuttor commented Jun 6, 2019

mkuron commented Jun 6, 2019 •

edited

Loading

jngrad commented Jun 11, 2019

fweik commented Jun 12, 2019

KaiSzuttor commented Jun 12, 2019

fweik Jun 12, 2019

KaiSzuttor Jun 12, 2019

fweik Jun 12, 2019

KaiSzuttor commented Jun 12, 2019 •

edited

Loading

fweik commented Jun 12, 2019

KaiSzuttor commented Jun 12, 2019

fweik commented Jun 12, 2019

fweik commented Jul 9, 2019

mkuron Jul 9, 2019

jngrad Jul 9, 2019

mkuron Jul 9, 2019

KaiSzuttor commented Jul 9, 2019

fweik commented Jul 9, 2019

mkuron commented Jul 9, 2019

mkuron commented Jul 9, 2019 •

edited

Loading

jngrad commented Jul 10, 2019

KaiSzuttor commented Jul 11, 2019

bors bot commented Jul 11, 2019

Lbgpu node vel #2878

Lbgpu node vel #2878

Conversation

fweik commented May 31, 2019

codecov bot commented May 31, 2019 • edited Loading

Codecov Report

fweik commented Jun 3, 2019

mkuron commented Jun 3, 2019

KaiSzuttor commented Jun 6, 2019

KaiSzuttor commented Jun 6, 2019

mkuron commented Jun 6, 2019 • edited Loading

jngrad commented Jun 11, 2019

fweik commented Jun 12, 2019

KaiSzuttor commented Jun 12, 2019

fweik Jun 12, 2019

Choose a reason for hiding this comment

KaiSzuttor Jun 12, 2019

Choose a reason for hiding this comment

fweik Jun 12, 2019

Choose a reason for hiding this comment

KaiSzuttor commented Jun 12, 2019 • edited Loading

fweik commented Jun 12, 2019

KaiSzuttor commented Jun 12, 2019

fweik commented Jun 12, 2019

fweik commented Jul 9, 2019

mkuron Jul 9, 2019

Choose a reason for hiding this comment

jngrad Jul 9, 2019

Choose a reason for hiding this comment

mkuron Jul 9, 2019

Choose a reason for hiding this comment

KaiSzuttor commented Jul 9, 2019

fweik commented Jul 9, 2019

mkuron commented Jul 9, 2019

mkuron commented Jul 9, 2019 • edited Loading

jngrad commented Jul 10, 2019

KaiSzuttor commented Jul 11, 2019

bors bot commented Jul 11, 2019

Build succeeded

codecov bot commented May 31, 2019 •

edited

Loading

mkuron commented Jun 6, 2019 •

edited

Loading

KaiSzuttor commented Jun 12, 2019 •

edited

Loading

mkuron commented Jul 9, 2019 •

edited

Loading