[Build] Enable CUDA C++20 with host code isolation

.github/workflows/H-Coverage.yml

@@ -199,7 +199,7 @@ jobs:
           mkdir -p build && cd build
           ccache -z
           cmake .. -DPY_VERSION=3.12 -DWITH_GPU=ON -DWITH_DISTRIBUTE=ON -DWITH_TESTING=ON -DCUDA_ARCH_NAME=Manual -DCUDA_ARCH_BIN="90" -DFA_JOB_POOLS_COMPILE=1 -DWITH_CUDNN_FRONTEND=ON -DON_INFER=OFF -DWITH_NVSHMEM=ON
-          make -j20
+          CCACHE_DISABLE=1 make -k -j20
           EXIT_CODE=$?
           ccache -s
           exit $EXIT_CODE

cmake/cuda.cmake

@@ -323,7 +323,9 @@ message(STATUS "NVCC_FLAGS_EXTRA: ${NVCC_FLAGS_EXTRA}")
 set(CUDA_PROPAGATE_HOST_FLAGS OFF)
 # Release/Debug flags set by cmake. Such as -O3 -g -DNDEBUG etc.
 # So, don't set these flags here.
-set(CMAKE_CUDA_STANDARD 17)
+
+# Windows builds are forced to use C++17 in cmake/flags.cmake.
+set(CMAKE_CUDA_STANDARD 20)
 
 # (Note) For windows, if delete /W[1-4], /W1 will be added defaultly and conflict with -w
 # So replace /W[1-4] with /W0

paddle/phi/kernels/CMakeLists.txt

@@ -351,6 +351,10 @@ set(cc_search_pattern
     "stride/*.cc"
     "fusion/cpu/*.cc")
 
+if(WITH_GPU OR WITH_ROCM)
+  list(APPEND cc_search_pattern "fusion/gpu/*.cc")
+endif()
+
 if(WITH_ONEDNN)
   set(cc_search_pattern ${cc_search_pattern} "legacy/onednn/*.cc" "onednn/*.cc"
                         "fusion/onednn/*.cc")
@@ -365,6 +369,16 @@ file(
   RELATIVE "${CMAKE_CURRENT_SOURCE_DIR}"
   ${cc_search_pattern})
 
