[Inference] Add cutlass gemm dequant op

.gitignore

@@ -125,5 +125,7 @@ FETCH_HEAD
 .vscode
 ./ppdiffusers/ppdiffusers/version.py
 
+# third party
+csrc/gpu/cutlass_kernels/cutlass
 dataset/
-output/
+output/

csrc/README.md

@@ -13,3 +13,16 @@ pip install -r requirements.txt
 ```shell
 python setup_cuda.py install
 ```
+
+### 手动安装 Cutlass 库
+1. 访问 Cutlass 仓库: [NVIDIA/cutlass](https://github.com/NVIDIA/cutlass)
+
+2. 拉取代码:
+    git clone -b v3.5.0 --single-branch https://github.com/NVIDIA/cutlass.git
+
+3. 将下载的 `cutlass` 目录放在 `csrc/gpu/cutlass_kernels/cutlass`下
+
+4. 重新编译 Cuda 算子
+```shell
+python setup_cuda.py install
+```

csrc/gpu/cutlass_kernels/cutlass_helper.h

@@ -0,0 +1,42 @@
+// Copyright (c) 2023 PaddlePaddle Authors. All Rights Reserved.
+// 
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+// 
+//     http://www.apache.org/licenses/LICENSE-2.0
+// 
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+#pragma once
+
+#include "cutlass/half.h"
+#include "cutlass/bfloat16.h"
+#include "paddle/extension.h"
+
+template <paddle::DataType D>
+class CutlassDtypeTraits;
+
+template <>
+class CutlassDtypeTraits<paddle::DataType::FLOAT32> {
+public:
+  typedef float DataType;
+  typedef float data_t;
+};
+
+template <>
+class CutlassDtypeTraits<paddle::DataType::FLOAT16> {
+public:
+  typedef cutlass::half_t DataType;
+  typedef paddle::float16 data_t;
+};
+
+template <>
+class CutlassDtypeTraits<paddle::DataType::BFLOAT16> {
+public:
+  typedef cutlass::bfloat16_t DataType;
+  typedef paddle::bfloat16 data_t;
+};