Skip to content

Commit f119a76

Browse files
TejaX-AlaghariTejaX-Alaghari
authored
[SYCLomatic] Added support 4 CUDA Driver APIs used in Blender (#2774)
1 parent 9af6b57 commit f119a76

File tree

9 files changed

+246
-26
lines changed

9 files changed

+246
-26
lines changed

‎clang/lib/DPCT/RuleInfra/APINames_removed.inc

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -5,6 +5,7 @@
55
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
66
//
77
//===----------------------------------------------------------------------===//
8+
// clang-format off
89
ENTRY(cublasInit, "this functionality is redundant in SYCL.")
910
ENTRY(cublasShutdown, "this functionality is redundant in SYCL.")
1011
ENTRY(cublasGetError, "this functionality is redundant in SYCL.")
@@ -42,4 +43,8 @@ ENTRY(cudaProfilerStop, "SYCL currently does not support this fun
4243
ENTRY(cuFuncSetAttribute, "SYCL currently does not support setting kernel function attributes")
4344
ENTRY(cuGetExportTable, "SYCL does not provide a standard API to export internal runtime or driver API. Check and implement the functionality corresponding to the function of the first parameter `const void **table` populated by the API.")
4445

46+
ENTRY(cuDevicePrimaryCtxSetFlags_v2, "SYCL currently does not support setting device context flags.")
47+
ENTRY(cuDevicePrimaryCtxGetState, "SYCL currently does not support querying device context flags.")
48+
4549
ENTRY(cudaGraphicsResourceSetMapFlags, "this functionality is deprecated in DX12 and hence is not supported in SYCL.")
50+
// clang-format on

‎clang/lib/DPCT/RuleInfra/MapNames.cpp

Lines changed: 37 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -713,6 +713,11 @@ void MapNames::setExplicitNamespaceMap(
713713
getLibraryHelperNamespace() + "fft::fft_engine_ptr",
714714
HelperFeatureEnum::device_ext)},
715715
{"CUdevice", std::make_shared<TypeNameRule>("int")},
716+
{"CUdevice_P2PAttribute",
717+
std::make_shared<TypeNameRule>(DpctGlobalInfo::usePeerAccess()
718+
? getClNamespace() +
719+
"ext::oneapi::peer_access"
720+
: "CUdevice_P2PAttribute")},
716721
{"CUarray_st",
717722
std::make_shared<TypeNameRule>(
718723
DpctGlobalInfo::useExtBindlessImages()
@@ -1243,10 +1248,42 @@ void MapNames::setExplicitNamespaceMap(
12431248
{"CU_DEVICE_ATTRIBUTE_MAX_PITCH",
12441249
std::make_shared<EnumNameRule>("get_max_pitch",
12451250
HelperFeatureEnum::device_ext)},
1251+
{"CU_DEVICE_P2P_ATTRIBUTE_ACCESS_SUPPORTED",
1252+
std::make_shared<EnumNameRule>(
1253+
DpctGlobalInfo::usePeerAccess()
1254+
? getClNamespace() + "ext::oneapi::peer_access::access_supported"
1255+
: "CU_DEVICE_P2P_ATTRIBUTE_ACCESS_SUPPORTED")},
1256+
{"CU_DEVICE_P2P_ATTRIBUTE_CUDA_ARRAY_ACCESS_SUPPORTED",
1257+
std::make_shared<EnumNameRule>(
1258+
DpctGlobalInfo::usePeerAccess()
1259+
? getClNamespace() + "ext::oneapi::peer_access::access_supported"
1260+
: "CU_DEVICE_P2P_ATTRIBUTE_CUDA_ARRAY_ACCESS_SUPPORTED")},
1261+
{"CU_DEVICE_P2P_ATTRIBUTE_ACCESS_ACCESS_SUPPORTED",
1262+
std::make_shared<EnumNameRule>(
1263+
DpctGlobalInfo::usePeerAccess()
1264+
? getClNamespace() + "ext::oneapi::peer_access::access_supported"
1265+
: "CU_DEVICE_P2P_ATTRIBUTE_ACCESS_ACCESS_SUPPORTED")},
1266+
{"CU_DEVICE_P2P_ATTRIBUTE_ARRAY_ACCESS_ACCESS_SUPPORTED",
1267+
std::make_shared<EnumNameRule>(
1268+
DpctGlobalInfo::usePeerAccess()
1269+
? getClNamespace() + "ext::oneapi::peer_access::access_supported"
1270+
: "CU_DEVICE_P2P_ATTRIBUTE_ARRAY_ACCESS_ACCESS_SUPPORTED")},
1271+
{"CU_DEVICE_P2P_ATTRIBUTE_NATIVE_ATOMIC_SUPPORTED",
1272+
std::make_shared<EnumNameRule>(
1273+
DpctGlobalInfo::usePeerAccess()
1274+
? getClNamespace() +
1275+
"ext::oneapi::peer_access::atomics_supported"
1276+
: "CU_DEVICE_P2P_ATTRIBUTE_NATIVE_ATOMIC_SUPPORTED")},
1277+
{"CU_CTX_BLOCKING_SYNC", std::make_shared<EnumNameRule>("0")},
1278+
{"CU_CTX_COREDUMP_ENABLE", std::make_shared<EnumNameRule>("0")},
12461279
{"CU_CTX_LMEM_RESIZE_TO_MAX", std::make_shared<EnumNameRule>("0")},
12471280
{"CU_CTX_MAP_HOST", std::make_shared<EnumNameRule>("0")},
1281+
{"CU_CTX_SCHED_AUTO", std::make_shared<EnumNameRule>("0")},
12481282
{"CU_CTX_SCHED_BLOCKING_SYNC", std::make_shared<EnumNameRule>("0")},
12491283
{"CU_CTX_SCHED_SPIN", std::make_shared<EnumNameRule>("0")},
1284+
{"CU_CTX_SCHED_YIELD", std::make_shared<EnumNameRule>("0")},
1285+
{"CU_CTX_SYNC_MEMOPS", std::make_shared<EnumNameRule>("0")},
1286+
{"CU_CTX_USER_COREDUMP_ENABLE", std::make_shared<EnumNameRule>("0")},
12501287
{"CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK",
12511288
std::make_shared<EnumNameRule>("get_device_info().get_local_mem_size",
12521289
HelperFeatureEnum::device_ext)},

‎clang/lib/DPCT/RulesLang/APINamesMemory.inc

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -875,6 +875,21 @@ CONDITIONAL_FACTORY_ENTRY(
875875
ARG("cuMemSetAccess"),
876876
ARG("--use-experimental-features=virtual_mem")))
877877

878+
CONDITIONAL_FACTORY_ENTRY(
879+
UsePeerAccess(),
880+
ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(
881+
"cuDeviceGetP2PAttribute", DEREF(ARG_WC(0)),
882+
MEMBER_CALL(
883+
CALL(MapNames::getDpctNamespace() + "get_device", ARG_WC(2)), false,
884+
"ext_oneapi_can_access_peer",
885+
CALL(MapNames::getDpctNamespace() + "get_device", ARG_WC(3)),
886+
ARG_WC(1)))),
887+
ASSIGNABLE_FACTORY(
888+
WARNING_FACTORY_ENTRY("cuDeviceGetP2PAttribute",
889+
ASSIGN_FACTORY_ENTRY("cuDeviceGetP2PAttribute",
890+
DEREF(ARG_WC(0)), ARG("0")),
891+
Diagnostics::EXPLICIT_PEER_ACCESS)))
892+
878893
CONDITIONAL_FACTORY_ENTRY(
879894
UsePeerAccess(),
880895
ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(

‎clang/lib/DPCT/RulesLang/RulesLang.cpp

Lines changed: 73 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -347,8 +347,8 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
347347
"cublasLtMatrixTransformDesc_t", "cudaGraphicsMapFlags",
348348
"cudaGraphicsRegisterFlags", "cudaExternalMemoryHandleType",
349349
"cudaExternalSemaphoreHandleType", "CUstreamCallback",
350-
"cudaHostFn_t", "__nv_half2", "__nv_half",
351-
"cudaGraphNodeType", "CUsurfref"))))))
350+
"cudaHostFn_t", "__nv_half2", "__nv_half", "cudaGraphNodeType",
351+
"CUsurfref", "CUdevice_P2PAttribute"))))))
352352
.bind("cudaTypeDef"),
353353
this);
354354

