Skip to content

[External Resource Interop][Semaphores] Added support for 2 more semaphore types #2751

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

Closed
wants to merge 1 commit into from
Closed

[External Resource Interop][Semaphores] Added support for 2 more semaphore types #2751

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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 clang/examples/DPCT/Runtime/cudaDestroyExternalMemory.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,7 @@
// Option: --use-experimental-features=bindless_images

void test(cudaExternalMemory_t extMem) {
// Start
cudaDestroyExternalMemory(extMem /*cudaExternalMemory_t*/);
// End
}
7 changes: 7 additions & 0 deletions clang/examples/DPCT/Runtime/cudaDestroyExternalSemaphore.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,7 @@
// Option: --use-experimental-features=bindless_images

void test(cudaExternalSemaphore_t extSem) {
// Start
cudaDestroyExternalSemaphore(extSem /*cudaExternalSemaphore_t*/);
// End
}
10 changes: 10 additions & 0 deletions clang/examples/DPCT/Runtime/cudaExternalMemoryGetMappedBuffer.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,10 @@
// Option: --use-experimental-features=bindless_images

void test(void **devPtr, cudaExternalMemory_t extMem,
const cudaExternalMemoryBufferDesc *bufferDesc) {
// Start
cudaExternalMemoryGetMappedBuffer(
devPtr /*void ***/, extMem /*cudaExternalMemory_t*/,
bufferDesc /*const cudaExternalMemoryBufferDesc **/);
// End
}
10 changes: 10 additions & 0 deletions clang/examples/DPCT/Runtime/cudaExternalMemoryGetMappedMipmappedArray.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,10 @@
// Option: --use-experimental-features=bindless_images

void test(cudaMipmappedArray_t *mipmap, cudaExternalMemory_t extMem,
const cudaExternalMemoryMipmappedArrayDesc *mipmapDesc) {
// Start
cudaExternalMemoryGetMappedMipmappedArray(
mipmap /*cudaMipmappedArray_t **/, extMem /*cudaExternalMemory_t*/,
mipmapDesc /*const cudaExternalMemoryMipmappedArrayDesc **/);
// End
}
10 changes: 10 additions & 0 deletions clang/examples/DPCT/Runtime/cudaImportExternalMemory.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,10 @@
// Option: --use-experimental-features=bindless_images

void test(cudaExternalMemory_t *extMem,
const cudaExternalMemoryHandleDesc *memHandleDesc) {
// Start
cudaImportExternalMemory(
extMem /*cudaExternalMemory_t **/,
memHandleDesc /*const cudaExternalMemoryHandleDesc **/);
// End
}
10 changes: 10 additions & 0 deletions clang/examples/DPCT/Runtime/cudaImportExternalSemaphore.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,10 @@
// Option: --use-experimental-features=bindless_images

void test(cudaExternalSemaphore_t *extSem,
const cudaExternalSemaphoreHandleDesc *semHandleDesc) {
// Start
cudaImportExternalSemaphore(
extSem /*cudaExternalSemaphore_t **/,
semHandleDesc /*const cudaExternalSemaphoreHandleDesc **/);
// End
}
16 changes: 16 additions & 0 deletions clang/examples/DPCT/Runtime/cudaSignalExternalSemaphoresAsync.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,16 @@
// Option: --use-experimental-features=bindless_images

void test(const cudaExternalSemaphore_t *extSemArray,
const cudaExternalSemaphoreSignalParams *paramsArray,
unsigned int numExtSems, cudaStream_t stream = 0) {
// Start
cudaSignalExternalSemaphoresAsync(
extSemArray /*const cudaExternalSemaphore_t **/,
paramsArray /*const cudaExternalSemaphoreSignalParams **/,
numExtSems /*unsigned int*/);
cudaSignalExternalSemaphoresAsync(
extSemArray /*const cudaExternalSemaphore_t **/,
paramsArray /*const cudaExternalSemaphoreSignalParams **/,
numExtSems /*unsigned int*/, stream /*cudaStream_t*/);
// End
}
16 changes: 16 additions & 0 deletions clang/examples/DPCT/Runtime/cudaWaitExternalSemaphoresAsync.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,16 @@
// Option: --use-experimental-features=bindless_images

