Skip to content

Using ibv_exp_reg_mr instead of ibv_reg_mr #4

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

Open
e-ago wants to merge 5 commits into
base: master
Choose a base branch
from
Open

Using ibv_exp_reg_mr instead of ibv_reg_mr #4

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
cc4bf62
mp_register() extended with implicit ODP
e-ago Jul 2, 2018
8cccb9a
mp_register() does cuPointerGetAttribute only without odp. comm_pingp…
e-ago Jul 2, 2018
57fb8bf
minor fix debug print
e-ago Jul 2, 2018
ea34f23
Adding error checks. Can't use ODP in combination with GMEM
e-ago Jul 2, 2018
e08034b
mp_register now uses ibv_exp_reg_mr instead of ibv_reg_mr
e-ago Jul 2, 2018
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
4 changes: 2 additions & 2 deletions benchmarks/mp_pingpong_all.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -842,8 +842,8 @@ int main (int argc, char *argv[])
CUDA_CHECK(cudaMemset(rbuf_d, 0, buf_size));
}

MP_CHECK(mp_register(sbuf_d, buf_size, &sreg));
MP_CHECK(mp_register(rbuf_d, buf_size, &rreg));
MP_CHECK(mp_register(sbuf_d, buf_size, &sreg, 0));
MP_CHECK(mp_register(rbuf_d, buf_size, &rreg, 0));

if (!my_rank) fprintf(stdout, "%10d", size);

