Skip to content

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

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

Merged
merged 2 commits into from
Apr 28, 2025
Merged

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

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
8ace255
Added support for 2 more semaphore types
TejaX-Alaghari Mar 28, 2025
d739d0f
Added API names without _v2 for supporting migration in 10.2
TejaX-Alaghari Apr 17, 2025
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 @@ -1608,6 +1608,18 @@ void MapNames::setExplicitNamespaceMap(
? getExpNamespace() +
"external_semaphore_handle_type::win32_nt_dx12_fence"
: "cudaExternalSemaphoreHandleTypeD3D12Fence")},
{"cudaExternalSemaphoreHandleTypeTimelineSemaphoreFd",
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
40 changes: 36 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 @@ -195,33 +195,65 @@ CONDITIONAL_FACTORY_ENTRY(
ARG("cudaImportExternalSemaphore"),
ARG("--use-experimental-features=bindless_images")))

CONDITIONAL_FACTORY_ENTRY(
UseExtBindlessImages,
CONDITIONAL_FACTORY_ENTRY(
makeCheckNot(CheckArgIsDefaultCudaStream(3)),
ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY("cudaSignalExternalSemaphoresAsync",
CALL(MapNames::getDpctNamespace() +
"experimental::signal_external_semaphores",
ARG(0), ARG(1), ARG(2), ARG(3)))),
ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY("cudaSignalExternalSemaphoresAsync",
CALL(MapNames::getDpctNamespace() +
"experimental::signal_external_semaphores",
ARG(0), ARG(1), ARG(2))))),
UNSUPPORT_FACTORY_ENTRY("cudaSignalExternalSemaphoresAsync", Diagnostics::TRY_EXPERIMENTAL_FEATURE,
ARG("cudaSignalExternalSemaphoresAsync"),
ARG("--use-experimental-features=bindless_images")))

CONDITIONAL_FACTORY_ENTRY(
UseExtBindlessImages,
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"),
ARG("--use-experimental-features=bindless_images")))

CONDITIONAL_FACTORY_ENTRY(
UseExtBindlessImages,
CONDITIONAL_FACTORY_ENTRY(
makeCheckNot(CheckArgIsDefaultCudaStream(3)),
ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY("cudaWaitExternalSemaphoresAsync",
CALL(MapNames::getDpctNamespace() +
"experimental::wait_external_semaphores",
ARG(0), ARG(1), ARG(2), ARG(3)))),
ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY("cudaWaitExternalSemaphoresAsync",
CALL(MapNames::getDpctNamespace() +
"experimental::wait_external_semaphores",
ARG(0), ARG(1), ARG(2))))),
UNSUPPORT_FACTORY_ENTRY("cudaWaitExternalSemaphoresAsync", Diagnostics::TRY_EXPERIMENTAL_FEATURE,
ARG("cudaWaitExternalSemaphoresAsync"),
ARG("--use-experimental-features=bindless_images")))

CONDITIONAL_FACTORY_ENTRY(
UseExtBindlessImages,
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
6 changes: 4 additions & 2 deletions clang/lib/DPCT/RulesLang/RulesLangGraphicsInterop.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -61,8 +61,10 @@ void GraphicsInteropRule::registerMatcher(ast_matchers::MatchFinder &MF) {
"cuGraphicsUnregisterResource", "cudaImportExternalMemory",
"cudaExternalMemoryGetMappedMipmappedArray",
"cudaExternalMemoryGetMappedBuffer", "cudaDestroyExternalMemory",
"cudaImportExternalSemaphore", "cudaSignalExternalSemaphoresAsync_v2",
"cudaWaitExternalSemaphoresAsync_v2", "cudaDestroyExternalSemaphore");
"cudaImportExternalSemaphore", "cudaSignalExternalSemaphoresAsync",
"cudaSignalExternalSemaphoresAsync_v2",
"cudaWaitExternalSemaphoresAsync", "cudaWaitExternalSemaphoresAsync_v2",
"cudaDestroyExternalSemaphore");
};
MF.addMatcher(
callExpr(callee(functionDecl(graphicsInteropAPI()))).bind("call"), this);
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 @@ -1343,31 +1343,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::
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 @@ -1377,31 +1386,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