From 73a5448a13fb29bc83d0af0690b19bd39484b0e2 Mon Sep 17 00:00:00 2001 From: qili93 Date: Tue, 20 Feb 2024 10:33:27 +0000 Subject: [PATCH 1/2] [DCU] fix compile error, test=develop --- cmake/generic.cmake | 14 ++++++---- paddle/phi/CMakeLists.txt | 7 +++-- paddle/phi/core/visit_type.h | 2 +- paddle/phi/kernels/CMakeLists.txt | 26 +++++++++++++++++++ paddle/phi/kernels/funcs/CMakeLists.txt | 5 ++++ paddle/phi/kernels/gpu/binomial_kernel.cu | 9 +++++++ .../phi/kernels/gpu/top_p_sampling_kernel.cu | 2 +- .../gpu/triangular_solve_grad_kernel.cu | 10 +++++++ .../kernels/gpu/triangular_solve_kernel.cu | 10 +++++++ paddle/phi/kernels/gpu/unique_kernel.cu | 5 ++++ paddle/phi/kernels/gpudnn/pool_gpudnn.h | 12 +++++++-- 11 files changed, 91 insertions(+), 11 deletions(-) diff --git a/cmake/generic.cmake b/cmake/generic.cmake index f31b82bfff0f7..1bc258ff1261d 100644 --- a/cmake/generic.cmake +++ b/cmake/generic.cmake @@ -764,11 +764,11 @@ function(hip_library TARGET_NAME) "${multiValueArgs}" ${ARGN}) if(hip_library_SRCS) # FindHIP.cmake defined hip_add_library, HIP_SOURCE_PROPERTY_FORMAT is requried if no .cu files found - if(NOT (${CMAKE_CURRENT_SOURCE_DIR} MATCHES ".*/operators" - OR ${CMAKE_CURRENT_SOURCE_DIR} MATCHES ".*/phi/kernels")) - set_source_files_properties(${hip_library_SRCS} - PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) - endif() + # if(NOT (${CMAKE_CURRENT_SOURCE_DIR} MATCHES ".*/operators" + # OR ${CMAKE_CURRENT_SOURCE_DIR} MATCHES ".*/phi/kernels")) + # set_source_files_properties(${hip_library_SRCS} + # PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) + # endif() if(hip_library_SHARED OR hip_library_shared) # build *.so hip_add_library(${TARGET_NAME} SHARED ${hip_library_SRCS}) else() @@ -782,6 +782,10 @@ function(hip_library TARGET_NAME) endif() # cpplint code style foreach(source_file ${hip_library_SRCS}) + if(NOT ${source_file} MATCHES "\\.cu$") + set_source_files_properties(${source_file} + PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) + endif() string(REGEX REPLACE "\\.[^.]*$" "" source ${source_file}) if(EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/${source}.h) list(APPEND hip_library_HEADERS diff --git a/paddle/phi/CMakeLists.txt b/paddle/phi/CMakeLists.txt index 5556b457883ae..7325aef2202b5 100644 --- a/paddle/phi/CMakeLists.txt +++ b/paddle/phi/CMakeLists.txt @@ -134,8 +134,11 @@ if(WITH_GPU) DEPS ${PHI_DEPS}) elseif(WITH_ROCM) - hip_add_library(phi ${PHI_BUILD_TYPE} ${PHI_SRCS}) - target_link_libraries(phi ${PHI_DEPS}) + hip_library( + phi ${PHI_BUILD_TYPE} + SRCS ${PHI_SRCS} + DEPS ${PHI_DEPS}) + elseif(WITH_XPU_KP) xpu_library( phi ${PHI_BUILD_TYPE} diff --git a/paddle/phi/core/visit_type.h b/paddle/phi/core/visit_type.h index 069e737d95607..7ee12e26d7d0e 100644 --- a/paddle/phi/core/visit_type.h +++ b/paddle/phi/core/visit_type.h @@ -355,7 +355,7 @@ namespace phi { "`"); \ } \ }() -#if defined(PADDLE_WITH_XPU) +#if defined(PADDLE_WITH_XPU) || defined(PADDLE_WITH_HIP) #define PD_VISIT_ALL_TYPES(TYPE, NAME, ...) \ [&] { \ const auto& __dtype__ = TYPE; \ diff --git a/paddle/phi/kernels/CMakeLists.txt b/paddle/phi/kernels/CMakeLists.txt index 892f9f0a4fe09..85697df11bc56 100644 --- a/paddle/phi/kernels/CMakeLists.txt +++ b/paddle/phi/kernels/CMakeLists.txt @@ -227,6 +227,32 @@ if(NOT WITH_CUDNN_FRONTEND) "fusion/gpu/max_pool2d_v2_kernel.cu") endif() +# Note(qili93): remove kernels not supported on DCU yet +if(WITH_ROCM) + list( + REMOVE_ITEM + kernel_cu + "gpu/affine_grid_grad_kernel.cu" + "gpu/apply_per_channel_scale_kernel.cu" + "gpu/cholesky_solve_kernel.cu" + "gpu/eigh_kernel.cu" + "gpu/eigvalsh_kernel.cu" + "gpu/lstsq_kernel.cu" + "gpu/lu_kernel.cu" + "gpu/matrix_rank_kernel.cu" + "gpu/matrix_rank_tol_kernel.cu" + "gpu/multiclass_nms3_kernel.cu" + "gpu/put_along_axis_grad_kernel.cu" + "gpu/put_along_axis_kernel.cu" + "gpu/qr_kernel.cu" + "gpu/svd_kernel.cu" + "gpudnn/mha_cudnn_frontend.cu" + "fusion/gpu/block_multi_head_attention_kernel.cu" + "fusion/gpu/fused_bn_add_activation_grad_kernel.cu" + "fusion/gpu/fused_bn_add_activation_kernel.cu" + "fusion/gpu/fusion_transpose_flatten_concat_kernel.cu") +endif() + set(cc_search_pattern "*.cc" "cpu/*.cc" diff --git a/paddle/phi/kernels/funcs/CMakeLists.txt b/paddle/phi/kernels/funcs/CMakeLists.txt index 999625cf3dfb4..d124e269e5c00 100644 --- a/paddle/phi/kernels/funcs/CMakeLists.txt +++ b/paddle/phi/kernels/funcs/CMakeLists.txt @@ -15,4 +15,9 @@ if(WITH_GPU OR WITH_ROCM) "*.cu") endif() +# Note(qili93): remove kernels not supported on DCU yet +if(WITH_ROCM) + list(REMOVE_ITEM func_cu_srcs "weight_only_gemv.cu") +endif() + collect_srcs(kernels_srcs SRCS ${func_cc_srcs} ${func_cu_srcs}) diff --git a/paddle/phi/kernels/gpu/binomial_kernel.cu b/paddle/phi/kernels/gpu/binomial_kernel.cu index 402de44a6c210..2686485dbeb65 100644 --- a/paddle/phi/kernels/gpu/binomial_kernel.cu +++ b/paddle/phi/kernels/gpu/binomial_kernel.cu @@ -74,8 +74,13 @@ __device__ int64_t btrs( const T m = std::floor((n + 1) * p); while (1) { +#ifdef __NVCC__ U = static_cast(curand_uniform(&state)) - 0.5; V = static_cast(curand_uniform(&state)); +#elif __HIPCC__ + U = static_cast(hiprand_uniform(&state)) - 0.5; + V = static_cast(hiprand_uniform(&state)); +#endif us = 0.5 - std::abs(U); k = static_cast(std::floor((2 * a / us + b) * U + c)); @@ -118,7 +123,11 @@ __device__ int64_t binomial_inversion( #endif while (1) { +#ifdef __NVCC__ unif = static_cast(curand_uniform(&state)); +#elif __HIPCC__ + unif = static_cast(hiprand_uniform(&state)); +#endif T geom = std::ceil(std::log(unif) / logprob); geom_sum += geom; if (geom_sum > n) { diff --git a/paddle/phi/kernels/gpu/top_p_sampling_kernel.cu b/paddle/phi/kernels/gpu/top_p_sampling_kernel.cu index fae94de3c6787..c86db04c9ba94 100644 --- a/paddle/phi/kernels/gpu/top_p_sampling_kernel.cu +++ b/paddle/phi/kernels/gpu/top_p_sampling_kernel.cu @@ -310,7 +310,7 @@ __device__ __forceinline__ void BlockReduce(Pair shared_max[], if (*beam >= MaxLength) break; } else { #ifdef PADDLE_WITH_HIP - uint64 mask = 0; + unsigned mask = 0u; mask = __ballot(true); if (tid_max / WARP_SIZE == wid) { if (__shfl_down(*beam, tid_max % WARP_SIZE, WARP_SIZE) == MaxLength) diff --git a/paddle/phi/kernels/gpu/triangular_solve_grad_kernel.cu b/paddle/phi/kernels/gpu/triangular_solve_grad_kernel.cu index 67861b282529b..bbb98e4c05b14 100644 --- a/paddle/phi/kernels/gpu/triangular_solve_grad_kernel.cu +++ b/paddle/phi/kernels/gpu/triangular_solve_grad_kernel.cu @@ -15,6 +15,7 @@ #include "paddle/phi/core/kernel_registry.h" #include "paddle/phi/kernels/impl/triangular_solve_grad_kernel_impl.h" +#ifdef PADDLE_WITH_CUDA PD_REGISTER_KERNEL(triangular_solve_grad, GPU, ALL_LAYOUT, @@ -23,3 +24,12 @@ PD_REGISTER_KERNEL(triangular_solve_grad, double, phi::dtype::complex, phi::dtype::complex) {} +#else // PADDLE_WITH_HIP +// blas_impl.hip.h not support CUBlas::TRSM for complex +PD_REGISTER_KERNEL(triangular_solve_grad, + GPU, + ALL_LAYOUT, + phi::TriangularSolveGradKernel, + float, + double) {} +#endif diff --git a/paddle/phi/kernels/gpu/triangular_solve_kernel.cu b/paddle/phi/kernels/gpu/triangular_solve_kernel.cu index 342b8e3885d7b..2c2bb299f8594 100644 --- a/paddle/phi/kernels/gpu/triangular_solve_kernel.cu +++ b/paddle/phi/kernels/gpu/triangular_solve_kernel.cu @@ -123,6 +123,7 @@ void TriangularSolveKernel(const Context& dev_ctx, } // namespace phi +#ifdef PADDLE_WITH_CUDA PD_REGISTER_KERNEL(triangular_solve, GPU, ALL_LAYOUT, @@ -131,3 +132,12 @@ PD_REGISTER_KERNEL(triangular_solve, double, phi::dtype::complex, phi::dtype::complex) {} +#else // PADDLE_WITH_HIP +// blas_impl.hip.h not support CUBlas::TRSM for complex +PD_REGISTER_KERNEL(triangular_solve, + GPU, + ALL_LAYOUT, + phi::TriangularSolveKernel, + float, + double) {} +#endif diff --git a/paddle/phi/kernels/gpu/unique_kernel.cu b/paddle/phi/kernels/gpu/unique_kernel.cu index c90766728a87d..18b23dcab1c39 100644 --- a/paddle/phi/kernels/gpu/unique_kernel.cu +++ b/paddle/phi/kernels/gpu/unique_kernel.cu @@ -26,7 +26,12 @@ #include #include +#ifdef PADDLE_WITH_CUDA #include "cub/cub.cuh" +#else +#include +namespace cub = hipcub; +#endif #include "paddle/phi/backends/gpu/gpu_context.h" #include "paddle/phi/common/memory_utils.h" #include "paddle/phi/core/kernel_registry.h" diff --git a/paddle/phi/kernels/gpudnn/pool_gpudnn.h b/paddle/phi/kernels/gpudnn/pool_gpudnn.h index cd2758109f28c..9bf7e4de20569 100644 --- a/paddle/phi/kernels/gpudnn/pool_gpudnn.h +++ b/paddle/phi/kernels/gpudnn/pool_gpudnn.h @@ -35,13 +35,21 @@ class CudnnIndexType; template <> class CudnnIndexType { public: - static const cudnnDataType_t type = CUDNN_DATA_INT32; +#ifdef PADDLE_WITH_CUDA + static const dnnDataType_t type = CUDNN_DATA_INT32; +#else + static const dnnDataType_t type = miopenInt32; +#endif }; template <> class CudnnIndexType { public: - static const cudnnDataType_t type = CUDNN_DATA_INT8; +#ifdef PADDLE_WITH_CUDA + static const dnnDataType_t type = CUDNN_DATA_INT8; +#else + static const dnnDataType_t type = miopenInt8; +#endif }; inline GPUDNNDataLayout GetLayoutFromStr(std::string data_format) { From 82f212615198e512594eb7b1d09412a1cdbc72bd Mon Sep 17 00:00:00 2001 From: qili93 Date: Tue, 20 Feb 2024 10:48:34 +0000 Subject: [PATCH 2/2] remove comments --- cmake/generic.cmake | 6 ------ 1 file changed, 6 deletions(-) diff --git a/cmake/generic.cmake b/cmake/generic.cmake index 1bc258ff1261d..4131b919ee8ce 100644 --- a/cmake/generic.cmake +++ b/cmake/generic.cmake @@ -763,12 +763,6 @@ function(hip_library TARGET_NAME) cmake_parse_arguments(hip_library "${options}" "${oneValueArgs}" "${multiValueArgs}" ${ARGN}) if(hip_library_SRCS) - # FindHIP.cmake defined hip_add_library, HIP_SOURCE_PROPERTY_FORMAT is requried if no .cu files found - # if(NOT (${CMAKE_CURRENT_SOURCE_DIR} MATCHES ".*/operators" - # OR ${CMAKE_CURRENT_SOURCE_DIR} MATCHES ".*/phi/kernels")) - # set_source_files_properties(${hip_library_SRCS} - # PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) - # endif() if(hip_library_SHARED OR hip_library_shared) # build *.so hip_add_library(${TARGET_NAME} SHARED ${hip_library_SRCS}) else()