Skip to content

Commit db4d985

Browse files
TejaX-AlaghariTejaX-Alaghari
committed
Added support for 2 more semaphore types
1 parent c1cb3cc commit db4d985

15 files changed

+225
-43
lines changed

Diff for: ‎clang/examples/DPCT/Runtime/cudaDestroyExternalMemory.cu

+7
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,7 @@
1+
// Option: --use-experimental-features=bindless_images
2+
3+
void test(cudaExternalMemory_t extMem) {
4+
// Start
5+
cudaDestroyExternalMemory(extMem /*cudaExternalMemory_t*/);
6+
// End
7+
}

Diff for: ‎clang/examples/DPCT/Runtime/cudaDestroyExternalSemaphore.cu

+7
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,7 @@
1+
// Option: --use-experimental-features=bindless_images
2+
3+
void test(cudaExternalSemaphore_t extSem) {
4+
// Start
5+
cudaDestroyExternalSemaphore(extSem /*cudaExternalSemaphore_t*/);
6+
// End
7+
}

Diff for: ‎clang/examples/DPCT/Runtime/cudaExternalMemoryGetMappedBuffer.cu

+9
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,9 @@
1+
// Option: --use-experimental-features=bindless_images
2+
3+
void test(void **devPtr, cudaExternalMemory_t extMem, const cudaExternalMemoryBufferDesc *bufferDesc) {
4+
// Start
5+
cudaExternalMemoryGetMappedBuffer(devPtr /*void ***/,
6+
extMem /*cudaExternalMemory_t*/,
7+
bufferDesc /*const cudaExternalMemoryBufferDesc **/);
8+
// End
9+
}

Diff for: ‎clang/examples/DPCT/Runtime/cudaExternalMemoryGetMappedMipmappedArray.cu

+9
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,9 @@
1+
// Option: --use-experimental-features=bindless_images
2+
3+
void test(cudaMipmappedArray_t *mipmap, cudaExternalMemory_t extMem, const cudaExternalMemoryMipmappedArrayDesc *mipmapDesc) {
4+
// Start
5+
cudaExternalMemoryGetMappedMipmappedArray(mipmap /*cudaMipmappedArray_t **/,
6+
extMem /*cudaExternalMemory_t*/,
7+
mipmapDesc /*const cudaExternalMemoryMipmappedArrayDesc **/);
8+
// End
9+
}

Diff for: ‎clang/examples/DPCT/Runtime/cudaImportExternalMemory.cu

+8
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,8 @@
1+
// Option: --use-experimental-features=bindless_images
2+
3+
void test(cudaExternalMemory_t *extMem, const cudaExternalMemoryHandleDesc *memHandleDesc) {
4+
// Start
5+
cudaImportExternalMemory(extMem /*cudaExternalMemory_t **/,
6+
memHandleDesc /*const cudaExternalMemoryHandleDesc **/);
7+
// End
8+
}

Diff for: ‎clang/examples/DPCT/Runtime/cudaImportExternalSemaphore.cu

+8
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,8 @@
1+
// Option: --use-experimental-features=bindless_images
2+
3+
void test(cudaExternalSemaphore_t *extSem, const cudaExternalSemaphoreHandleDesc *semHandleDesc) {
4+
// Start
5+
cudaImportExternalSemaphore(extSem /*cudaExternalSemaphore_t **/,
6+
semHandleDesc /*const cudaExternalSemaphoreHandleDesc **/);
7+
// End
8+
}

Diff for: ‎clang/examples/DPCT/Runtime/cudaSignalExternalSemaphoresAsync.cu

+13
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,13 @@
1+
// Option: --use-experimental-features=bindless_images
2+
3+
void test(const cudaExternalSemaphore_t *extSemArray, const cudaExternalSemaphoreSignalParams *paramsArray, unsigned int numExtSems, cudaStream_t stream = 0) {
4+
// Start
5+
cudaSignalExternalSemaphoresAsync(extSemArray /*const cudaExternalSemaphore_t **/,
6+
paramsArray /*const cudaExternalSemaphoreSignalParams **/,
7+
numExtSems /*unsigned int*/);
8+
cudaSignalExternalSemaphoresAsync(extSemArray /*const cudaExternalSemaphore_t **/,
9+
paramsArray /*const cudaExternalSemaphoreSignalParams **/,
10+
numExtSems /*unsigned int*/,
11+
stream /*cudaStream_t*/);
12+
// End
13+
}

