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

fix fused_seqpool_cvm if branch bug #118

Closed
wants to merge 9 commits into from
Closed

fix fused_seqpool_cvm if branch bug #118

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
9 commits
Select commit Hold shift + click to select a range
8072027
update bkcl and xdnn.
tiancaitzp Aug 6, 2024
e889882
modify xctr version
guojunwork Aug 7, 2024
ea430f3
sync xpu's fused_seqpool_cvm_op with gpu
xymyeah Aug 16, 2024
7339575
modify PushCopy lm overflow
guojunwork Aug 29, 2024
a512603
add xpu boxps in cmake
jack603047588 Sep 5, 2024
218e8f9
add xpu kl2 boxps in cmake
jack603047588 Sep 5, 2024
9a034d1
fix parquet bug
guojunwork Sep 9, 2024
e7c5cb0
modify parquet bug
guojunwork Sep 9, 2024
7d8c53f
fix fused_seqpool_cvm logic
YaoCheng8667 Sep 13, 2024
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
7 changes: 7 additions & 0 deletions cmake/external/box_ps.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -22,6 +22,13 @@ IF((NOT DEFINED BOX_PS_VER) OR (NOT DEFINED BOX_PS_URL))
#SET(BOX_PS_URL "http://box-ps.gz.bcebos.com/box_ps.tar.gz" CACHE STRING "" FORCE)
SET(BOX_PS_URL "data-im.baidu.com:/home/work/var/CI_DATA/im/static/box_ps.tar.gz/box_ps.tar.gz.43" CACHE STRING "" FORCE)
ENDIF()
IF(WITH_XPU OR WITH_XPU_KP)
SET(BOX_PS_URL "data-im.baidu.com:/home/work/var/CI_DATA/im/static/box_ps.xpu.kl2.tar.gz/box_ps.xpu.kl2.tar.gz.1")
IF(WITH_XPU_XRE5)
message(WARNING "WITH_XPU_XRE5 is ON, but boxps only configuired with kl2")
ENDIF()
ENDIF()

MESSAGE(STATUS "BOX_PS_NAME: ${BOX_PS_NAME}, BOX_PS_URL: ${BOX_PS_URL}")
SET(BOX_PS_SOURCE_DIR "${THIRD_PARTY_PATH}/box_ps")
SET(BOX_PS_DOWNLOAD_DIR "${BOX_PS_SOURCE_DIR}/src/${BOX_PS_PROJECT}")
Expand Down
8 changes: 4 additions & 4 deletions cmake/external/xpu.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -92,9 +92,9 @@ set(XPU_XCCL_URL
"${XPU_XCCL_BASE_URL}/${XPU_XCCL_DIR_NAME}.tar.gz"
CACHE STRING "" FORCE)
#https://klx-sdk-release-public.su.bcebos.com/xdnn_train/dev/20240312/xctr.tar.gz
#https://klx-sdk-release-public.su.bcebos.com/xdnn_train/dev/20221130/xctr.tar.gz
#https://klx-sdk-release-public.su.bcebos.com/xdnn_train/dev/20240807/xctr.tar.gz
set(XPU_XCTR_URL
https://klx-sdk-release-public.su.bcebos.com/xdnn_train/dev/20240701/xctr.tar.gz
https://klx-sdk-release-public.su.bcebos.com/xdnn_train/dev/20240909/xctr.tar.gz
CACHE STRING "" FORCE)

