Improve ell performance

[yhmtsai]

This PR improves the ell performance.

To read the memory contiguously, make the number of row in block to be the multiple of warp_size
(the warp_size is maybe not the minimal)
the block size is set to 512
In CUDA, the max number per row in one thread block is 16 (3216)
In AMD, the max number per row in one thread block is 8 (648)

I do not do warp reduction. I use the shared memory to accumulate the result from different threads by atomicAdd in the same row because the threads in the same rows are not in the same warp.
I try non-square transpose and then do warp reduction, whose performance is similar with this PR in ND/nd12.

I test on 30 matrices. I sort the speedup of hipsp_ell/ell and take 10 in the head, 10 in the end and 10 in the middle.
i.e.

1. 10 matrices which hipsp_ell has largest speed up against original ell
2. 10 matrices which original ell has largest speed up against hipsp_ell
3. 10 matrices I pick up in the middle to check the implementation in general

• this pr ell vs original ell
  
• original speedup of ell vs hipsp_ell
  
• this pr speedup of ell vs hipsp_ell
  

Note: I still need to run this PR on the whole dataset

[upsj]

LGTM!

[upsj]

nit:

Suggested change

	static_cast<size_type>(blockDim.x) * blockIdx.x + threadIdx.x;
	size_type(blockDim.x) * blockIdx.x + threadIdx.x;

[thoasm]

I would actually prefer static_cast here since it is in fact a cast.
Usually, I would use size_type(12) with actual numbers (for me, this is using a constructor vs. casting).

Since I don't think we have any guidelines for it, feel free to use what you want @yhmtsai, since both options should result in the exact same behavior in this case.

[upsj]

I would agree that this is a matter of taste. My reasons for proposing all these changes are

1. they are semantically equivalent (except for composite type names like unsigned int, they can be used interchangably) and
2. the "constructor" solution is shorter.

I would however not agree that there is a difference between construction and conversion (i.e., casting) for primitive types. Maybe we can discuss this question in general and possibly find a common guideline before I start spamming all pull requests with these change suggestions ;)

[thoasm]

For primitive types, you are correct, there is no difference. However, for more complex types, it can make a difference.
I just like to have an explicit cast word, so everyone knows you are casting here.

I think adding either solution to the guidelines might be a bit too detailed since the guideline page is currently quite long. But we can discuss this tomorrow.

[tcojean]

For more complex types, I'm not so sure there is even a difference. See Section 5.2.3 and early 5.4:
http://www.open-std.org/jtc1/sc22/wg21/docs/papers/2011/n3242.pdf

If I understand all of this properly, this notation would even dynamic cast/reinterpret cast transparently? I never tested this but they link to section 5.4 which is for all casting operations.
We are going into the details of C++ here though.

[tcojean]

Also, a minor detail: in C++ standard jargon, they rather call this type of cast "functional notation" rather than "constructor notation".

[thoasm]

I have one question, otherwise, it looks good.

[tcojean]

LGTM!

[pratikvn]

LGTM!

[thoasm]

LGTM!

[yhmtsai]

For V100,
V100_ell_upd: this is the setting now. (the number of threads per row in one block is <= 32)
V100_ell_fix: the number of threads per row in one block is <= 16

x is the max number per row / #rows

• vs original ell
  
• V100_ell_32 vs cusparse
  
• V100_ell_fix vs cusparse
  

I choose 32 as the max number of threads per worker (the number of threads per worker is <= 32)

For radeon VII, I need to do full benchmark.
I try different limit for the maximum number of threads per worker in one block 8, 16, 32, 64.
There is a trade off among global memory coalescing, the number of work per threads, the atomic collision, and etc.
In the test, I choose the 32 as the maximum number of threads per row in one block which is same as cuda's.
Please ignore the value x is smaller than 1e-2, they should be the similar.
I set ratio 0 to see the performance with these parameters under 1e-2. the ratio 1e-2 still seems to be a good choice.

Also, I tried __builtin_nontemporal_load/__builtin_nontemporal_stored, but it looks like similar.

[yhmtsai]

It uses max_thread_per_worker is 32 in CUDA/AMD, specialize the kernel when num_thread_per_worker=1 without __builtin_nontemporal_load/__builtin_nontemporal_store.

x is max_nnz_per_row/nrows

radeon VII
I delete the max_nnz_per_row/nrows <= 1e-2 and max_nnz_per_row < 2 plot which uses num_thread_per_worker=1 because the kernel is wrong during benchmarking. Fixed

• ell_all_upd: max_thread_per_worker=32 and special 1
• ell_all_fix: max_thread_per_worker=8
• radeonVII: original

1. ell_all_upd, ell_all_fix, radeonVII ell
   
2. ell_all_upd, ell_all_fix
   
3. ell_all_upd, ell_all_fix, radeonVII hipsp_ell
   

V100
This version is correct

• V100_ell_upd: max_thread_per_worker=32 and special 1
• V100_ell_fix: max_thread_per_worker=16
• V100: original

1. V100_ell_upd, V100_ell_fix, V100 ell
   
2. V100_ell_upd, V100_ell_fix
   
3. V100_ell_upd, V100_ell_fix, V100 cusp_ell
   

I change some kernel staff and the variable name, so I re-request the reviews.

[thoasm]

LGTM!