+if(WITH_ROCM)
+  list(
+    REMOVE_ITEM
+    kernel_cc
+    "fusion/gpu/fused_attention_kernel.cc"
+    "fusion/gpu/fused_attention_grad_kernel.cc"
+    "fusion/gpu/fused_feedforward_kernel.cc"
+    "fusion/gpu/fused_feedforward_grad_kernel.cc")
+endif()
+
 if(NOT
    (WITH_AVX
     AND AVX512F_FOUND

paddle/phi/kernels/funcs/fc_functor.cu

@@ -12,18 +12,11 @@ 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. */
 
-#include <algorithm>
-
-#include "paddle/phi/backends/all_context.h"
+// Keep the host functor interface out of NVCC's translation unit.
+#include "paddle/phi/backends/gpu/gpu_launch_config.h"  // NOLINT(build/include)
 #include "paddle/phi/kernels/funcs/aligned_vector.h"
-#include "paddle/phi/kernels/funcs/blas/blas.h"
-#include "paddle/phi/kernels/funcs/fc_functor.h"
-
-#include "paddle/phi/backends/gpu/gpu_launch_config.h"
-#include "paddle/phi/core/dense_tensor.h"
-#include "paddle/phi/kernels/funcs/blas/blaslt_impl.cu.h"
+#include "paddle/phi/kernels/funcs/fc_functor_gpu.h"
 #include "paddle/phi/kernels/funcs/quant_dequant.h"
-#include "paddle/phi/kernels/matmul_kernel.h"
 
 namespace phi {
 namespace funcs {
@@ -335,132 +328,78 @@ void AddReluKernel(gpuStream_t stream,
 }
 #endif
 
-template <typename DeviceContext, typename T>
-void FCFunctor<DeviceContext, T>::operator()(const DeviceContext& dev_ctx,
-                                             const int M,
-                                             const int N,
-                                             const int K,
-                                             const T* X,
-                                             const T* W,
-                                             T* Y,
-                                             const T* B,
-                                             bool relu,
-                                             bool padding_weights) {
-  PADDLE_ENFORCE_EQ(padding_weights,
-                    false,
-                    errors::PermissionDenied(
-                        "Weight padding in fc can not be used in GPU scope."));
-  auto blas = funcs::GetBlas<DeviceContext, T>(dev_ctx);
-  blas.GEMM(CblasNoTrans,
-            CblasNoTrans,
-            M,
-            N,
-            K,
-            static_cast<T>(1.0),
-            X,
-            W,
-            static_cast<T>(0.0),
-            Y);
-  if (B == NULL) {
-    return;
-  }
-
-  // M * N
-  AddReluKernel(dev_ctx.stream(), M, N, Y, B, relu);
+template <typename T>
+void LaunchFcQuantKernel(const T* input,
+                         int8_t* output,
+                         float scale,
+                         int m,
+                         int n,
+                         int round_type,
+                         float max_bound,
+                         float min_bound,
+                         gpuStream_t stream) {
+  ::phi::LaunchQuantKernelWithVecSize<T>(
+      input, output, scale, m, n, round_type, max_bound, min_bound, stream);
 }
 
-template class FCFunctor<GPUContext, float16>;
-template class FCFunctor<GPUContext, float>;
-template class FCFunctor<GPUContext, double>;
-
-template <typename DeviceContext, typename T>
-void FCInt8Functor<DeviceContext, T>::operator()(
-    const DeviceContext& dev_ctx,
-    const int M,
-    const int N,
-    const int K,
-    const T* X,
-    const DenseTensor* w_tensor,
-    T* Y,
-    float scale_in,
-    std::vector<float> scale_weights,
-    int quant_round_type,
-    float quant_max_bound,
-    float quant_min_bound,
-    const T* B,
-    bool relu,
-    bool padding_weights) {
-  PADDLE_ENFORCE_EQ(padding_weights,
-                    false,
-                    errors::PermissionDenied(
-                        "Weight padding in fc can not be used in GPU scope."));
-  const int8_t* W = w_tensor->data<int8_t>();
-
-  DenseTensor quant_x_tensor, quant_y_tensor;
-  quant_x_tensor.Resize({M, K});
-  quant_y_tensor.Resize({M, N});
-  dev_ctx.template Alloc<int8_t>(&quant_x_tensor,
-                                 quant_x_tensor.numel() * sizeof(int8_t));
-  dev_ctx.template Alloc<int32_t>(&quant_y_tensor,
-                                  quant_y_tensor.numel() * sizeof(int32_t));
-  LaunchQuantKernelWithVecSize<T>(X,
-                                  quant_x_tensor.data<int8_t>(),
-                                  scale_in,
-                                  M,
-                                  K,
-                                  quant_round_type,
-                                  quant_max_bound,
-                                  quant_min_bound,
-                                  dev_ctx.stream());
-
-  MatmulKernel<int8_t, GPUContext>(
-      dev_ctx, quant_x_tensor, *w_tensor, false, false, &quant_y_tensor);
-
-  DenseTensor scale_weights_dev;
-  scale_weights_dev.Resize({N});
-  dev_ctx.template Alloc<float>(&scale_weights_dev,
-                                scale_weights_dev.numel() * sizeof(float));
-  float* scale_weights_dev_ptr = scale_weights_dev.data<float>();
-#ifdef PADDLE_WITH_HIP
-  hipMemcpyAsync(scale_weights_dev_ptr,
-                 scale_weights.data(),
-                 N * sizeof(float),
-                 hipMemcpyHostToDevice);
-#else
-  cudaMemcpyAsync(scale_weights_dev_ptr,
-                  scale_weights.data(),
-                  N * sizeof(float),
-                  cudaMemcpyHostToDevice);
-#endif
-
-  phi::backends::gpu::GpuLaunchConfig config;
-  if (N % DequantKernelVecSize == 0) {
-    config = phi::backends::gpu::GetGpuLaunchConfig1D(
-        dev_ctx, M * N, DequantKernelVecSize);
-  } else {
-    config = phi::backends::gpu::GetGpuLaunchConfig1D(dev_ctx, M * N, 1);
-  }
-  LaunchDequantKernelWithScaleOfInputAndWeight(quant_y_tensor.data<int32_t>(),
-                                               Y,
-                                               M,
-                                               N,
-                                               dev_ctx.stream(),
-                                               &config,
-                                               scale_in,
-                                               scale_weights_dev_ptr,
-                                               quant_max_bound);
-
-  if (B == NULL) {
-    return;
-  }
-
-  // M * N
-  AddReluKernel(dev_ctx.stream(), M, N, Y, B, relu);
+template <typename T>
+void LaunchFcDequantKernel(const GPUContext& dev_ctx,
+                           const int32_t* input,
+                           T* output,
+                           int m,
+                           int n,
+                           float quant_in_scale,
+                           const float* quant_weight_scale,
+                           float quant_max_bound) {
+  const int vec_size = n % DequantKernelVecSize == 0 ? DequantKernelVecSize : 1;
+  auto config = backends::gpu::GetGpuLaunchConfig1D(dev_ctx, m * n, vec_size);
+  ::phi::LaunchDequantKernelWithScaleOfInputAndWeight(input,
+                                                      output,
+                                                      m,
+                                                      n,
+                                                      dev_ctx.stream(),
+                                                      &config,
+                                                      quant_in_scale,
+                                                      quant_weight_scale,
+                                                      quant_max_bound);
 }
 
-template class FCInt8Functor<GPUContext, float16>;
-template class FCInt8Functor<GPUContext, float>;
-template class FCInt8Functor<GPUContext, double>;
+template void AddReluKernel<float>(
+    gpuStream_t, int, int, float*, const float*, bool);
+template void AddReluKernel<double>(
+    gpuStream_t, int, int, double*, const double*, bool);
+
+template void LaunchFcQuantKernel<float16>(
+    const float16*, int8_t*, float, int, int, int, float, float, gpuStream_t);
+template void LaunchFcQuantKernel<float>(
+    const float*, int8_t*, float, int, int, int, float, float, gpuStream_t);
+template void LaunchFcQuantKernel<double>(
+    const double*, int8_t*, float, int, int, int, float, float, gpuStream_t);
+
+template void LaunchFcDequantKernel<float16>(const GPUContext&,
+                                             const int32_t*,
+                                             float16*,
+                                             int,
+                                             int,
+                                             float,
+                                             const float*,
+                                             float);
+template void LaunchFcDequantKernel<float>(const GPUContext&,
+                                           const int32_t*,
+                                           float*,
+                                           int,
+                                           int,
+                                           float,
+                                           const float*,
+                                           float);
+template void LaunchFcDequantKernel<double>(const GPUContext&,
+                                            const int32_t*,
+                                            double*,
+                                            int,
+                                            int,
+                                            float,
+                                            const float*,
+                                            float);
 
 }  // namespace funcs
 }  // namespace phi