#"data-im.baidu.com:/home/work/var/CI_DATA/im/static/pack_paddle_depence.sh/pack_paddle_depence.sh"
Expand All @@ -105,7 +105,7 @@ set(XPU_PACK_DEPENCE_URL
if(WITH_XPU_XRE5)
add_definitions(-DPADDLE_WITH_XPU_XRE5)
set(XPU_XRE_BASE_VERSION "5.0.9.1")
set(XPU_XCCL_BASE_VERSION "1.2.7.2")
set(XPU_XCCL_BASE_VERSION "1.2.8.2")
set(XPU_XRE_BASE_URL
"https://klx-sdk-release-public.su.bcebos.com/xre/kl3-release/${XPU_XRE_BASE_VERSION}"
)
Expand Down Expand Up @@ -133,7 +133,7 @@ if (WITH_BOX_PS OR WITH_XPU_KP)
CACHE STRING "" FORCE)
#"https://klx-sdk-release-public.su.bcebos.com/xdnn/release/2.6.0.1/${XPU_XDNN_DIR_NAME}.tar.gz"
set(XPU_XDNN_URL
"https://klx-sdk-release-public.su.bcebos.com/xdnn_train/dev/paddlebox/xdnn_master/20240718/${XPU_XDNN_DIR_NAME}.tar.gz"
"https://klx-sdk-release-public.su.bcebos.com/xdnn_train/dev/paddlebox/xdnn_master/20240731/${XPU_XDNN_DIR_NAME}.tar.gz"
CACHE STRING "" FORCE)
set(SCALOPUS_URL
"https://klx-sdk-release-public.su.bcebos.com/xdnn_train/dev/paddlebox/20230306/scalopus.tar.gz"
Expand Down
20 changes: 15 additions & 5 deletions paddle/fluid/framework/fleet/box_wrapper_kernel.kps
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -935,7 +935,8 @@ __global__ void PullCopy(const TEmbedxOp* op,
int thread_id = cluster_id() * ncores + cid;
int nthreads = cluster_num() * ncores;

const int buf_length = 30;
const int total_local_mem = 6500;
const int buf_length = total_local_mem / ((hidden_size + pull_float_num + 3) * 4);
int per_thread_len = roundup_div(total_length, nthreads);
int per_thread_loop_count = roundup_div(per_thread_len, buf_length);
int per_thread_per_loop_len = roundup_div(per_thread_len, per_thread_loop_count);
Expand Down Expand Up @@ -1016,7 +1017,7 @@ __global__ void PullCopyMerged(const TEmbedxOp* op,
int thread_id = cluster_id() * ncores + cid;
int nthreads = cluster_num() * ncores;

const int total_local_mem = 7000;
const int total_local_mem = 6500;
const int buf_length = total_local_mem / ((hidden_size + pull_float_num + 3) * 4);
const int same_key_idx_buf_len = buf_length;
int per_thread_len = roundup_div(merged_length, nthreads);
Expand Down Expand Up @@ -1127,6 +1128,8 @@ inline void FeaturePullCopy(const paddle::platform::Place& place,
break;
}
}

CHECK(6500 / ((hidden_size + pull_float_num + 3) * 4) > 0) << "xpu lm is overflow hidden_size " << hidden_size << " pull_float_num " << pull_float_num;
if (xpu_merged_idx == nullptr) {
PullCopy<TEmbedxOp><<<8, 64, stream>>>(d_op,
scale,
Expand Down Expand Up @@ -1579,7 +1582,12 @@ __global__ void PushCopyNNCrossDiscontinuous(const TExpandPushGetOp* op,
GM2LM((__global_ptr__ float*)sm_src_expand_vals_ptr[lm_key2slot[j]] + lm_slot_inner_offset[j] * expand_embed_dim,
lm_src_expand_vals, expand_embed_dim * sizeof(float));
for (int k = 1 + skip_offset + hidden_size; k < push_float_num; ++k) {
lm_dst_vals[j * push_float_num + k] = lm_src_expand_vals[k - skip_offset - 1] * -bs;
// lm_dst_vals[j * push_float_num + k] = lm_src_expand_vals[k - skip_offset - 1] * -bs;
lm_dst_vals[j * push_float_num + k] = lm_op[0].get(lm_src_expand_vals,
0,
k - skip_offset - 1 - hidden_size,
hidden_size, // hidden_size is useless
expand_embed_dim) * -bs;
}
mfence();
} else {
Expand Down Expand Up @@ -1748,7 +1756,9 @@ __global__ void PushCopy(unsigned long long* src_vals,
int thread_id = cluster_id() * ncores + cid;
int nthreads = cluster_num() * ncores;

const int buf_length = 40; // max 2048 float
//const int buf_length = 40; // max 2048 float
const int total_local_mem = 6500;
const int buf_length = total_local_mem / ((hidden_size + push_float_num + 3) * 4);
int per_thread_len = roundup_div(total_length, nthreads);
int per_thread_loop_count = roundup_div(per_thread_len, buf_length);
int per_thread_per_loop_len =
Expand Down Expand Up @@ -1911,7 +1921,7 @@ void BoxWrapperKernel::CopyForPush(
CHECK(false) << "FeaturePushCopyVariable not implement";
} else {
// FeaturePushCopy
// TODO:
CHECK(6500 / ((hidden_size + push_float_num + 3) * 4) > 0) << "xpu lm is overflow hidden_size " << hidden_size << " push_float_num " << push_float_num;
PushCopy<<<8, 64, stream>>>(reinterpret_cast<unsigned long long*>(gm_src_ptr), push_grad_values, push_offset,
push_float_num_, c_total_length, hidden_size, batch_size, total_dims,
skip_offset, cvm_offset, key2slot, slots, slot_num);
Expand Down
Loading