void test(const cudaExternalSemaphore_t *extSemArray,
const cudaExternalSemaphoreWaitParams *paramsArray,
unsigned int numExtSems, cudaStream_t stream = 0) {
// Start
cudaWaitExternalSemaphoresAsync(
extSemArray /*const cudaExternalSemaphore_t **/,
paramsArray /*const cudaExternalSemaphoreWaitParams **/,
numExtSems /*unsigned int*/);
cudaWaitExternalSemaphoresAsync(
extSemArray /*const cudaExternalSemaphore_t **/,
paramsArray /*const cudaExternalSemaphoreWaitParams **/,
numExtSems /*unsigned int*/, stream /*cudaStream_t*/);
// End
}
12 changes: 12 additions & 0 deletions clang/lib/DPCT/RuleInfra/MapNames.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1601,6 +1601,18 @@ void MapNames::setExplicitNamespaceMap(
? getExpNamespace() +
"external_semaphore_handle_type::win32_nt_dx12_fence"
: "cudaExternalSemaphoreHandleTypeD3D12Fence")},
{"cudaExternalSemaphoreHandleTypeTimelineSemaphoreFd",
Copy link
Contributor

Choose a reason for hiding this comment

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

Where is the migration?

Copy link
Contributor Author

@TejaX-Alaghari TejaX-Alaghari Apr 4, 2025 •
edited
Loading

Choose a reason for hiding this comment

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

The migration for cudaExternalSemaphoreHandleTypeTimelineSemaphoreFd is external_semaphore_handle_type::timeline_fd here

std::make_shared<EnumNameRule>(
DpctGlobalInfo::useExtBindlessImages()
? getExpNamespace() +
"external_semaphore_handle_type::timeline_fd"
: "cudaExternalSemaphoreHandleTypeTimelineSemaphoreFd")},
{"cudaExternalSemaphoreHandleTypeTimelineSemaphoreWin32",
std::make_shared<EnumNameRule>(
DpctGlobalInfo::useExtBindlessImages()
? getExpNamespace() +
"external_semaphore_handle_type::timeline_win32_nt_handle"
: "cudaExternalSemaphoreHandleTypeTimelineSemaphoreWin32")},
{"NVSHMEM_TEAM_WORLD",
std::make_shared<EnumNameRule>("ISHMEM_TEAM_WORLD")},
{"NVSHMEM_TEAM_SHARED",
Expand Down
8 changes: 4 additions & 4 deletions clang/lib/DPCT/RulesLang/APINamesGraphicsInterop.inc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -201,11 +201,11 @@ CONDITIONAL_FACTORY_ENTRY(
makeCheckNot(CheckArgIsDefaultCudaStream(3)),
ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY("cudaSignalExternalSemaphoresAsync_v2",
CALL(MapNames::getDpctNamespace() +
"experimental::signal_external_semaphore",
"experimental::signal_external_semaphores",
ARG(0), ARG(1), ARG(2), ARG(3)))),
ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY("cudaSignalExternalSemaphoresAsync_v2",
CALL(MapNames::getDpctNamespace() +
"experimental::signal_external_semaphore",
"experimental::signal_external_semaphores",
ARG(0), ARG(1), ARG(2))))),
UNSUPPORT_FACTORY_ENTRY("cudaSignalExternalSemaphoresAsync_v2", Diagnostics::TRY_EXPERIMENTAL_FEATURE,
ARG("cudaSignalExternalSemaphoresAsync_v2"),
Expand All @@ -217,11 +217,11 @@ CONDITIONAL_FACTORY_ENTRY(
makeCheckNot(CheckArgIsDefaultCudaStream(3)),
ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY("cudaWaitExternalSemaphoresAsync_v2",
CALL(MapNames::getDpctNamespace() +
"experimental::wait_external_semaphore",
"experimental::wait_external_semaphores",
ARG(0), ARG(1), ARG(2), ARG(3)))),
ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY("cudaWaitExternalSemaphoresAsync_v2",
CALL(MapNames::getDpctNamespace() +
"experimental::wait_external_semaphore",
"experimental::wait_external_semaphores",
ARG(0), ARG(1), ARG(2))))),
UNSUPPORT_FACTORY_ENTRY("cudaWaitExternalSemaphoresAsync_v2", Diagnostics::TRY_EXPERIMENTAL_FEATURE,
ARG("cudaWaitExternalSemaphoresAsync_v2"),
Expand Down
78 changes: 48 additions & 30 deletions clang/runtime/dpct-rt/include/dpct/bindless_images.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1307,31 +1307,40 @@ inline void import_external_semaphore(external_sem_wrapper **extSem,
/// parameters.
/// \param [in] numExtSems Number of external semaphores to signal.
/// \param [in] q_ptr The queue used to signal the external semaphore resource.
inline void signal_external_semaphore(external_sem_wrapper **extSem,
external_sem_params *semSignalParams,
unsigned int numExtSems,
queue_ptr q_ptr = &get_default_queue()) {
inline void signal_external_semaphores(external_sem_wrapper **extSem,
external_sem_params *semSignalParams,
unsigned int numExtSems,
queue_ptr q_ptr = &get_default_queue()) {
for (int i = 0; i < numExtSems; i++) {
auto extSemType = extSem[i]->get_handle_type();
switch (extSemType) {
#ifdef _WIN32
if (extSemType == sycl::ext::oneapi::experimental::
external_semaphore_handle_type::win32_nt_dx12_fence) {
case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
Copy link
Contributor

@ziranzha ziranzha Mar 31, 2025 •
edited
Loading

Choose a reason for hiding this comment

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

Can be merge WIN32 together?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Okay, I'll refactor the switch case

win32_nt_dx12_fence:
case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
timeline_win32_nt_handle:
q_ptr->ext_oneapi_signal_external_semaphore(
extSem[i]->get(), semSignalParams[i].get_value());
} else if (extSemType ==
sycl::ext::oneapi::experimental::external_semaphore_handle_type::
win32_nt_handle) {
break;
case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
win32_nt_handle:
q_ptr->ext_oneapi_signal_external_semaphore(extSem[i]->get());
}
#else
if (extSemType == sycl::ext::oneapi::experimental::
external_semaphore_handle_type::opaque_fd) {
break;
#else // _WIN32
case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
timeline_fd:
q_ptr->ext_oneapi_signal_external_semaphore(
extSem[i]->get(), semSignalParams[i].get_value());
break;
case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
opaque_fd:
q_ptr->ext_oneapi_signal_external_semaphore(extSem[i]->get());
}
break;
#endif // _WIN32
else {
default:
throw std::runtime_error(
"Unsupported external semaphore resource handle type!");
break;
}
}
}
Expand All @@ -1341,31 +1350,40 @@ inline void signal_external_semaphore(external_sem_wrapper **extSem,
/// \param [in] semWaitParams Pointer to the external semaphore wait parameters.
/// \param [in] numExtSems Number of external semaphores to wait.
/// \param [in] q_ptr The queue used to wait on the external semaphore resource.
inline void wait_external_semaphore(external_sem_wrapper **extSem,
external_sem_params *semWaitParams,
unsigned int numExtSems,
queue_ptr q_ptr = &get_default_queue()) {
inline void wait_external_semaphores(external_sem_wrapper **extSem,
external_sem_params *semWaitParams,
unsigned int numExtSems,
queue_ptr q_ptr = &get_default_queue()) {
for (int i = 0; i < numExtSems; i++) {
auto extSemType = extSem[i]->get_handle_type();
switch (extSemType) {
#ifdef _WIN32
if (extSemType == sycl::ext::oneapi::experimental::
external_semaphore_handle_type::win32_nt_dx12_fence) {
case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
win32_nt_dx12_fence:
case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
timeline_win32_nt_handle:
q_ptr->ext_oneapi_wait_external_semaphore(extSem[i]->get(),
semWaitParams[i].get_value());
} else if (extSemType ==
sycl::ext::oneapi::experimental::external_semaphore_handle_type::
win32_nt_handle) {
break;
case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
win32_nt_handle:
q_ptr->ext_oneapi_wait_external_semaphore(extSem[i]->get());
}
#else
if (extSemType == sycl::ext::oneapi::experimental::
external_semaphore_handle_type::opaque_fd) {
break;
#else // _WIN32
case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
timeline_fd:
q_ptr->ext_oneapi_wait_external_semaphore(extSem[i]->get(),
semWaitParams[i].get_value());
break;
case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
opaque_fd:
q_ptr->ext_oneapi_wait_external_semaphore(extSem[i]->get());
}
break;
#endif // _WIN32
else {
default:
throw std::runtime_error(
"Unsupported external semaphore resource handle type!");
break;
}
}
}
Expand Down
14 changes: 9 additions & 5 deletions clang/test/dpct/externalResInterop.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -109,10 +109,14 @@ int main() {
memHandleDesc.type = cudaExternalMemoryHandleTypeOpaqueWin32Kmt;
#endif // !NO_BUILD_TEST

// CHECK: semHandleDesc.set_handle_type(sycl::ext::oneapi::experimental::external_semaphore_handle_type::win32_nt_handle);
// CHECK: semHandleDesc.set_handle_type(sycl::ext::oneapi::experimental::external_semaphore_handle_type::timeline_win32_nt_handle);
// CHECK-NEXT: semHandleDesc.set_handle_type(sycl::ext::oneapi::experimental::external_semaphore_handle_type::timeline_fd);
// CHECK-NEXT: semHandleDesc.set_handle_type(sycl::ext::oneapi::experimental::external_semaphore_handle_type::win32_nt_handle);
// CHECK-NEXT: semHandleDesc.set_handle_type(sycl::ext::oneapi::experimental::external_semaphore_handle_type::win32_nt_dx12_fence);
// CHECK-NEXT: semHandleDesc.set_handle_type(sycl::ext::oneapi::experimental::external_semaphore_handle_type::opaque_fd);
// CHECK-NEXT: semHandleDesc.set_handle_type(semHandleType);
semHandleDesc.type = cudaExternalSemaphoreHandleTypeTimelineSemaphoreWin32;
semHandleDesc.type = cudaExternalSemaphoreHandleTypeTimelineSemaphoreFd;
semHandleDesc.type = cudaExternalSemaphoreHandleTypeOpaqueWin32;
semHandleDesc.type = cudaExternalSemaphoreHandleTypeD3D12Fence;
semHandleDesc.type = cudaExternalSemaphoreHandleTypeOpaqueFd;
Expand Down Expand Up @@ -365,10 +369,10 @@ int main() {
// CHECK-WINDOWS-NEXT: DPCT1136:{{[0-9]+}}: SYCL Bindless Images extension only supports importing external resource using NT handle on Windows. If assert(semHandleDesc.get_win32_handle()) fails, you may need to adjust the code to use (semHandleDesc.get_win32_handle()).
// CHECK-WINDOWS-NEXT: */
// CHECK: dpct::experimental::import_external_semaphore(&extSem, &semHandleDesc);
// CHECK-NEXT: dpct::experimental::signal_external_semaphore(&extSem, &signalParams, 1);
// CHECK-NEXT: dpct::experimental::signal_external_semaphore(extSemArr, signalParamsArr, numExtSems, stream);
// CHECK-NEXT: dpct::experimental::wait_external_semaphore(&extSem, &waitParams, 1);
// CHECK-NEXT: dpct::experimental::wait_external_semaphore(extSemArr, waitParamsArr, numExtSems, stream);
// CHECK-NEXT: dpct::experimental::signal_external_semaphores(&extSem, &signalParams, 1);
// CHECK-NEXT: dpct::experimental::signal_external_semaphores(extSemArr, signalParamsArr, numExtSems, stream);
// CHECK-NEXT: dpct::experimental::wait_external_semaphores(&extSem, &waitParams, 1);
// CHECK-NEXT: dpct::experimental::wait_external_semaphores(extSemArr, waitParamsArr, numExtSems, stream);
// CHECK-NEXT: delete extSem;
cudaImportExternalSemaphore(&extSem, &semHandleDesc);
cudaSignalExternalSemaphoresAsync(&extSem, &signalParams, 1);
Expand Down
Loading