Diff for: ‎clang/examples/DPCT/Runtime/cudaWaitExternalSemaphoresAsync.cu

+13
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,13 @@
1+
// Option: --use-experimental-features=bindless_images
2+
3+
void test(const cudaExternalSemaphore_t *extSemArray, const cudaExternalSemaphoreWaitParams *paramsArray, unsigned int numExtSems, cudaStream_t stream = 0) {
4+
// Start
5+
cudaWaitExternalSemaphoresAsync(extSemArray /*const cudaExternalSemaphore_t **/,
6+
paramsArray /*const cudaExternalSemaphoreWaitParams **/,
7+
numExtSems /*unsigned int*/);
8+
cudaWaitExternalSemaphoresAsync(extSemArray /*const cudaExternalSemaphore_t **/,
9+
paramsArray /*const cudaExternalSemaphoreWaitParams **/,
10+
numExtSems /*unsigned int*/,
11+
stream /*cudaStream_t*/);
12+
// End
13+
}

Diff for: ‎clang/lib/DPCT/RuleInfra/MapNames.cpp

+12
Original file line numberDiff line numberDiff line change
@@ -1550,6 +1550,18 @@ void MapNames::setExplicitNamespaceMap(
15501550
? getExpNamespace() +
15511551
"external_semaphore_handle_type::win32_nt_dx12_fence"
15521552
: "cudaExternalSemaphoreHandleTypeD3D12Fence")},
1553+
{"cudaExternalSemaphoreHandleTypeTimelineSemaphoreFd",
1554+
std::make_shared<EnumNameRule>(
1555+
DpctGlobalInfo::useExtBindlessImages()
1556+
? getExpNamespace() +
1557+
"external_semaphore_handle_type::timeline_fd"
1558+
: "cudaExternalSemaphoreHandleTypeTimelineSemaphoreFd")},
1559+
{"cudaExternalSemaphoreHandleTypeTimelineSemaphoreWin32",
1560+
std::make_shared<EnumNameRule>(
1561+
DpctGlobalInfo::useExtBindlessImages()
1562+
? getExpNamespace() +
1563+
"external_semaphore_handle_type::timeline_win32_nt_handle"
1564+
: "cudaExternalSemaphoreHandleTypeTimelineSemaphoreWin32")},
15531565
// ...
15541566
};
15551567

Diff for: ‎clang/lib/DPCT/RulesLang/APINamesGraphicsInterop.inc