Expand Down
4 changes: 2 additions & 2 deletions benchmarks/mp_pingpong_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -232,8 +232,8 @@ double sr_exchange (MPI_Comm comm, int size, int iter_count, int validate, doubl
CUDA_CHECK(cudaMalloc((void **)&rbuf_d, size*iter_count));
CUDA_CHECK(cudaMemset(rbuf_d, 0, size*iter_count));

MP_CHECK(mp_register(sbuf_d, size*iter_count, &sreg));
MP_CHECK(mp_register(rbuf_d, size*iter_count, &rreg));
MP_CHECK(mp_register(sbuf_d, size*iter_count, &sreg, 0));
MP_CHECK(mp_register(rbuf_d, size*iter_count, &rreg, 0));

if (validate) {
mp_dbg_msg("initializing the buffer \n");
Expand Down
4 changes: 2 additions & 2 deletions benchmarks/mp_pingpong_kernel_stream.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -455,8 +455,8 @@ int main (int argc, char *argv[])
CUDA_CHECK(cudaMalloc((void **)&rbuf_d, buf_size));
CUDA_CHECK(cudaMemset(rbuf_d, 0, buf_size));

MP_CHECK(mp_register(sbuf_d, buf_size, &sreg));
MP_CHECK(mp_register(rbuf_d, buf_size, &rreg));
MP_CHECK(mp_register(sbuf_d, buf_size, &sreg, 0));
MP_CHECK(mp_register(rbuf_d, buf_size, &rreg, 0));

if (!my_rank) {
if (prof_init(&prof_normal, 1000, 1000, "1us", 100, 1, tags)) {
Expand Down
4 changes: 2 additions & 2 deletions benchmarks/mp_pingpong_kernel_stream_latency.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -539,8 +539,8 @@ int main (int argc, char *argv[])
CUDA_CHECK(cudaMalloc((void **)&rbuf_d, buf_size));
CUDA_CHECK(cudaMemset(rbuf_d, 0, buf_size));

MP_CHECK(mp_register(sbuf_d, buf_size, &sreg));
MP_CHECK(mp_register(rbuf_d, buf_size, &rreg));
MP_CHECK(mp_register(sbuf_d, buf_size, &sreg, 0));
MP_CHECK(mp_register(rbuf_d, buf_size, &rreg, 0));

if (!my_rank) fprintf(stdout, "%10d", size);

Expand Down
4 changes: 2 additions & 2 deletions benchmarks/mp_pingpong_kernel_stream_latency_mpi.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -719,8 +719,8 @@ int main (int argc, char *argv[])
CUDA_CHECK(cudaMemset(rbuf_d, 0, buf_size));
}

MP_CHECK(mp_register(sbuf_d, buf_size, &sreg));
MP_CHECK(mp_register(rbuf_d, buf_size, &rreg));
MP_CHECK(mp_register(sbuf_d, buf_size, &sreg, 0));
MP_CHECK(mp_register(rbuf_d, buf_size, &rreg, 0));

if (!my_rank) fprintf(stdout, "%10d", size);

Expand Down
4 changes: 2 additions & 2 deletions benchmarks/mp_pingpong_kernel_stream_mpi.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -641,8 +641,8 @@ int main (int argc, char *argv[])
CUDA_CHECK(cudaMemset(rbuf_d, 0, buf_size));
}

MP_CHECK(mp_register(sbuf_d, buf_size, &sreg));
MP_CHECK(mp_register(rbuf_d, buf_size, &rreg));
MP_CHECK(mp_register(sbuf_d, buf_size, &sreg, 0));
MP_CHECK(mp_register(rbuf_d, buf_size, &rreg, 0));

if (my_rank == 0) {
if (prof_init(&prof_normal, 1000, 1000, "1us", 100, 1, tags)) {
Expand Down
4 changes: 2 additions & 2 deletions benchmarks/mp_pingpong_kernel_stream_wait_send.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -388,8 +388,8 @@ int main (int c, char *v[])
CUDA_CHECK(cudaMalloc((void **)&rbuf_d, buf_size));
CUDA_CHECK(cudaMemset(rbuf_d, 0, buf_size));

MP_CHECK(mp_register(sbuf_d, buf_size, &sreg));
MP_CHECK(mp_register(rbuf_d, buf_size, &rreg));
MP_CHECK(mp_register(sbuf_d, buf_size, &sreg, 0));
MP_CHECK(mp_register(rbuf_d, buf_size, &rreg, 0));

if (!my_rank) {
if (prof_init(&prof_normal, 1000, 1000, "1us", 100, 1, tags)) {
Expand Down
4 changes: 2 additions & 2 deletions benchmarks/mp_producer_consumer_kernel_stream.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -373,8 +373,8 @@ int main (int c, char *v[])
CUDA_CHECK(cudaMalloc((void **)&rbuf_d, buf_size));
CUDA_CHECK(cudaMemset(rbuf_d, 0, buf_size));

MP_CHECK(mp_register(sbuf_d, buf_size, &sreg));
MP_CHECK(mp_register(rbuf_d, buf_size, &rreg));
MP_CHECK(mp_register(sbuf_d, buf_size, &sreg, 0));
MP_CHECK(mp_register(rbuf_d, buf_size, &rreg, 0));

if (!my_rank) {
if (prof_init(&prof_normal, 1000, 1000, "1us", 100, 1, tags)) {
Expand Down
4 changes: 2 additions & 2 deletions benchmarks/mp_sendrecv_kernel_stream.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -359,8 +359,8 @@ int main (int c, char *v[])
CUDA_CHECK(cudaMalloc((void **)&rbuf_d, buf_size));
CUDA_CHECK(cudaMemset(rbuf_d, 0, buf_size));

MP_CHECK(mp_register(sbuf_d, buf_size, &sreg));
MP_CHECK(mp_register(rbuf_d, buf_size, &rreg));
MP_CHECK(mp_register(sbuf_d, buf_size, &sreg, 0));
MP_CHECK(mp_register(rbuf_d, buf_size, &rreg, 0));

if (!my_rank) {
if (prof_init(&prof_normal, 1000, 1000, "1us", 100, 1, tags)) {
Expand Down
23 changes: 20 additions & 3 deletions comm_library/comm.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -246,11 +246,11 @@ int comm_init(MPI_Comm comm, int gpuId)
iomb();

DBG("registering ready_table size=%zd\n", table_size);
MP_CHECK(mp_register(ready_table, table_size, &ready_table_reg));
MP_CHECK(mp_register(ready_table, table_size, &ready_table_reg, 0));
DBG("creating ready_table window\n");
MP_CHECK(mp_window_create(ready_table, table_size, &ready_table_win));
DBG("registering remote_ready_table\n");
MP_CHECK(mp_register(remote_ready_values, table_size, &remote_ready_values_reg));
MP_CHECK(mp_register(remote_ready_values, table_size, &remote_ready_values_reg, 0));

comm_initialized = 1;

Expand Down Expand Up @@ -706,7 +706,24 @@ int comm_register(void *buf, size_t size, comm_reg_t *creg)

if (!*reg) {
DBG("registering buffer %p\n", buf);
MP_CHECK(mp_register(buf, size, reg));
MP_CHECK(mp_register(buf, size, reg, 0));
}

out:
return ret;
}

int comm_register_odp(comm_reg_t *creg)
Copy link
Collaborator

@drossetti drossetti Jul 2, 2018

Choose a reason for hiding this comment

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

The name of this API is not expressive enough.
You are effectively exposing all/most of the local process memory.

Explicit/implicit ODP support has limitations, see https://community.mellanox.com/docs/DOC-2898, e.g. have to check for capability.

{
assert(comm_initialized);
int ret = 0;
int retcode;
mp_reg_t *reg = (mp_reg_t*)creg;
assert(reg);

if (!*reg) {
DBG("registering implicit ODP MR\n");
MP_CHECK(mp_register(NULL, 0, reg, IBV_EXP_ACCESS_ON_DEMAND));
Copy link
Collaborator

@drossetti drossetti Jul 2, 2018

Choose a reason for hiding this comment

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

The documentations says:
"To register an Implicit ODP MR, in addition to the IBV_EXP_ACCESS_ON_DEMAND access flag, use in->addr = 0 and in->length = IBV_EXP_IMPLICIT_MR_SIZE."
so 0 is not a good size.
I would provide a high level API in libmp, one which does all the appropriate tests.

Copy link
Collaborator

@drossetti drossetti Jul 3, 2018

Choose a reason for hiding this comment

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

I noticed that the size parameter is ignored by mp_register() when ODP is enabled. So technically this code is correct

Copy link
Collaborator

@drossetti drossetti Jul 3, 2018

Choose a reason for hiding this comment

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

I wonder what happens if the app calls comm_register_odp() multiple times. Does ibv_reg_mr() succeed? Does it return the same mr / lkey ?

}

out:
Expand Down
1 change: 1 addition & 0 deletions comm_library/comm.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -125,6 +125,7 @@ extern "C" {
int comm_prepare_wait_all(int count, comm_request_t *creqs);
comm_dev_descs_t comm_prepared_requests();
int comm_register(void *buf, size_t size, comm_reg_t *creg);
int comm_register_odp(comm_reg_t *creg);
int comm_deregister(comm_reg_t *creg);
int comm_select_device(int mpiRank);

Expand Down
Loading