Perf: Optimize hsolver GPU code (useful information of GPU optimization: __syncwarp() should be used instead of __syncthreads())

[OldDriver233]

What's changed?

• Unrolling loops. As code in a wrap uses SIMD, no __syncthreads() is needed.
• Adjusting vector operations. Currently the operations used too less blocks(typically < 20), causing waste of SMs.

[OldDriver233]

@dyzheng the issues mentioned above are solved in commit 40f973f.

[dyzheng]

The remain question is that can you list some test result to verify your code can improve the performance in PR description?

[OldDriver233]

The remain question is that can you list some test result to verify your code can improve the performance in PR description?

Yes. For vector functions the change had been obvious.

• Before the change:
  

• After the change:
  

[OldDriver233]

For line_minimize and calc_grad the difference is minor.

• Before:
  

• After:
  

[OldDriver233]

N.B. the testing case is ./examples/gpu/si16_pw, running on NVIDIA GeForce RTX 3070 Mobile / Max-Q.

[caic99]

@OldDriver233 What tools are you using for profiling & visualizing those results? Is that nsight compute? That looks really cool.

[OldDriver233]

@OldDriver233 What tools are you using for profiling & visualizing those results? Is that nsight compute? That looks really cool.

Yes. In nsight compute, you can filter out kernel functions you wish to profile, which is a useful feature.

[haozhihan]

LGTM

[haozhihan]

By the way, because we are working on first-principles calculation software, the cases of calculations are often very different, including large cases and small cases.

When we talk about performance improvement, if we only focus on one case (or even one kernel), it is not very convincing.

When @denghuilu and I (@haozhihan) were developing GPU code before, we needed to conduct system testing (many cases covering various systems) to convince @mohanchen that our code performance had greatly improved.

Considering this, should we specify a process to prove performance improvement so that external developers can better participate in the performance improvement of ABACUS?

There really was no external developer contributing code for performance improvements in ABACUS before.

So thanks again to @OldDriver233.

[OldDriver233]

testcase.zip
Here is the testing result of 10 cases. Note that my computer(With NVIDIA GeForce RTX 3070 Mobile / Max-Q) reported out-of-memory on testcase 009, so this case is not tested.

For the 9 cases, vector related functions appeared to be faster than the original program in most cases. In case 007 the difference appears to be the biggest: 495008 vs. 523616, about 5% difference.

But for calc_grad and line_minimize, situation becomes different. For small cases they are slightly faster, but for bigger cases they became slower. I suppose that my card does not have enough fp64 units, resulting in the difference.

[mohanchen]

LGTM