@@ -944,6 +944,14 @@ void TypeInDeclRule::runRule(const MatchFinder::MatchResult &Result) {
944944
}
945945
}
946946

947+
if (CanonicalTypeStr == "CUdevice_P2PAttribute") {
948+
if (!DpctGlobalInfo::usePeerAccess()) {
949+
report(TL->getBeginLoc(), Diagnostics::API_NOT_MIGRATED, false,
950+
CanonicalTypeStr);
951+
return;
952+
}
953+
}
954+
947955
if (CanonicalTypeStr == "cooperative_groups::__v1::thread_group" ||
948956
CanonicalTypeStr == "cooperative_groups::__v1::thread_block") {
949957
if (auto ETL = TL->getUnqualifiedLoc().getAs<ElaboratedTypeLoc>()) {
@@ -1926,7 +1934,7 @@ void EnumConstantRule::registerMatcher(MatchFinder &MF) {
19261934
"cufftType", "cudaMemoryType", "CUctx_flags_enum",
19271935
"CUpointer_attribute_enum", "CUmemorytype_enum",
19281936
"cudaGraphicsMapFlags", "cudaGraphicsRegisterFlags",
1929-
"cudaGraphNodeType"))),
1937+
"cudaGraphNodeType", "CUdevice_P2PAttribute_enum"))),
19301938
matchesName("CUDNN_.*"), matchesName("CUSOLVER_.*")))))
19311939
.bind("EnumConstant"),
19321940
this);
@@ -2006,7 +2014,8 @@ void EnumConstantRule::runRule(const MatchFinder::MatchResult &Result) {
20062014
EnumName == "cudaGraphNodeTypeExtSemaphoreWait" ||
20072015
EnumName == "cudaGraphNodeTypeMemAlloc" ||
20082016
EnumName == "cudaGraphNodeTypeMemFree" ||
2009-
EnumName == "cudaGraphNodeTypeConditional") {
2017+
EnumName == "cudaGraphNodeTypeConditional" ||
2018+
EnumName == "CU_DEVICE_P2P_ATTRIBUTE_PERFORMANCE_RANK") {
20102019
report(E->getBeginLoc(), Diagnostics::API_NOT_MIGRATED, false, EnumName);
20112020
return;
20122021
} else if (EnumName == "cudaComputeModeDefault" ||
@@ -5945,7 +5954,8 @@ void MemoryMigrationRule::memcpyMigration(
59455954
handleAsync(C, 7, Result);
59465955
} else if (NameRef.rfind("cudaMemcpy3D", 0) == 0 ||
59475956
NameRef.rfind("cuMemcpy3D", 0) == 0 ||
5948-
NameRef.rfind("cuMemcpy2D", 0) == 0) {
5957+
NameRef.rfind("cuMemcpy2D", 0) == 0 ||
5958+
NameRef.rfind("cuMemcpy2DUnaligned", 0) == 0) {
59495959
handleAsync(C, 1, Result);
59505960
std::string Replacement;
59515961
llvm::raw_string_ostream OS(Replacement);
@@ -6770,14 +6780,14 @@ void MemoryMigrationRule::registerMatcher(MatchFinder &MF) {
67706780
"cudaGetChannelDesc", "cuMemHostAlloc", "cuMemFreeHost",
67716781
"cuMemGetInfo_v2", "cuMemAlloc_v2", "cuMemcpyHtoD_v2",
67726782
"cuMemcpyDtoH_v2", "cuMemcpyHtoDAsync_v2", "cuMemcpyDtoHAsync_v2",
6773-
"cuMemcpy2D_v2", "cuMemcpy2DAsync_v2", "cuMemcpy3D_v2",
6774-
"cuMemcpy3DAsync_v2", "cuMemcpy3DPeer", "cuMemcpy3DPeerAsync",
6775-
"cudaMemGetInfo", "cuMemAllocManaged", "cuMemAllocHost_v2",
6776-
"cuMemHostGetDevicePointer_v2", "cuMemcpyDtoDAsync_v2",
6777-
"cuMemcpyDtoD_v2", "cuMemAllocPitch_v2", "cuMemPrefetchAsync",
6778-
"cuMemFree_v2", "cuDeviceTotalMem_v2", "cuMemHostGetFlags",
6779-
"cuMemHostRegister_v2", "cuMemHostUnregister", "cuMemcpy",
6780-
"cuMemcpyAsync", "cuMemcpyHtoA_v2", "cuMemcpyAtoH_v2",
6783+
"cuMemcpy2D_v2", "cuMemcpy2DAsync_v2", "cuMemcpy2DUnaligned_v2",
6784+
"cuMemcpy3D_v2", "cuMemcpy3DAsync_v2", "cuMemcpy3DPeer",
6785+
"cuMemcpy3DPeerAsync", "cudaMemGetInfo", "cuMemAllocManaged",
6786+
"cuMemAllocHost_v2", "cuMemHostGetDevicePointer_v2",
6787+
"cuMemcpyDtoDAsync_v2", "cuMemcpyDtoD_v2", "cuMemAllocPitch_v2",
6788+
"cuMemPrefetchAsync", "cuMemFree_v2", "cuDeviceTotalMem_v2",
6789+
"cuMemHostGetFlags", "cuMemHostRegister_v2", "cuMemHostUnregister",
6790+
"cuMemcpy", "cuMemcpyAsync", "cuMemcpyHtoA_v2", "cuMemcpyAtoH_v2",
67816791
"cuMemcpyHtoAAsync_v2", "cuMemcpyAtoHAsync_v2", "cuMemcpyDtoA_v2",
67826792
"cuMemcpyAtoD_v2", "cuMemcpyAtoA_v2", "cuMemsetD16_v2",
67836793
"cuMemsetD16Async", "cuMemsetD2D16_v2", "cuMemsetD2D16Async",
@@ -6868,7 +6878,6 @@ void MemoryMigrationRule::runRule(const MatchFinder::MatchResult &Result) {
68686878
Name.compare("cuMemcpyDtoD_v2") && Name.compare("cuMemAdvise") &&
68696879
Name.compare("cuMemPrefetchAsync") &&
68706880
Name.compare("cuMemcpyHtoDAsync_v2") &&
6871-
Name.compare("cuMemcpyDtoD_v2") &&
68726881
Name.compare("cuMemHostUnregister") &&
68736882
Name.compare("cuMemHostRegister_v2") &&
68746883
Name.compare("cudaHostGetFlags") && Name.compare("cuMemHostGetFlags") &&
@@ -6968,6 +6977,7 @@ MemoryMigrationRule::MemoryMigrationRule() {
69686977
&MemoryMigrationRule::memcpySymbolMigration},
69696978
{"cudaMemcpy2D", &MemoryMigrationRule::memcpyMigration},
69706979
{"cuMemcpy2D_v2", &MemoryMigrationRule::memcpyMigration},
6980+
{"cuMemcpy2DUnaligned_v2", &MemoryMigrationRule::memcpyMigration},
69716981
{"cuMemcpy2DAsync_v2", &MemoryMigrationRule::memcpyMigration},
69726982
{"cudaMemcpy3D", &MemoryMigrationRule::memcpyMigration},
69736983
{"cudaMemcpy3DPeer", &MemoryMigrationRule::memcpyMigration},
@@ -8344,11 +8354,11 @@ void DriverModuleAPIRule::runRule(
83448354
void DriverDeviceAPIRule::registerMatcher(ast_matchers::MatchFinder &MF) {
83458355

83468356
auto DriverDeviceAPI = [&]() {
8347-
return hasAnyName("cuDeviceGet", "cuDeviceComputeCapability",
8348-
"cuDriverGetVersion", "cuDeviceGetCount",
8349-
"cuDeviceGetAttribute", "cuDeviceGetName",
8350-
"cuDeviceGetUuid", "cuDeviceGetUuid_v2",
8351-
"cuGetErrorString", "cuGetErrorName");
8357+
return hasAnyName(
8358+
"cuDeviceGet", "cuDeviceComputeCapability", "cuDriverGetVersion",
8359+
"cuDeviceGetCount", "cuDeviceGetAttribute", "cuDeviceGetName",
8360+
"cuDeviceGetUuid", "cuDeviceGetUuid_v2", "cuGetErrorString",
8361+
"cuGetErrorName", "cuDeviceGetP2PAttribute");
83528362
};
83538363

83548364
MF.addMatcher(
@@ -8524,6 +8534,39 @@ void DriverDeviceAPIRule::runRule(
85248534
false);
85258535
return;
85268536
}
8537+
} else if (APIName == "cuDeviceGetP2PAttribute") {
8538+
if (!DpctGlobalInfo::usePeerAccess()) {
8539+
report(CE->getBeginLoc(), Diagnostics::API_NOT_MIGRATED, false, APIName);
8540+
return;
8541+
}
8542+
auto SecArg = CE->getArg(1);
8543+
if (auto DRE = dyn_cast<DeclRefExpr>(SecArg->IgnoreImpCasts())) {
8544+
std::string warningMessage = "";
8545+
std::string AttributeName = "";
8546+
8547+
if (const VarDecl *VD = dyn_cast<VarDecl>(DRE->getDecl())) {
8548+
if (VD->hasInit()) {
8549+
// get the attribute name from definition
8550+
if (auto Init = dyn_cast<DeclRefExpr>(VD->getInit())) {
8551+
AttributeName = Init->getNameInfo().getName().getAsString();
8552+
warningMessage = "initialized value of ";
8553+
}
8554+
}
8555+
} else {
8556+
AttributeName = DRE->getNameInfo().getAsString();
8557+
warningMessage = "parameter ";
8558+
}
8559+
8560+
if (!AttributeName.empty()) {
8561+
auto Search = MapNames::EnumNamesMap.find(AttributeName);
8562+
if (Search == MapNames::EnumNamesMap.end()) {
8563+
report(CE->getBeginLoc(), Diagnostics::NOT_SUPPORTED_PARAMETER, false,
8564+
APIName,
8565+
warningMessage + getStmtSpelling(SecArg) + " is unsupported");
8566+
return;
8567+
}
8568+
}
8569+
}
85278570
}
85288571
auto Itr = CallExprRewriterFactoryBase::RewriterMap->find(APIName);
85298572
if (Itr != CallExprRewriterFactoryBase::RewriterMap->end()) {
@@ -8543,7 +8586,9 @@ void DriverContextAPIRule::registerMatcher(ast_matchers::MatchFinder &MF) {
85438586
"cuCtxDestroy_v2", "cuDevicePrimaryCtxRetain",
85448587
"cuDevicePrimaryCtxRelease_v2", "cuDevicePrimaryCtxRelease",
85458588
"cuCtxGetDevice", "cuCtxGetApiVersion", "cuCtxGetLimit",
8546-
"cuCtxPushCurrent_v2", "cuCtxPopCurrent_v2");
8589+
"cuCtxPushCurrent_v2", "cuCtxPopCurrent_v2",
8590+
"cuDevicePrimaryCtxSetFlags", "cuDevicePrimaryCtxSetFlags_v2",
8591+
"cuDevicePrimaryCtxGetState");
85478592
};
85488593

85498594
MF.addMatcher(
@@ -8604,7 +8649,10 @@ void DriverContextAPIRule::runRule(
86048649
return;
86058650
} else if (APIName == "cuCtxDestroy_v2" ||
86068651
APIName == "cuDevicePrimaryCtxRelease_v2" ||
8607-
APIName == "cuDevicePrimaryCtxRelease") {
8652+
APIName == "cuDevicePrimaryCtxRelease" ||
8653+
APIName == "cuDevicePrimaryCtxSetFlags_v2" ||
8654+
APIName == "cuDevicePrimaryCtxSetFlags" ||
8655+
APIName == "cuDevicePrimaryCtxGetState") {
86088656
SourceLocation CallBegin(CE->getBeginLoc());
86098657
SourceLocation CallEnd(CE->getEndLoc());
86108658

@@ -8627,6 +8675,10 @@ void DriverContextAPIRule::runRule(
86278675
CallEnd = CallEnd.getLocWithOffset(1);
86288676

86298677
std::string Msg = "this functionality is redundant in SYCL.";
8678+
if (auto WarnMsg = MapNames::RemovedAPIWarningMessage.find(APIName);
8679+
WarnMsg != MapNames::RemovedAPIWarningMessage.end()) {
8680+
Msg = WarnMsg->second;
8681+
}
86308682
if (IsAssigned) {
86318683
report(CE->getBeginLoc(), Diagnostics::FUNC_CALL_REMOVED_0, false,
86328684
APIName, Msg);

‎clang/lib/DPCT/SrcAPI/APINames.inc

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1567,11 +1567,11 @@ ENTRY(cuDeviceComputeCapability, cuDeviceComputeCapability, true, NO_FLAG, P4, "
15671567
ENTRY(cuDeviceGetProperties, cuDeviceGetProperties, false, NO_FLAG, P4, "comment")
15681568

15691569
// Primary Context Management
1570-
ENTRY(cuDevicePrimaryCtxGetState, cuDevicePrimaryCtxGetState, false, NO_FLAG, P4, "comment")
1570+
ENTRY(cuDevicePrimaryCtxGetState, cuDevicePrimaryCtxGetState, true, NO_FLAG, P4, "DPCT1026/DPCT1027")
15711571
ENTRY(cuDevicePrimaryCtxRelease, cuDevicePrimaryCtxRelease_v2, true, NO_FLAG, P4, "DPCT1026/DPCT1027")
15721572
ENTRY(cuDevicePrimaryCtxReset, cuDevicePrimaryCtxReset_v2, false, NO_FLAG, P4, "comment")
15731573
ENTRY(cuDevicePrimaryCtxRetain, cuDevicePrimaryCtxRetain, true, NO_FLAG, P4, "comment")
1574-
ENTRY(cuDevicePrimaryCtxSetFlags, cuDevicePrimaryCtxSetFlags_v2, false, NO_FLAG, P4, "comment")
1574+
ENTRY(cuDevicePrimaryCtxSetFlags, cuDevicePrimaryCtxSetFlags_v2, true, NO_FLAG, P4, "DPCT1026/DPCT1027")
15751575

15761576
// Context Management
15771577
ENTRY(cuCtxCreate, cuCtxCreate, true, NO_FLAG, P4, "comment")
@@ -1679,7 +1679,7 @@ ENTRY(cuMemHostUnregister, cuMemHostUnregister, true, NO_FLAG, P4, "DPCT1026/DPC
16791679
ENTRY(cuMemcpy, cuMemcpy, true, NO_FLAG, P4, "Successful")
16801680
ENTRY(cuMemcpy2D, cuMemcpy2D_v2, true, NO_FLAG, P4, "Successful")
16811681
ENTRY(cuMemcpy2DAsync, cuMemcpy2DAsync_v2, true, NO_FLAG, P4, "Successful")
1682-
ENTRY(cuMemcpy2DUnaligned, cuMemcpy2DUnaligned_v2, false, NO_FLAG, P4, "comment")
1682+
ENTRY(cuMemcpy2DUnaligned, cuMemcpy2DUnaligned_v2, true, NO_FLAG, P4, "Successful")
16831683
ENTRY(cuMemcpy3D, cuMemcpy3D_v2, true, NO_FLAG, P4, "Successful")
16841684
ENTRY(cuMemcpy3DAsync, cuMemcpy3DAsync_v2, true, NO_FLAG, P4, "comment")
16851685
ENTRY(cuMemcpy3DPeer, cuMemcpy3DPeer, true, NO_FLAG, P4, "Successful")
@@ -2005,7 +2005,7 @@ ENTRY(cuTensorMapReplaceAddress, cuTensorMapReplaceAddress, false, NO_FLAG, P4,
20052005
ENTRY(cuCtxDisablePeerAccess, cuCtxDisablePeerAccess, false, NO_FLAG, P4, "comment")
20062006
ENTRY(cuCtxEnablePeerAccess, cuCtxEnablePeerAccess, true, NO_FLAG, P0, "DPCT1026/DPCT1027")
20072007
ENTRY(cuDeviceCanAccessPeer, cuDeviceCanAccessPeer, true, NO_FLAG, P4, "DPCT1031")
2008-
ENTRY(cuDeviceGetP2PAttribute, cuDeviceGetP2PAttribute, false, NO_FLAG, P4, "comment")
2008+
ENTRY(cuDeviceGetP2PAttribute, cuDeviceGetP2PAttribute, true, NO_FLAG, P4, "Partial: all the attr types except perf rank, DPCT1007")
20092009

20102010
// Graphics Interoperability
20112011
ENTRY(cuGraphicsMapResources, cuGraphicsMapResources, true, NO_FLAG, P4, "successful/DPCT1119")

‎clang/test/dpct/driver-mem-usm-none.cu

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -100,6 +100,8 @@ int main(){
100100

101101
// CHECK: dpct::dpct_memcpy(cpy);
102102
cuMemcpy2D(&cpy);
103+
// CHECK: dpct::dpct_memcpy(cpy);
104+
cuMemcpy2DUnaligned(&cpy);
103105
// CHECK: dpct::async_dpct_memcpy(cpy, *stream);
104106
cuMemcpy2DAsync(&cpy, stream);
105107

‎clang/test/dpct/driver-mem.cu

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -162,6 +162,8 @@ int main(){
162162

163163
// CHECK: dpct::dpct_memcpy(cpy);
164164
cuMemcpy2D(&cpy);
165+
// CHECK: dpct::dpct_memcpy(cpy);
166+
cuMemcpy2DUnaligned(&cpy);
165167
// CHECK: dpct::async_dpct_memcpy(cpy, *stream);
166168
cuMemcpy2DAsync(&cpy, stream);
167169

‎clang/test/dpct/driver_context.cu

Lines changed: 41 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -43,6 +43,47 @@ int main(){
4343
// CHECK-NEXT: */
4444
cuDevicePrimaryCtxRelease(device);
4545

46+
// CHECK: unsigned int flags = 0;
47+
// CHECK-NEXT: flags = 0;
48+
// CHECK-NEXT: flags = 0;
49+
unsigned int flags = CU_CTX_BLOCKING_SYNC;
50+
flags = CU_CTX_SCHED_AUTO;
51+
flags = CU_CTX_SCHED_YIELD;
52+
53+
#ifndef NO_BUILD_TEST
54+
// CHECK: flags = 0;
55+
#if (CUDA_VERSION >= 12010)
56+
flags = CU_CTX_COREDUMP_ENABLE;
57+
#else
58+
flags = CU_CTX_BLOCKING_SYNC;
59+
#endif
60+
61+
// CHECK: flags = 0;
62+
#if (CUDA_VERSION >= 12010)
63+
flags = CU_CTX_USER_COREDUMP_ENABLE;
64+
#else
65+
flags = CU_CTX_SCHED_AUTO;
66+
#endif
67+
68+
// CHECK: flags = 0;
69+
#if (CUDA_VERSION >= 12030)
70+
flags = CU_CTX_SYNC_MEMOPS;
71+
#else
72+
flags = CU_CTX_SCHED_YIELD;
73+
#endif
74+
#endif
75+
int active;
76+
77+
// CHECK: /*
78+
// CHECK-NEXT: DPCT1026:{{[0-9]+}}: The call to cuDevicePrimaryCtxSetFlags_v2 was removed because SYCL currently does not support setting device context flags.
79+
// CHECK-NEXT: */
80+
cuDevicePrimaryCtxSetFlags(device, flags);
81+
82+
// CHECK: /*
83+
// CHECK-NEXT: DPCT1026:{{[0-9]+}}: The call to cuDevicePrimaryCtxGetState was removed because SYCL currently does not support querying device context flags.
84+
// CHECK-NEXT: */
85+
cuDevicePrimaryCtxGetState(device, &flags, &active);
86+
4687
// CHECK: MY_SAFE_CALL(DPCT_CHECK_ERROR(ctx = dpct::push_device_for_curr_thread(device)));
4788
MY_SAFE_CALL(cuCtxCreate(&ctx, CU_CTX_LMEM_RESIZE_TO_MAX, device));
4889

0 commit comments

Comments
 (0)