"...git@developer.sourcefind.cn:renzhc/diffusers_dcu.git" did not exist on "4e0984db6cde2b9585be17464c8fec5ddfd61c89"
Unverified Commit 40365904 authored by Illia Silin's avatar Illia Silin Committed by GitHub
Browse files

fix clang format (#740)

parent e2ebc8e7
...@@ -611,11 +611,12 @@ struct DeviceGroupedGemmSoftmaxGemmPermute_Xdl_CShuffle ...@@ -611,11 +611,12 @@ struct DeviceGroupedGemmSoftmaxGemmPermute_Xdl_CShuffle
some_has_main_k_block_loop |= y; some_has_main_k_block_loop |= y;
} }
hipGetErrorString(hipMemcpyWithStream(arg.p_workspace_, hipGetErrorString(
arg.group_kernel_args_.data(), hipMemcpyWithStream(arg.p_workspace_,
arg.group_kernel_args_.size() * sizeof(GroupKernelArg), arg.group_kernel_args_.data(),
hipMemcpyHostToDevice, arg.group_kernel_args_.size() * sizeof(GroupKernelArg),
stream_config.stream_id_)); hipMemcpyHostToDevice,
stream_config.stream_id_));
float ave_time = 0; float ave_time = 0;
......
...@@ -655,7 +655,7 @@ struct DeviceGroupedContractionMultipleD_Xdl_CShuffle ...@@ -655,7 +655,7 @@ struct DeviceGroupedContractionMultipleD_Xdl_CShuffle
hipGetErrorString(hipMemcpyWithStream(arg.p_workspace_, hipGetErrorString(hipMemcpyWithStream(arg.p_workspace_,
arg.contraction_multi_d_kernel_args_.data(), arg.contraction_multi_d_kernel_args_.data(),
arg.contraction_multi_d_kernel_args_.size() * arg.contraction_multi_d_kernel_args_.size() *
sizeof(ContractionMultiDKernelArg), sizeof(ContractionMultiDKernelArg),
hipMemcpyHostToDevice, hipMemcpyHostToDevice,
stream_config.stream_id_)); stream_config.stream_id_));
......
...@@ -597,11 +597,12 @@ struct DeviceGroupedGemmMultipleD_Dl : public DeviceGroupedGemm<ALayout, ...@@ -597,11 +597,12 @@ struct DeviceGroupedGemmMultipleD_Dl : public DeviceGroupedGemm<ALayout,
} }
} }
hipGetErrorString(hipMemcpyWithStream(arg.p_workspace_, hipGetErrorString(
arg.gemm_desc_kernel_arg_.data(), hipMemcpyWithStream(arg.p_workspace_,
arg.gemm_desc_kernel_arg_.size() * sizeof(GemmKernelArg), arg.gemm_desc_kernel_arg_.data(),
hipMemcpyHostToDevice, arg.gemm_desc_kernel_arg_.size() * sizeof(GemmKernelArg),
stream_config.stream_id_)); hipMemcpyHostToDevice,
stream_config.stream_id_));
auto launch_kernel = [&](auto has_main_k_block_loop, auto launch_kernel = [&](auto has_main_k_block_loop,
auto has_double_tail_k_block_loop) { auto has_double_tail_k_block_loop) {
......
...@@ -548,12 +548,12 @@ struct DeviceGroupedGemm_Xdl : public DeviceGroupedGemm<ALayout, ...@@ -548,12 +548,12 @@ struct DeviceGroupedGemm_Xdl : public DeviceGroupedGemm<ALayout,
} }
} }
hipGetErrorString( hipGetErrorString(hipMemcpyWithStream(arg.p_workspace_,
hipMemcpyWithStream(arg.p_workspace_, arg.gemm_desc_kernel_arg_.data(),
arg.gemm_desc_kernel_arg_.data(), arg.gemm_desc_kernel_arg_.size() *
arg.gemm_desc_kernel_arg_.size() * sizeof(GemmBiasTransKernelArg), sizeof(GemmBiasTransKernelArg),
hipMemcpyHostToDevice, hipMemcpyHostToDevice,
stream_config.stream_id_)); stream_config.stream_id_));
float ave_time = 0; float ave_time = 0;
......
...@@ -406,11 +406,12 @@ struct DeviceGroupedGemmXdlSplitKCShuffle : public DeviceGroupedGemmSplitK<ALayo ...@@ -406,11 +406,12 @@ struct DeviceGroupedGemmXdlSplitKCShuffle : public DeviceGroupedGemmSplitK<ALayo
} }
} }
hip_check_error(hipMemcpyWithStream(arg.p_workspace_, hip_check_error(
arg.gemm_kernel_args_.data(), hipMemcpyWithStream(arg.p_workspace_,
arg.gemm_kernel_args_.size() * sizeof(GemmTransKernelArg), arg.gemm_kernel_args_.data(),
hipMemcpyHostToDevice, arg.gemm_kernel_args_.size() * sizeof(GemmTransKernelArg),
stream_config.stream_id_)); hipMemcpyHostToDevice,
stream_config.stream_id_));
float ave_time = 0; float ave_time = 0;
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment