Skip to content

Commit 1af5a50

Browse files
committed
Remove dead code of ht struct
Signed-off-by: Ke Wen <kwen@nvidia.com>
1 parent 3619159 commit 1af5a50

1 file changed

Lines changed: 2 additions & 33 deletions

File tree

contrib/nccl_ep/nccl_ep.cc

Lines changed: 2 additions & 33 deletions
Original file line numberDiff line numberDiff line change
@@ -1366,35 +1366,6 @@ struct ncclEpHandle {
13661366
int hidden_int4;
13671367

13681368
union {
1369-
struct {
1370-
// Both intranode and internode
1371-
int* recv_counter;
1372-
int* recv_counter_device;
1373-
int* internal_recv_expert_counter_host = nullptr;
1374-
int received_token_count = -1;
1375-
ncclNDTensor_t rank_token_counts;
1376-
ncclNDTensor_t expert_token_counts;
1377-
ncclNDTensor_t token_rank_mask;
1378-
ncclNDTensor_t global_channel_prefix;
1379-
ncclNDTensor_t nvl_send_head;
1380-
ncclNDTensor_t recv_global_channel_prefix;
1381-
1382-
// Internode only
1383-
int* rdma_recv_counter;
1384-
int* rdma_recv_counter_device;
1385-
int rdma_received_token_count = -1;
1386-
std::optional<ncclNDTensor_t> rdma_rank_token_counts;
1387-
ncclNDTensor_t rdma_channel_prefix;
1388-
ncclNDTensor_t recv_rdma_rank_prefix;
1389-
ncclNDTensor_t recv_global_rank_prefix;
1390-
ncclNDTensor_t rdma_send_head;
1391-
ncclNDTensor_t recv_source_metadata;
1392-
ncclNDTensor_t recv_rdma_channel_prefix;
1393-
1394-
// Intranode only
1395-
ncclNDTensor_t inter_rank_token_offsets;
1396-
ncclNDTensor_t recv_token_source_map;
1397-
} ht;
13981369
struct {
13991370
// packed tensors for LL
14001371
ncclNDTensor_t expert_recv_source_indices;
@@ -1539,10 +1510,8 @@ struct ncclEpHandle {
15391510
cached_mode(false),
15401511
num_scales(0),
15411512
hidden_int4(0) {
1542-
// Zero the entire union (ht, ll, hybridep share memory)
1543-
// Use max size to ensure all members are zeroed
1544-
constexpr size_t union_size = std::max({sizeof(ht), sizeof(ll), sizeof(hybridep)});
1545-
memset(static_cast<void*>(&ht), 0, union_size);
1513+
constexpr size_t union_size = std::max(sizeof(ll), sizeof(hybridep));
1514+
memset(static_cast<void*>(&ll), 0, union_size);
15461515
}
15471516

15481517
~ncclEpHandle() {

0 commit comments

Comments
 (0)