+4-4
Original file line numberDiff line numberDiff line change
@@ -201,11 +201,11 @@ CONDITIONAL_FACTORY_ENTRY(
201201
makeCheckNot(CheckArgIsDefaultCudaStream(3)),
202202
ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY("cudaSignalExternalSemaphoresAsync_v2",
203203
CALL(MapNames::getDpctNamespace() +
204-
"experimental::signal_external_semaphore",
204+
"experimental::signal_external_semaphores",
205205
ARG(0), ARG(1), ARG(2), ARG(3)))),
206206
ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY("cudaSignalExternalSemaphoresAsync_v2",
207207
CALL(MapNames::getDpctNamespace() +
208-
"experimental::signal_external_semaphore",
208+
"experimental::signal_external_semaphores",
209209
ARG(0), ARG(1), ARG(2))))),
210210
UNSUPPORT_FACTORY_ENTRY("cudaSignalExternalSemaphoresAsync_v2", Diagnostics::TRY_EXPERIMENTAL_FEATURE,
211211
ARG("cudaSignalExternalSemaphoresAsync_v2"),
@@ -217,11 +217,11 @@ CONDITIONAL_FACTORY_ENTRY(
217217
makeCheckNot(CheckArgIsDefaultCudaStream(3)),
218218
ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY("cudaWaitExternalSemaphoresAsync_v2",
219219
CALL(MapNames::getDpctNamespace() +
220-
"experimental::wait_external_semaphore",
220+
"experimental::wait_external_semaphores",
221221
ARG(0), ARG(1), ARG(2), ARG(3)))),
222222
ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY("cudaWaitExternalSemaphoresAsync_v2",
223223
CALL(MapNames::getDpctNamespace() +
224-
"experimental::wait_external_semaphore",
224+
"experimental::wait_external_semaphores",
225225
ARG(0), ARG(1), ARG(2))))),
226226
UNSUPPORT_FACTORY_ENTRY("cudaWaitExternalSemaphoresAsync_v2", Diagnostics::TRY_EXPERIMENTAL_FEATURE,
227227
ARG("cudaWaitExternalSemaphoresAsync_v2"),

Diff for: ‎clang/runtime/dpct-rt/include/dpct/bindless_images.hpp

+45-31
Original file line numberDiff line numberDiff line change
@@ -1290,31 +1290,38 @@ inline void import_external_semaphore(external_sem_wrapper **extSem,
12901290
/// parameters.
12911291
/// \param [in] numExtSems Number of external semaphores to signal.
12921292
/// \param [in] q_ptr The queue used to signal the external semaphore resource.
1293-
inline void signal_external_semaphore(external_sem_wrapper **extSem,
1293+
inline void signal_external_semaphores(external_sem_wrapper **extSem,
12941294
external_sem_params *semSignalParams,
12951295
unsigned int numExtSems,
12961296
queue_ptr q_ptr = &get_default_queue()) {
12971297
for (int i = 0; i < numExtSems; i++) {
12981298
auto extSemType = extSem[i]->get_handle_type();
1299+
switch (extSemType) {
12991300
#ifdef _WIN32
1300-
if (extSemType == sycl::ext::oneapi::experimental::
1301-
external_semaphore_handle_type::win32_nt_dx12_fence) {
1302-
q_ptr->ext_oneapi_signal_external_semaphore(
1301+
case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
1302+
win32_nt_dx12_fence:
1303+
case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
1304+
timeline_win32_nt_handle:
1305+
#else // _WIN32
1306+
case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
1307+
timeline_fd:
1308+
#endif // _WIN32
1309+
q_ptr->ext_oneapi_signal_external_semaphores(
13031310
extSem[i]->get(), semSignalParams[i].get_value());
1304-
} else if (extSemType ==
1305-
sycl::ext::oneapi::experimental::external_semaphore_handle_type::
1306-
win32_nt_handle) {
1307-
q_ptr->ext_oneapi_signal_external_semaphore(extSem[i]->get());
1308-
}
1309-
#else
1310-
if (extSemType == sycl::ext::oneapi::experimental::
1311-
external_semaphore_handle_type::opaque_fd) {
1312-
q_ptr->ext_oneapi_signal_external_semaphore(extSem[i]->get());
1313-
}
1311+
break;
1312+
#ifdef _WIN32
1313+
case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
1314+
win32_nt_handle:
1315+
#else // _WIN32
1316+
case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
1317+
opaque_fd:
13141318
#endif // _WIN32
1315-
else {
1319+
q_ptr->ext_oneapi_signal_external_semaphores(extSem[i]->get());
1320+
break;
1321+
default:
13161322
throw std::runtime_error(
13171323
"Unsupported external semaphore resource handle type!");
1324+
break;
13181325
}
13191326
}
13201327
}
@@ -1324,31 +1331,38 @@ inline void signal_external_semaphore(external_sem_wrapper **extSem,
13241331
/// \param [in] semWaitParams Pointer to the external semaphore wait parameters.
13251332
/// \param [in] numExtSems Number of external semaphores to wait.
13261333
/// \param [in] q_ptr The queue used to wait on the external semaphore resource.
1327-
inline void wait_external_semaphore(external_sem_wrapper **extSem,
1334+
inline void wait_external_semaphores(external_sem_wrapper **extSem,
13281335
external_sem_params *semWaitParams,
13291336
unsigned int numExtSems,
13301337
queue_ptr q_ptr = &get_default_queue()) {
13311338
for (int i = 0; i < numExtSems; i++) {
13321339
auto extSemType = extSem[i]->get_handle_type();
1340+
switch (extSemType) {
13331341
#ifdef _WIN32
1334-
if (extSemType == sycl::ext::oneapi::experimental::
1335-
external_semaphore_handle_type::win32_nt_dx12_fence) {
1336-
q_ptr->ext_oneapi_wait_external_semaphore(extSem[i]->get(),
1337-
semWaitParams[i].get_value());
1338-
} else if (extSemType ==
1339-
sycl::ext::oneapi::experimental::external_semaphore_handle_type::
1340-
win32_nt_handle) {
1341-
q_ptr->ext_oneapi_wait_external_semaphore(extSem[i]->get());
1342-
}
1343-
#else
1344-
if (extSemType == sycl::ext::oneapi::experimental::
1345-
external_semaphore_handle_type::opaque_fd) {
1346-
q_ptr->ext_oneapi_wait_external_semaphore(extSem[i]->get());
1347-
}
1342+
case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
1343+
win32_nt_dx12_fence:
1344+
case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
1345+
timeline_win32_nt_handle:
1346+
#else // _WIN32
1347+
case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
1348+
timeline_fd:
13481349
#endif // _WIN32
1349-
else {
1350+
q_ptr->ext_oneapi_wait_external_semaphores(
1351+
extSem[i]->get(), semWaitParams[i].get_value());
1352+
break;
1353+
#ifdef _WIN32
1354+
case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
1355+
win32_nt_handle:
1356+
#else // _WIN32
1357+
case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
1358+
opaque_fd:
1359+
#endif // _WIN32
1360+
q_ptr->ext_oneapi_wait_external_semaphores(extSem[i]->get());
1361+
break;
1362+
default:
13501363
throw std::runtime_error(
13511364
"Unsupported external semaphore resource handle type!");
1365+
break;
13521366
}
13531367
}
13541368
}

Diff for: ‎clang/test/dpct/externalResInterop.cu

+9-5
Original file line numberDiff line numberDiff line change
@@ -109,10 +109,14 @@ int main() {
109109
memHandleDesc.type = cudaExternalMemoryHandleTypeOpaqueWin32Kmt;
110110
#endif // !NO_BUILD_TEST
111111

112-
// CHECK: semHandleDesc.set_handle_type(sycl::ext::oneapi::experimental::external_semaphore_handle_type::win32_nt_handle);
112+
// CHECK: semHandleDesc.set_handle_type(sycl::ext::oneapi::experimental::external_semaphore_handle_type::timeline_win32_nt_handle);
113+
// CHECK-NEXT: semHandleDesc.set_handle_type(sycl::ext::oneapi::experimental::external_semaphore_handle_type::timeline_fd);
114+
// CHECK-NEXT: semHandleDesc.set_handle_type(sycl::ext::oneapi::experimental::external_semaphore_handle_type::win32_nt_handle);
113115
// CHECK-NEXT: semHandleDesc.set_handle_type(sycl::ext::oneapi::experimental::external_semaphore_handle_type::win32_nt_dx12_fence);
114116
// CHECK-NEXT: semHandleDesc.set_handle_type(sycl::ext::oneapi::experimental::external_semaphore_handle_type::opaque_fd);
115117
// CHECK-NEXT: semHandleDesc.set_handle_type(semHandleType);
118+
semHandleDesc.type = cudaExternalSemaphoreHandleTypeTimelineSemaphoreWin32;
119+
semHandleDesc.type = cudaExternalSemaphoreHandleTypeTimelineSemaphoreFd;
116120
semHandleDesc.type = cudaExternalSemaphoreHandleTypeOpaqueWin32;
117121
semHandleDesc.type = cudaExternalSemaphoreHandleTypeD3D12Fence;
118122
semHandleDesc.type = cudaExternalSemaphoreHandleTypeOpaqueFd;
@@ -365,10 +369,10 @@ int main() {
365369
// 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()).
366370
// CHECK-WINDOWS-NEXT: */
367371
// CHECK: dpct::experimental::import_external_semaphore(&extSem, &semHandleDesc);
368-
// CHECK-NEXT: dpct::experimental::signal_external_semaphore(&extSem, &signalParams, 1);
369-
// CHECK-NEXT: dpct::experimental::signal_external_semaphore(extSemArr, signalParamsArr, numExtSems, stream);
370-
// CHECK-NEXT: dpct::experimental::wait_external_semaphore(&extSem, &waitParams, 1);
371-
// CHECK-NEXT: dpct::experimental::wait_external_semaphore(extSemArr, waitParamsArr, numExtSems, stream);
372+
// CHECK-NEXT: dpct::experimental::signal_external_semaphores(&extSem, &signalParams, 1);
373+
// CHECK-NEXT: dpct::experimental::signal_external_semaphores(extSemArr, signalParamsArr, numExtSems, stream);
374+
// CHECK-NEXT: dpct::experimental::wait_external_semaphores(&extSem, &waitParams, 1);
375+
// CHECK-NEXT: dpct::experimental::wait_external_semaphores(extSemArr, waitParamsArr, numExtSems, stream);
372376
// CHECK-NEXT: delete extSem;
373377
cudaImportExternalSemaphore(&extSem, &semHandleDesc);
374378
cudaSignalExternalSemaphoresAsync(&extSem, &signalParams, 1);

0 commit comments

Comments
 (0)