[SYCL][UR][HIP] Hip adapter multi device ctx #11105

hdelan · 2023-09-07T16:25:30Z

Allow multi device context in HIP plugin. This makes a default context contain all of the devices in a given platform.

hdelan · 2023-09-07T16:26:57Z

Same as #10737 but for HIP.

This move to a multi device context makes context interop tricky for reasons outlined here #10975 . I have suggested that we deprecate context interop for CUDA and HIP plugins.

Images are still not supported in AMDGPU, so I will most likely need to revisit some of this code relating to images once the tests are working.

hdelan · 2023-09-07T16:34:59Z

sycl/plugins/unified_runtime/ur/adapters/hip/memory.cpp

@@ -279,29 +270,16 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemGetInfo(ur_mem_handle_t hMemory,
 /// \return UR_RESULT_SUCCESS
 UR_APIEXPORT ur_result_t UR_APICALL
 urMemGetNativeHandle(ur_mem_handle_t hMem, ur_native_handle_t *phNativeMem) {


Note that I have made this entry point UNSUPPORTED for the reasons listed in comment.

d1c92cb changed ScopedContext to take a device instead of a context thus sematically changing its meaning. This rename makes it clear what the intented usages are.

jandres742

L0 review requested here, but no L0 code.

+1 just for to allow merging.

This reverts commit c2a3cf7.

npmiller · 2023-10-05T10:46:27Z

sycl/plugins/unified_runtime/ur/adapters/hip/context.cpp

@@ -121,8 +119,10 @@ urContextRetain(ur_context_handle_t hContext) {

 UR_APIEXPORT ur_result_t UR_APICALL urContextGetNativeHandle(
    ur_context_handle_t hContext, ur_native_handle_t *phNativeContext) {
+  // FIXME this only returns the native context of the first device in the
+  // SYCL context. This entry point should be deprecated.


Linking the PR doing it:

[SYCL][CUDA][HIP] Deprecate context interop for CUDA and HIP #10975

npmiller · 2023-10-05T10:49:18Z

sycl/plugins/unified_runtime/ur/adapters/hip/context.hpp

+/// aforementioned custom `hipCtx_t`s.
+/// The HIP documentation, confirmed with performance analysis, suggest using
+/// the Primary context whenever possible. The Primary context is also used by
+/// the HIP Runtime API. For UR applications to interop with HIP Runtime API,


Is this actually true for HIP as well or just copied over from CUDA? Is there a HIP runtime API and a HIP driver API?

Hmm yeah I think this is not relevant to HIP. Will remove

sycl/plugins/unified_runtime/ur/adapters/hip/device.hpp

npmiller · 2023-10-05T14:34:33Z

sycl/plugins/unified_runtime/ur/adapters/hip/usm.cpp

@@ -68,7 +68,7 @@ UR_APIEXPORT ur_result_t UR_APICALL USMFreeImpl(ur_context_handle_t hContext,
                                                void *pMem) {
  ur_result_t Result = UR_RESULT_SUCCESS;
  try {
-    ScopedContext Active(hContext->getDevice());
+    // Free does not require a scopedDevice


Suggested change

// Free does not require a scopedDevice

// Free does not require a ScopedDevice

npmiller · 2023-10-05T14:36:45Z

sycl/plugins/unified_runtime/ur/adapters/hip/usm.cpp

@@ -151,7 +152,6 @@ urUSMGetMemAllocInfo(ur_context_handle_t hContext, const void *pMem,
  UrReturnHelper ReturnValue(propValueSize, pPropValue, pPropValueSizeRet);

  try {
-    ScopedContext Active(hContext->getDevice());


Isn't this needed for the hipPointerGetAttributes?

No apparently not

npmiller · 2023-10-05T15:45:17Z

sycl/plugins/unified_runtime/ur/adapters/hip/memory.cpp

+      // First allocation will be made at urMemBufferCreate if context only
+      // has one device
+      if (PerformInitialCopy && hContext->NumDevices == 1) {
+        // Operates on the default stream of the current CUDA context.


Suggested change

// Operates on the default stream of the current CUDA context.

// Operates on the default stream of the current HIP context.

npmiller · 2023-10-05T15:46:10Z

sycl/plugins/unified_runtime/ur/adapters/hip/memory.cpp

+          // to make buffer data available on device before any other UR
+          // call uses it.
+          hipStream_t defaultStream = 0;
+          UR_CHECK_ERROR(hipStreamSynchronize(defaultStream));


Isn't hipMemcpyHtoD synchronous? Also the comment says cuMemCpyHtoD

npmiller · 2023-10-05T15:48:24Z

sycl/plugins/unified_runtime/ur/adapters/hip/memory.cpp

+    if (URMemObj != nullptr) {
+      // First allocation will be made at urMemBufferCreate if context only
+      // has one device
+      if (PerformInitialCopy && hContext->NumDevices == 1) {


We might need to do the initial copy even if there's more devices, if COPY_HOST_PTR is used

npmiller · 2023-10-05T15:54:39Z

sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp

    }

-    hipArray *SrcArray = hImageSrc->Mem.SurfaceMem.getArray();
+    hipArray *SrcArray = ImageSrc->getArray(hQueue->getDevice());


Would it be cleaner to collapse allocateMemObjOnDeviceIfNeeded into the getters? You could have them do the allocation if needed when we try to get a device pointer.

We could even potentially have getters do the migration too, but that might not work in all cases like kernel arguments.

That's a good idea I think. Let me try to figure that out a bit

npmiller · 2023-10-05T16:16:09Z

sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp

+    // Use the default stream if copying from another device
+    hipStream_t HIPStream = DeviceToCopyFrom == hQueue->getDevice()
+                                ? hQueue->getNextTransferStream()
+                                : hipStream_t{0};


How does this work for hipMemcpyDToD then? does the stream have to be for the origin device and not for the destination device?

EnqueueMemBufferRead is only to read a buffer to the host, so a DtoD is never needed, only a DtoH from the most up to date device

So DtoH requires a stream from that particular device?

hdelan requested review from a team as code owners September 7, 2023 16:25

hdelan requested a review from sergey-semenov September 7, 2023 16:25

hdelan force-pushed the hip-adapter-multi-device-ctx branch 2 times, most recently from 9402e7a to 97c04eb Compare September 7, 2023 16:33

hdelan commented Sep 7, 2023

View reviewed changes

hdelan force-pushed the hip-adapter-multi-device-ctx branch from 97c04eb to e32e26e Compare September 7, 2023 16:51

hdelan requested a review from a team as a code owner September 7, 2023 16:51

hdelan force-pushed the hip-adapter-multi-device-ctx branch 3 times, most recently from b93c0fa to d6ea231 Compare September 13, 2023 20:43

ldrumm and others added 5 commits September 14, 2023 15:30

[SYCL][UR][HIP] s/ScopedContext/ScopedDevice

3605f74

d1c92cb changed ScopedContext to take a device instead of a context thus sematically changing its meaning. This rename makes it clear what the intented usages are.

Multi dev ctx for hip adapter

3c73acc

Delete interop-buffer-hip

3b51fac

Add some description and add further checks for Images

d07903e

Format

0de2831

hdelan force-pushed the hip-adapter-multi-device-ctx branch from e707ee2 to 0de2831 Compare September 14, 2023 14:31

jandres742 approved these changes Sep 22, 2023

View reviewed changes

Move base event into device

c2a3cf7

hdelan force-pushed the hip-adapter-multi-device-ctx branch from fb6a43f to c2a3cf7 Compare September 25, 2023 19:01

hdelan added 2 commits October 2, 2023 10:06

Remove mem dep in kernel

f2ede30

Revert "Move base event into device"

5c105ce

This reverts commit c2a3cf7.

hdelan mentioned this pull request Oct 4, 2023

[SYCL][CUDA] Cuda adapter multi device context #10737

Closed

npmiller reviewed Oct 5, 2023

View reviewed changes

hdelan mentioned this pull request Oct 26, 2023

[HIP] Hip adapter multi dev ctx oneapi-src/unified-runtime#999

Merged

hdelan closed this Nov 3, 2023

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[SYCL][UR][HIP] Hip adapter multi device ctx #11105

[SYCL][UR][HIP] Hip adapter multi device ctx #11105

hdelan commented Sep 7, 2023

hdelan commented Sep 7, 2023 •

edited

Loading

hdelan Sep 7, 2023 •

edited

Loading

jandres742 left a comment

npmiller Oct 5, 2023

npmiller Oct 5, 2023

hdelan Oct 6, 2023

npmiller Oct 5, 2023

npmiller Oct 5, 2023

hdelan Oct 23, 2023

npmiller Oct 5, 2023

npmiller Oct 5, 2023

npmiller Oct 5, 2023

npmiller Oct 5, 2023

hdelan Oct 6, 2023

npmiller Oct 5, 2023

hdelan Oct 6, 2023

npmiller Oct 10, 2023

	// Free does not require a scopedDevice
	// Free does not require a ScopedDevice

	// Operates on the default stream of the current CUDA context.
	// Operates on the default stream of the current HIP context.

[SYCL][UR][HIP] Hip adapter multi device ctx #11105

[SYCL][UR][HIP] Hip adapter multi device ctx #11105

Conversation

hdelan commented Sep 7, 2023

hdelan commented Sep 7, 2023 • edited Loading

hdelan Sep 7, 2023 • edited Loading

Choose a reason for hiding this comment

jandres742 left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

hdelan commented Sep 7, 2023 •

edited

Loading

hdelan Sep 7, 2023 •

edited

Loading