Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

[multi-node] Fix sage infer hang #287

Merged
merged 10 commits into from
May 17, 2023
Merged

[multi-node] Fix sage infer hang #287

Show file tree
Hide file tree
Changes from 7 commits
Commits
Show all changes
10 commits
Select commit Hold shift + click to select a range
e98f5e7
try to fix infer
DesmonDay May 9, 2023
5cf8c8c
Merge branch 'gpugraph_new' of https://github.com/xuewujiao/Paddle in…
DesmonDay May 10, 2023
639d590
temp commit
DesmonDay May 11, 2023
58671ba
fix 60B infer
DesmonDay May 16, 2023
fdaefb2
Merge branch 'gpugraph_new' of https://github.com/xuewujiao/Paddle in…
DesmonDay May 16, 2023
bfd76d1
bug fix
DesmonDay May 16, 2023
49a5188
delete unused config
DesmonDay May 16, 2023
d83d897
update completion
DesmonDay May 17, 2023
32398c7
change function name
DesmonDay May 17, 2023
04a2b18
fix buf_size
DesmonDay May 17, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
64 changes: 57 additions & 7 deletions paddle/fluid/framework/data_feed.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3563,11 +3563,20 @@ int GraphDataGenerator::FillInferBuf() {
}

size_t device_key_size = h_device_keys_len_[0][infer_cursor];
total_row_[0] =
(global_infer_node_type_start[infer_cursor] + conf_.buf_size <=
device_key_size)
? conf_.buf_size
: device_key_size - global_infer_node_type_start[infer_cursor];
if (conf_.is_multi_node) {
int get_type_end = dynamic_adjust_total_row();
if (get_type_end) {
total_row_[0] = device_key_size - global_infer_node_type_start[infer_cursor];
} else {
total_row_[0] = conf_.buf_size;
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

再提一个建议，这里最好不要修改buf_size_。
如果后面还有其他类型节点需要infer，这里buf_size_得不到reset，不友好。
建议这样，
if (global_reach_end)
total_row_[0] = device_key_size - global_infer_node_type_start[infer_cursor];
else
remain = device_key_size - global_infer_node_type_start[infer_cursor];
if local_reach_end
total_row_[0] = remain / 2;
else
total_row_[0] = remain;

Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

done，不过最后那个else那里，应该是total_row_[0] = conf_.buf_size;

}
} else {
total_row_[0] =
(global_infer_node_type_start[infer_cursor] + conf_.buf_size <=
device_key_size)
? conf_.buf_size
: device_key_size - global_infer_node_type_start[infer_cursor];
}

uint64_t *d_type_keys =
reinterpret_cast<uint64_t *>(d_device_keys_[0][infer_cursor]->ptr());
Expand All @@ -3580,9 +3589,11 @@ int GraphDataGenerator::FillInferBuf() {
sample_stream_);
cudaStreamSynchronize(sample_stream_);
}
VLOG(1) << "cursor: " << infer_cursor
VLOG(1) << "gpuid: " << conf_.gpuid
<< " cursor: " << infer_cursor
<< " start: " << global_infer_node_type_start[infer_cursor]
<< " num: " << total_row_[0];
<< " num: " << total_row_[0]
<< " device_key_size: " << device_key_size;
infer_node_start_ = global_infer_node_type_start[infer_cursor];
global_infer_node_type_start[infer_cursor] += total_row_[0];
infer_node_end_ = global_infer_node_type_start[infer_cursor];
Expand Down Expand Up @@ -4042,6 +4053,45 @@ int GraphDataGenerator::dynamic_adjust_batch_num_for_sage() {
return new_batch_size;
}

int GraphDataGenerator::dynamic_adjust_total_row() {
platform::CUDADeviceGuard guard(conf_.gpuid);
auto gpu_graph_ptr = GraphGpuWrapper::GetInstance();
auto &global_infer_node_type_start =
gpu_graph_ptr->global_infer_node_type_start_[conf_.gpuid];
auto &infer_cursor = gpu_graph_ptr->infer_cursor_[conf_.thread_id];
size_t device_key_size = h_device_keys_len_[0][infer_cursor];
int reach_end = global_infer_node_type_start[infer_cursor] + conf_.buf_size >=
device_key_size;
auto send_buff = memory::Alloc(
place_,
2 * sizeof(int),
phi::Stream(reinterpret_cast<phi::StreamId>(sample_stream_)));
int *send_buff_ptr = reinterpret_cast<int *>(send_buff->ptr());
cudaMemcpyAsync(send_buff_ptr,
&reach_end,
sizeof(int),
cudaMemcpyHostToDevice,
sample_stream_);
cudaStreamSynchronize(sample_stream_);
auto comm =
platform::NCCLCommContext::Instance().Get(0, place_.GetDeviceId());
PADDLE_ENFORCE_GPU_SUCCESS(platform::dynload::ncclAllReduce(&send_buff_ptr[0],
&send_buff_ptr[1],
1,
ncclInt,
ncclMax,
comm->comm(),
sample_stream_));
int get_all_reach_end = 0;
cudaMemcpyAsync(&get_all_reach_end,
&send_buff_ptr[1],
sizeof(int),
cudaMemcpyDeviceToHost,
sample_stream_);
cudaStreamSynchronize(sample_stream_);
return get_all_reach_end;
}

} // namespace framework
} // namespace paddle
#endif
1 change: 1 addition & 0 deletions paddle/fluid/framework/data_feed.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -981,6 +981,7 @@ class GraphDataGenerator {
int get_pass_end() { return pass_end_; }
void clear_gpu_mem();
int dynamic_adjust_batch_num_for_sage();
int dynamic_adjust_total_row();

protected:
void DoSageForInfer();
Expand Down