Fix gemm_algorithm_picker bug for cublasLt.
Hi @cheshire Can you please review this PR ? Thank you! Makes sense ! I think we do not notice it because we do not use cublasLT by default. @SandSnip3r could you take a look? Hmm, yeah. As far as I can tell, it looks like we never set the algorithm for cublasLt. However, I disagree with the way that this PR goes about setting it.

The code inside of `GetBestAlgorithm` already sets the algorithm via
```
result.mutable_gemm()->set_algorithm(profile_result.algorithm());
```

Instead, we should be setting the algorithm inside of the profile result similarly to how it's done for cublas. For cublas, the `blas::ProfileResult`'s algorithm is set inside of `cuda_blas.cc`. The same should be done inside `cuda_blas_lt_.cc` for cublasLt. This PR would further complicate the already complicated "blas" abstraction. > Hmm, yeah. As far as I can tell, it looks like we never set the algorithm for cublasLt. However, I disagree with the way that this PR goes about setting it.
> 
> The code inside of `GetBestAlgorithm` already sets the algorithm via
> 
> ```
> result.mutable_gemm()->set_algorithm(profile_result.algorithm());
> ```
> 
> Instead, we should be setting the algorithm inside of the profile result similarly to how it's done for cublas. For cublas, the `blas::ProfileResult`'s algorithm is set inside of `cuda_blas.cc`. The same should be done inside `cuda_blas_lt_.cc` for cublasLt. This PR would further complicate the already complicated "blas" abstraction.

Thanks for the review @cheshire @SandSnip3r 

The algorithm in cublasLt is a little different from cublas. 

In cublas, we can specify the best algorithm by set a integer value (algorithm id).
However, in cublasLt, the algorithm is repsented by cublasLtMatmulAlgo_t, which is more complex[ref].(https://gitlab.com/nvidia/headers/cuda-individual/cublas/-/blob/main/cublasLt.h#L119)

In order to use the "blas" abstraction for gemm algorithm,  I think we can use cublasLtMatmulAlgoGetHeuristic in gemm_algorithm_picker, get the best algorithm, save it's index in blas::ProfileResult. And later in RunGemm(), we can first call cublasLtMatmulAlgoGetHeuristic(), and then use the saved index find the best algorithm.  > In cublas, we can specify the best algorithm by set a integer value (algorithm id). However, in cublasLt, the algorithm is repsented by cublasLtMatmulAlgo_t, which is more complex[ref].(https://gitlab.com/nvidia/headers/cuda-individual/cublas/-/blob/main/cublasLt.h#L119)
> 
> In order to use the "blas" abstraction for gemm algorithm, I think we can use cublasLtMatmulAlgoGetHeuristic in gemm_algorithm_picker, get the best algorithm, save it's index in blas::ProfileResult. And later in RunGemm(), we can first call cublasLtMatmulAlgoGetHeuristic(), and then use the saved index find the best algorithm.

I suppose you could do that, just to avoid needing to make the algorithm a bit more abstract in the profile result.

However, you now have an additional call to GetHeuristic. Also is that going to be called each time you call RunGemm? Also, that's trusting that the algorithms are returned in the same order every time, which could lead to a very hard-to-spot performance regression if that were no longer the case.

To do it right, I think you ought to adapt the interface to allow for the different types used by cublas and cublasLt for their algorithms. BTW also note that past Ampere, the select algorithm is ignored entirely in cuBLAS from what I recall. So ironically storing or not storing the algorithm on Ampere+ GPUs probably makes no difference. Hi @kimbaol Any update on this PR? Please. Thank you! This PR is stale because it has been open for 14 days with no activity. It will be closed if no further activity occurs. Thank you. Hi @kimbaol Any update on this PR? Please. Thank you! Hi @kimbaol I'm going to go ahead and close this PR, because it seems to have stalled. If you're still interested in pursing this (and responding to my comments), please feel free to reopen!  Thank you for you contribution!
In [cublas_lt_matmul.cc](https://github.com/tensorflow/tensorflow/blob/305c2c09dd998397f52da79650b36c35ff211225/tensorflow/compiler/xla/service/gpu/runtime/cublas_lt_matmul.cc#L140)，we will get all valid algorithms from cublasLtMatmulAlgoGetHeuristic()，and then select the i-th from  parameter "algorithm". However, it seems this "algorithm" parameter is not correctly set in gemm_algorithm_picker. 

This pr fix the bug, and I run several tests on A800, it can help to improve the gemm performance. 

(A800 with cuda12.2)
(XLA_FLAGS='--xla_gpu_enable_cublaslt', TF_CPP_VMODULE=gemm_algorithm_picker=10)
(m=1024, k=4096,n=2048)

Before

- logs
<img width="917" alt="image" src="https://github.com/tensorflow/tensorflow/assets/5410381/f8808d34-efb5-4c62-9699-3b75b44f1a37">

- result
<img width="1537" alt="image" src="https://github.com/tensorflow/tensorflow/assets/5410381/adc09542-9a27-445b-8e09-e40a0c2d25d4">




After 

- logs
<img width="914" alt="image" src="https://github.com/tensorflow/tensorflow/assets/5410381/433d1fbc-c816-44e3-b9fb-c84a0bcb7aa7">

- result
<img width="1526" alt="image" src="https://github.com/tensorflow/tensorflow/assets/5410381/c48acb15-f06d-44ab-8dfc-9d4866f92819">
