[CINN][New Hardware Update] SYCL first PR #69554

feixi21 · 2024-11-20T10:26:56Z

PR Category

CINN

PR Types

New features

Description

SYCL后端第一个PR

添加HygonDCUArchSYCL，并实现HygonDCUArchSYCL相应处理流程

paddle-bot · 2024-11-20T10:27:01Z

你的PR提交成功，感谢你对开源项目的贡献!
请关注后续CI自动化测试结果，详情请参考Paddle-CI手册。
Your PR has been submitted. Thanks for your contribution!
Please wait for the result of CI firstly. See Paddle CI Manual for details.

fix_bugs fix_bugs

ZhangX-21 · 2024-12-03T14:03:14Z

paddle/cinn/hlir/framework/pir/op_lowering_impl.cc

@@ -536,6 +542,14 @@ std::vector<ir::LoweredFunc> OpLowererImpl::DoOpLower(
          } else {
            op_func_arg_tensors->push_back(expr.as_tensor_ref());
          }
+        },


代码实现没有差异性，std::variant<common::HygonDCUArchHIP, common::HygonDCUArchSYCL>处理？

已参照修改意见修改。

ZhangX-21 · 2024-12-03T14:03:29Z

paddle/cinn/hlir/op/contrib/gather_nd.cc

@@ -222,6 +222,9 @@ std::shared_ptr<framework::OpStrategy> StrategyForGatherNd(
          },
          [&](common::HygonDCUArchHIP) {
            pe::IRGpuScheduleInjective(ir_sch, output_shapes.front(), target);
+          },


已参照修改意见修改。

ZhangX-21 · 2024-12-03T14:03:56Z

paddle/cinn/hlir/op/contrib/repeat.cc

@@ -224,6 +224,9 @@ std::shared_ptr<framework::OpStrategy> StrategyForRepeat(
          },
          [&](common::HygonDCUArchHIP) {
            pe::IRGpuScheduleInjective(ir_sch, output_shapes.front(), target);
+          },


已参照修改意见修改。

ZhangX-21 · 2024-12-03T14:14:03Z

paddle/cinn/runtime/sycl/sycl_backend_api.cc

+                                ::sycl::info::device_type::gpu)[0]
+                                .get_backend());
+      },
+      [&](common::X86Arch) { LOG(FATAL) << "SYCL Not supported this arch \n"; },


同上, arch 为sycl /hip 的backend为都是:sycl::backend::ext_oneapi_hip，不会歧义吗？

已参照修改意见修改。

ZhangX-21 · 2024-12-03T14:14:55Z

paddle/cinn/runtime/sycl/sycl_backend_api.cc

