Segment creation for CUDA MK

[Mert Hidayetoglu]

Hello,

 

I am working on a simple program that creates GPU segments and performs simply put and get operations.

 

First, I developed a CPU version using gex_Segment_Attach() function which does not require explicit memory allocation by the user. This function internally allocates memory and returns the segment handle where I can obtain the pointer using gex_Segment_QueryAddr(). I can use the pointer for put and get operations.

 

Now I am trying to do the same with CUDA memory kind. As far as I know, I should use gex_Segment_Create() function to create segments that lives in GPUs.

 

int gex_Segment_Create(

            gex_Segment_t   *segment_p,  // OUT

            gex_Client_t    client,

            gex_Addr_t      address,

            uintptr_t       length,

            gex_MK_t        kind,

            gex_Flags_t     flags);

      // Valid flags:

      //   NONE - zero is currently required

 

I wonder what the third argument means. Is it a pointer to a device buffer that is allocated by the user (i.e., using cudaMalloc())?

 

If so, I have to set device myself, using cudaSetDevice(), right? Also, should the allocation size match with the fourth argument and should it be a multiple GASNET_PAGESIZE? In this case, gex_Segment_QueryAddr() becomes redundant unless the user loses the pointer to the  device buffer, is that right?

 

Best regareds,

Mert

[Paul H. Hargrove]

Mert,

The `address` argument can be a device buffer allocated by the user, as you have described. 

I do not believe we've specified that the length be a multiple of `GASNET_PAGESIZE` for device memory, and would appreciate you letting us know if use of non-aligned sizes fails for you.
The CUDA device used by GASNet-EX will be the one used in the `gex_MK_Create()` call used to generate the `kind` argument to `gex_Segment_Create()`.

However, it is also permissible to pass `NULL` for the `address` argument to `gex_Segment_Create()`.
In that case, GASNet-EX will allocate the memory, using the device corresponding to the `kind` argument.

This case will require your code to use `gex_Segment_QueryAddr()` to determine the buffer's address.

Regardless of how the memory is allocated, be aware that `gex_EP_BindSegment()` and `gex_EP_PublishBoundSegment()` will be required before you can use the segment for put and get operations.  Those are implicit in `gex_Segment_Attach()`, but must be done manually otherwise.

Let us know if you have any other questions, or if something above is not clear.

-Paul

--
You received this message because you are subscribed to the Google Groups "gasnet-users" group.
To unsubscribe from this group and stop receiving emails from it, send an email to gasnet-users...@lbl.gov.
To view this discussion on the web visit https://groups.google.com/a/lbl.gov/d/msgid/gasnet-users/SJ0PR02MB73579CDC9E023A1BF6AD12A8D1EB2%40SJ0PR02MB7357.namprd02.prod.outlook.com.

--

Paul H. Hargrove <PHHar...@lbl.gov>
Pronouns: he, him, his

Computer Languages & Systems Software (CLaSS) Group

Computer Science Department

Lawrence Berkeley National Laboratory

[Mert Hidayetoglu]

Thanks, Paul!

 

I also included gex_EP_BindSegment() and gex_EP_PublishBoundSegment().

 

extern int gex_EP_PublishBoundSegment(

                gex_TM_t               tm,

                gex_EP_t               *eps,    // IN

                size_t                 num_eps,

                gex_Flags_t            flags);

 

Above, does *eps represent the endpoint associated with the caller, all endpoints, or only remote endpoints that I wish to communicate from the caller?

 

My code currently crashes, and I am not sure this is the right place to paste the code (83 lines). Shall I send you in a separate email?

 

Best regards,

Mert

[Paul H. Hargrove]

Mert,

The `eps` argument points to an array of `num_eps` elements.
It should include all of the local `gex_EP_t`'s which remote callers will wish to communicate with.

If you have code, error messages, or other things that you are not comfortable sharing in the public forum (even if just to be kind to other readers INBOXes), then feel free to send them to gasnet...@lbl.gov 

-Paul

[Mert Hidayetoglu]

Thanks, Paul. That’s what I thought!

 

I am still getting an error message. Just to be kind to other readers, I followed up the conversation at gasnet...@lbl.gov.

 

Thanks for your help.