Auto tune for cutlass #50809
Auto tune for cutlass #50809
Conversation
|
你的PR提交成功，感谢你对开源项目的贡献! |
| GatherGemmScatterGetCache() { | ||
| return autotune::AutoTuneCache::Instance().Get( | ||
| AlgorithmType::kGatherGemmScatterFP16NN); | ||
| } |
There was a problem hiding this comment.
GatherGemmScatterGetCache 函数移动到 cache_base.h目录下更合适.
There was a problem hiding this comment.
我将GatherGemmScatterGetCache重命名为GetGatherGemmScatter移动到cache.h，作为AutoTuneCache的成员函数了。
| const int32_t* b_indices, | ||
| const int32_t* c_d_indices, | ||
| T alpha, | ||
| T beta) { |
There was a problem hiding this comment.
既然类的模板参数中已经含有了变参模板 typename... Args，run 函数的书写可以仿照class MatmulAutoTuner 做简化，函数体内部也可以做同样地简化.
| static_cast<const int32_t*>(c_d_indices), | ||
| static_cast<cutlass::half_t>(1), | ||
| static_cast<cutlass::half_t>(1)); | ||
| GatherGemmScatter(dev_ctx, |
There was a problem hiding this comment.
static void dispatchKernel(const GPUContext& dev_ctx, 函数已经加入了模板参数 template <typename T>，const phi::DataType type应该不用传入了，配合模板参数if (type == phi::DataType::FLOAT16) { 这样的分支语句也可替换掉了.
There was a problem hiding this comment.
已简化，去掉了dispatchKernel，使用对GatherGemmScatterDriver（即原来的GatherGemmScatter)做partial template specialization的写法。
| nullptr, | ||
| static_cast<const int32_t*>(c_d_indices), | ||
| static_cast<float>(alpha), | ||
| static_cast<float>(beta)); |
There was a problem hiding this comment.
FLOAT64 的分支删除后，会采用什么措施补充吗？
There was a problem hiding this comment.
这个pr的做法是，fp64不走融合分支。因为模型不用fp64。
| for (auto i = 1; i < fp16_kernels.size(); i++) | ||
| tuner->AddCallBack(fp16_kernels[i]); | ||
|
|
||
| size_t key = autotune::GenKey(m / features_num_range, n, k); |
There was a problem hiding this comment.
建议把 T 模板参数转换为 phi::DataType 也传入到GenKey中，因为看到这个调用函数应该是同时针对 fp16和fp32，可能同样的输入M,N.K, 针对fp16和fp32两种类型，最佳计算Implement 也不相同的情况出现.
There was a problem hiding this comment.
cutlass提供的fp16和fp32的kernel互不通用，因此这里没有传入phi::DataType，而是用了partial template specialization，fp16和fp32分别在各自的候选kernel中搜索。
| #include "paddle/phi/backends/gpu/gpu_context.h" | ||
| #include "paddle/phi/common/data_type.h" | ||
| #include "paddle/phi/kernels/autotune/auto_tune_base.h" | ||
| #include "paddle/phi/kernels/sparse/gpu/cutlass_generator/build/generated/gemm/all_gemm_operations.h" |
There was a problem hiding this comment.
all_gemm_operations.h 这个头文件是否忘记在PR中提交上来？
There was a problem hiding this comment.
这个头文件是在编译期生成的，见PR50364
| void Run(const phi::GPUContext& ctx, | ||
| const size_t key, | ||
| T const alpha, | ||
| T const beta, |
There was a problem hiding this comment.
alpha, beta 两个参数的性能，看着也能被 Args... args 存储了.
There was a problem hiding this comment.
@JamesLim-sy
GatherGemmScatter会做一次matrix_c += alpha*matrix_a*matrix_b + beta*matrix_c。为避免在PickBestKernel时改变matrix_c的值，需要令alpha=0; beta=1;。因此此处将alpha, betaunpack。
PR types
New features
PR changes
OPs
Describe