+  this->queues.resize(this->devices.size());
+  // sycl::backend -> Target::Arch
+  switch (backend) {
+    case ::sycl::backend::ext_oneapi_hip:


在更早的时候是不是就拦截掉了？

之前函数返回SYCLBackendAPI类内的私有变量arch，因此之前此处需要根据backend初始化this->arch，已参照修改意见修改。

ZhangX-21 · 2024-12-03T14:16:18Z

paddle/cinn/runtime/sycl/sycl_backend_api.cc

+              .get_info<::sycl::info::device::sub_group_sizes>();
+      size_t max_sub_group_size =
+          *max_element(std::begin(sub_group_sizes), std::end(sub_group_sizes));
+      rv = static_cast<int>(max_sub_group_size);


缺少break?

已参照修改意见修改。

ZhangX-21 · 2024-12-03T14:20:33Z

paddle/cinn/runtime/sycl/sycl_module.cc

+SYCLModule::~SYCLModule() {
+  VLOG(3) << "destructor for SYCLModule";
+  if (so_handler_ != nullptr) {
+    // dlclose(so_handler_);


冗余代码？

已参照修改意见修改。

modify according to comment modify according to comment modify according to comment

ZhangX-21 · 2024-12-04T12:05:37Z

paddle/cinn/backends/codegen_cuda_host.h

@@ -60,6 +60,13 @@ class CodeGenGpuHost : public CodeGenHost {
          } else {
            return CodeGenHost::Visit(op);
          }
+        },


这个也是？

HygonDCUArchHIP需要判定是否为call_hip_kernel，HygonDCUArchSYCL需要判定是否为call_sycl_kernel，所需要判定的op不同，所以这里不能合并。

ZhangX-21 · 2024-12-04T12:06:07Z

paddle/cinn/backends/codegen_device_util.cc

@@ -218,6 +218,11 @@ void detail::CollectBucketStrategyHostFunctionVisitor::ProcessLoweredFunc(
      [&](common::HygonDCUArchHIP) {
 #ifdef CINN_WITH_HIP
        shared_mem_bytes = CalculateSharedMemory(func);
+#endif


这个也是吧

修改了一下HygonDCUArchSYCL的对应逻辑。

ZhangX-21 · 2024-12-04T12:07:15Z

paddle/cinn/backends/compiler.cc

@@ -286,6 +287,10 @@ std::string Compiler::GetSourceCode(const ir::Module& module) {
 #else
        CINN_NOT_IMPLEMENTED
 #endif
+      },
+      [&](common::HygonDCUArchSYCL) -> std::string {
+        PADDLE_THROW(phi::errors::Unimplemented(


CINN_NOT_IMPLEMENTED？

已参照修改意见修改。

ZhangX-21 · 2024-12-04T12:08:03Z

paddle/cinn/common/target.cc

@@ -79,6 +86,11 @@ int GetRuntimeArchImpl(HygonDCUArchHIP) {
      "HygonDCUArchHIP not supported GetRuntimeArch!"));
 }

+int GetRuntimeArchImpl(HygonDCUArchSYCL) {
+  PADDLE_THROW(::common::errors::InvalidArgument(
+      "HygonDCUArchSYCL not supported GetRuntimeArch!"));


CINN_NOT_IMPLEMENTED

已参照修改意见修改。

ZhangX-21 · 2024-12-04T12:08:57Z

paddle/cinn/hlir/framework/pir/op_lowering_impl.cc

@@ -372,6 +372,12 @@ std::vector<ir::LoweredFunc> OpLowererImpl::PostProcess(
 #ifdef CINN_WITH_HIP
            optim::EliminateCommonGlobalMemoryRead(&(func_body));
            optim::OptimizeExprGPU(&(func_body));
+#endif


已参照修改意见修改。

ZhangX-21 · 2024-12-04T12:17:31Z

paddle/cinn/runtime/sycl/sycl_util.cc

+    ::sycl::range<3> Grid(grid_z, grid_y, grid_x);
+    ::sycl::range<3> Block(block_z, block_y, block_x);
+    // need malloc_shared
+    // LOG(INFO) << "kernel args :" << (float* )(*(void **)(kernel_args[0]))[0]


冗余代码

已参照修改意见修改。

fix_bugs fix_bugs

ZhangX-21 · 2024-12-09T06:19:54Z

paddle/cinn/optim/vectorize_loops.cc

+                .Visit(&new_forloop->body);
+          },
+          [&](std::variant<common::HygonDCUArchHIP, common::HygonDCUArchSYCL>) {
+            setNvHygon();


setNvHygon会不会有歧义那？海光的设备为啥和NV的有关系那？

这里是在NV设备与海光设备上均调用相同的函数，无需修改

ZhangX-21 · 2024-12-09T06:23:09Z

paddle/cinn/runtime/sycl/sycl_backend_api.cc

+      [&](common::ARMArch) { LOG(FATAL) << "SYCL Not supported this arch \n"; },
+      [&](common::NVGPUArch) { backend = ::sycl::backend::ext_oneapi_cuda; },
+      [&](std::variant<common::HygonDCUArchHIP, common::HygonDCUArchSYCL>) {
+        backend = ::sycl::backend::ext_oneapi_hip;


common::HygonDCUArchHIP对应:hip::backend::ext_oneapi_hip?
common::HygonDCUArchSYCL对应::sycl::backend::ext_oneapi_sycl?

已参照修改意见修改。

fix_bugs fix_bugs fix_bugs

phlrain · 2024-12-17T10:45:22Z

paddle/cinn/pybind/runtime.cc

+cinn_buffer_t *CreateBufferFromNumpyImpl(common::HygonDCUArchSYCL,
+                                         py::array data) {
+  PADDLE_THROW(::common::errors::Unimplemented("CINN old obsolete code!"));
+}


这部分代码是不是可以清理了

phlrain · 2024-12-17T10:48:24Z

paddle/cinn/runtime/sycl/cinn_sycl_runtime_source.h

+
+#undef CINN_CUDA_INDEX_ADD
+
+int cinn_sycl_resize_bilinear(const int *buf,


这些函数是不是没有地方在调用的，是不是可以清理下

phlrain · 2024-12-17T10:49:40Z

paddle/cinn/runtime/sycl/sycl_util.h

+                           int b3,
+                           int b4);
+
+void cinn_call_cnnl_conv2d_forward(void* v_args,


类似这种接口，是不是也不在需要了

XiaoguangHu01

LGTM

feixi21 added 5 commits November 20, 2024 18:05

add HygonDCU_SYCL

535f9a6

DPCPP_CMAKE

382fb80

fix_merge_conflict

f30d3f7

add HygonDCUSYCL part two

2f3b708

add sycl runtime

83aa06c

paddle-bot bot added the contributor External developers label Nov 20, 2024

feixi21 added 3 commits November 21, 2024 10:31

fix_bugs

3aadea4

fix_bugs

15bd597

fix_bugs

4334c66

feixi21 force-pushed the paddle-sycl-PR1 branch from bce8bd8 to 4334c66 Compare November 28, 2024 08:34

feixi21 added 2 commits November 29, 2024 14:49

fix_codestyle

54d4824

add define sycl fp16 bf16 judgment

2dc4989

fix_bugs fix_bugs

feixi21 force-pushed the paddle-sycl-PR1 branch from ac7d13f to 2dc4989 Compare December 3, 2024 02:42

ZhangX-21 reviewed Dec 3, 2024

View reviewed changes

feixi21 force-pushed the paddle-sycl-PR1 branch 2 times, most recently from 2b9205d to bbd123a Compare December 4, 2024 07:51

modify according to comment

980cfbe

modify according to comment modify according to comment modify according to comment

feixi21 force-pushed the paddle-sycl-PR1 branch from bbd123a to 980cfbe Compare December 4, 2024 08:17

ZhangX-21 reviewed Dec 4, 2024

View reviewed changes

modify according to comment

f11c934

fix_bugs fix_bugs

feixi21 force-pushed the paddle-sycl-PR1 branch from 90cb5a8 to f11c934 Compare December 5, 2024 03:30

ZhangX-21 reviewed Dec 9, 2024

View reviewed changes

ZhangX-21 approved these changes Dec 10, 2024

View reviewed changes

feixi21 force-pushed the paddle-sycl-PR1 branch from 5310255 to 46fdb22 Compare December 11, 2024 02:51

fix_bugs

16b5843

fix_bugs fix_bugs fix_bugs

feixi21 force-pushed the paddle-sycl-PR1 branch 2 times, most recently from e7b6aab to 16b5843 Compare December 13, 2024 01:34

feixi21 and others added 2 commits December 13, 2024 09:50

fix_merge_conflicts

bb1c8ce

Merge branch 'develop' into paddle-sycl-PR1

3fb9665

phlrain approved these changes Dec 17, 2024

View reviewed changes

XiaoguangHu01 approved these changes Dec 18, 2024

View reviewed changes

phlrain merged commit aea56b7 into PaddlePaddle:develop Dec 18, 2024
28 checks passed

feixi21 deleted the paddle-sycl-PR1 branch April 3, 2025 08:17


		#undef CINN_CUDA_INDEX_ADD

		int cinn_sycl_resize_bilinear(const int *buf,

[CINN][New Hardware Update] SYCL first PR #69554

[CINN][New Hardware Update] SYCL first PR #69554

Uh oh!

Conversation

feixi21 commented Nov 20, 2024

PR Category

PR Types

Description

Uh oh!

paddle-bot bot commented Nov 20, 2024

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

XiaoguangHu01 left a comment

Choose a reason for hiding this comment

Uh oh!