From 81af8212870cffedb4efb0d4adb1b14c78ab3d4d Mon Sep 17 00:00:00 2001 From: intwanghao Date: Tue, 8 Oct 2024 16:46:07 +0800 Subject: [PATCH 01/10] fix Signed-off-by: intwanghao --- clang/include/clang/DPCT/DPCTOptions.inc | 5 + clang/lib/DPCT/APINames.inc | 16 +- clang/lib/DPCT/APINamesMemory.inc | 101 +++++++++ clang/lib/DPCT/ASTTraversal.cpp | 76 +++++++ clang/lib/DPCT/AnalysisInfo.h | 3 + clang/lib/DPCT/CallExprRewriterCommon.h | 4 + clang/lib/DPCT/MapNames.cpp | 57 ++++++ clang/lib/DPCT/TypeNames.inc | 5 - clang/lib/DPCT/ValidateArguments.h | 1 + clang/runtime/dpct-rt/include/dpct/memory.hpp | 191 ++++++++++++++++++ clang/test/dpct/virtual_memory.cu | 81 ++++++++ 11 files changed, 527 insertions(+), 13 deletions(-) create mode 100644 clang/test/dpct/virtual_memory.cu diff --git a/clang/include/clang/DPCT/DPCTOptions.inc b/clang/include/clang/DPCT/DPCTOptions.inc index 63cfdfea4778..a3b0de8eb321 100644 --- a/clang/include/clang/DPCT/DPCTOptions.inc +++ b/clang/include/clang/DPCT/DPCTOptions.inc @@ -787,6 +787,11 @@ DPCT_ENUM_OPTION( "be accessed within a kernel using syntax similar to C++ global " "variables.\n", false), + DPCT_OPTION_ENUM_VALUE( + "virtual_memory", int(ExperimentalFeatures::Exp_VirtualMemory), + "Experimental extension that allows map an address range onto " + "multiple allocations of physical memory.", + false), DPCT_OPTION_ENUM_VALUE( "all", int(ExperimentalFeatures::Exp_All), "Enable all experimental extensions listed in this option.\n", diff --git a/clang/lib/DPCT/APINames.inc b/clang/lib/DPCT/APINames.inc index 530ad2aa277c..a1cd9cd4b537 100644 --- a/clang/lib/DPCT/APINames.inc +++ b/clang/lib/DPCT/APINames.inc @@ -1717,20 +1717,20 @@ ENTRY(cuMipmappedArrayGetMemoryRequirements, cuMipmappedArrayGetMemoryRequiremen ENTRY(cuMipmappedArrayGetSparseProperties, cuMipmappedArrayGetSparseProperties, false, NO_FLAG, P7, "comment") // Virtual Memory Management -ENTRY(cuMemAddressFree, cuMemAddressFree, false, NO_FLAG, P4, "comment") -ENTRY(cuMemAddressReserve, cuMemAddressReserve, false, NO_FLAG, P4, "comment") -ENTRY(cuMemCreate, cuMemCreate, false, NO_FLAG, P4, "comment") +ENTRY(cuMemAddressFree, cuMemAddressFree, true, NO_FLAG, P4, "comment") +ENTRY(cuMemAddressReserve, cuMemAddressReserve, true, NO_FLAG, P4, "comment") +ENTRY(cuMemCreate, cuMemCreate, true, NO_FLAG, P4, "comment") ENTRY(cuMemExportToShareableHandle, cuMemExportToShareableHandle, false, NO_FLAG, P4, "comment") ENTRY(cuMemGetAccess, cuMemGetAccess, false, NO_FLAG, P4, "comment") -ENTRY(cuMemGetAllocationGranularity, cuMemGetAllocationGranularity, false, NO_FLAG, P4, "comment") +ENTRY(cuMemGetAllocationGranularity, cuMemGetAllocationGranularity, true, NO_FLAG, P4, "comment") ENTRY(cuMemGetAllocationPropertiesFromHandle, cuMemGetAllocationPropertiesFromHandle, false, NO_FLAG, P4, "comment") ENTRY(cuMemImportFromShareableHandle, cuMemImportFromShareableHandle, false, NO_FLAG, P4, "comment") -ENTRY(cuMemMap, cuMemMap, false, NO_FLAG, P4, "comment") +ENTRY(cuMemMap, cuMemMap, true, NO_FLAG, P4, "comment") ENTRY(cuMemMapArrayAsync, cuMemMapArrayAsync, false, NO_FLAG, P7, "comment") -ENTRY(cuMemRelease, cuMemRelease, false, NO_FLAG, P4, "comment") +ENTRY(cuMemRelease, cuMemRelease, true, NO_FLAG, P4, "comment") ENTRY(cuMemRetainAllocationHandle, cuMemRetainAllocationHandle, false, NO_FLAG, P7, "comment") -ENTRY(cuMemSetAccess, cuMemSetAccess, false, NO_FLAG, P4, "comment") -ENTRY(cuMemUnmap, cuMemUnmap, false, NO_FLAG, P4, "comment") +ENTRY(cuMemSetAccess, cuMemSetAccess, true, NO_FLAG, P4, "comment") +ENTRY(cuMemUnmap, cuMemUnmap, true, NO_FLAG, P4, "comment") // Stream Ordered Memory Allocator ENTRY(cuMemAllocAsync, cuMemAllocAsync, false, NO_FLAG, P7, "comment") diff --git a/clang/lib/DPCT/APINamesMemory.inc b/clang/lib/DPCT/APINamesMemory.inc index ed5576aee733..6a9e8003fa46 100644 --- a/clang/lib/DPCT/APINamesMemory.inc +++ b/clang/lib/DPCT/APINamesMemory.inc @@ -698,6 +698,107 @@ ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY( ARG(0), ARG(1), ARG(2), ARG(3), ARG(4), DEREF(makeCallArgCreatorWithCall(5)))))) +CONDITIONAL_FACTORY_ENTRY( + UseExpVirtualMemory, + ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY( + HelperFeatureEnum::device_ext, + CALL_FACTORY_ENTRY("cuMemCreate", + CALL(MapNames::getDpctNamespace() + + "experimental::mem_create", + ARG(0), ARG(1), ARG(2), ARG(3))))), + UNSUPPORT_FACTORY_ENTRY("cuMemCreate", + Diagnostics::TRY_EXPERIMENTAL_FEATURE, + ARG("cuMemCreate"), + ARG("--use-experimental-features=virtual_memory"))) + +CONDITIONAL_FACTORY_ENTRY( + UseExpVirtualMemory, + ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY( + HelperFeatureEnum::device_ext, + CALL_FACTORY_ENTRY("cuMemAddressReserve", + CALL(MapNames::getDpctNamespace() + + "experimental::mem_address_reserve", + ARG(0), ARG(1), ARG(2), ARG(3), ARG(4))))), + UNSUPPORT_FACTORY_ENTRY("cuMemAddressReserve", + Diagnostics::TRY_EXPERIMENTAL_FEATURE, + ARG("cuMemAddressReserve"), + ARG("--use-experimental-features=virtual_memory"))) + +CONDITIONAL_FACTORY_ENTRY( + UseExpVirtualMemory, + ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY( + HelperFeatureEnum::device_ext, + CALL_FACTORY_ENTRY("cuMemAddressFree", + CALL(MapNames::getDpctNamespace() + + "experimental::mem_address_free", + ARG(0), ARG(1))))), + UNSUPPORT_FACTORY_ENTRY("cuMemAddressFree", + Diagnostics::TRY_EXPERIMENTAL_FEATURE, + ARG("cuMemAddressFree"), + ARG("--use-experimental-features=virtual_memory"))) + +CONDITIONAL_FACTORY_ENTRY( + UseExpVirtualMemory, + ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY( + HelperFeatureEnum::device_ext, + CALL_FACTORY_ENTRY( + "cuMemGetAllocationGranularity", + CALL(MapNames::getDpctNamespace() + + "experimental::mem_get_allocation_granularity", + ARG(0), ARG(1), ARG(2))))), + UNSUPPORT_FACTORY_ENTRY("cuMemGetAllocationGranularity", + Diagnostics::TRY_EXPERIMENTAL_FEATURE, + ARG("cuMemGetAllocationGranularity"), + ARG("--use-experimental-features=virtual_memory"))) + +CONDITIONAL_FACTORY_ENTRY( + UseExpVirtualMemory, + ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY( + HelperFeatureEnum::device_ext, + CALL_FACTORY_ENTRY("cuMemRelease", CALL(MapNames::getDpctNamespace() + + "experimental::mem_release", + ARG(0))))), + UNSUPPORT_FACTORY_ENTRY("cuMemRelease", + Diagnostics::TRY_EXPERIMENTAL_FEATURE, + ARG("cuMemRelease"), + ARG("--use-experimental-features=virtual_memory"))) + +CONDITIONAL_FACTORY_ENTRY( + UseExpVirtualMemory, + ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY( + HelperFeatureEnum::device_ext, + CALL_FACTORY_ENTRY("cuMemMap", + CALL(MapNames::getDpctNamespace() + + "experimental::mem_map", + ARG(0), ARG(1), ARG(2), ARG(3), ARG(4))))), + UNSUPPORT_FACTORY_ENTRY("cuMemMap", Diagnostics::TRY_EXPERIMENTAL_FEATURE, + ARG("cuMemMap"), + ARG("--use-experimental-features=virtual_memory"))) + +CONDITIONAL_FACTORY_ENTRY( + UseExpVirtualMemory, + ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY( + HelperFeatureEnum::device_ext, + CALL_FACTORY_ENTRY("cuMemUnmap", CALL(MapNames::getDpctNamespace() + + "experimental::mem_unmap", + ARG(0), ARG(1))))), + UNSUPPORT_FACTORY_ENTRY("cuMemUnmap", Diagnostics::TRY_EXPERIMENTAL_FEATURE, + ARG("cuMemUnmap"), + ARG("--use-experimental-features=virtual_memory"))) + +CONDITIONAL_FACTORY_ENTRY( + UseExpVirtualMemory, + ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY( + HelperFeatureEnum::device_ext, + CALL_FACTORY_ENTRY("cuMemSetAccess", + CALL(MapNames::getDpctNamespace() + + "experimental::mem_set_access", + ARG(0), ARG(1), ARG(2), ARG(3))))), + UNSUPPORT_FACTORY_ENTRY("cuMemSetAccess", + Diagnostics::TRY_EXPERIMENTAL_FEATURE, + ARG("cuMemSetAccess"), + ARG("--use-experimental-features=virtual_memory"))) + CONDITIONAL_FACTORY_ENTRY( UsePeerAccess(), ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY( diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index 569219a48e55..23b373731eab 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/clang/lib/DPCT/ASTTraversal.cpp @@ -11228,6 +11228,30 @@ void MemoryMigrationRule::registerMatcher(MatchFinder &MF) { hasParent(callExpr(parentStmt()).bind("callExpr"))) .bind("unresolvedCall"), this); + auto virtualmemoryAPI = [&]() { + return hasAnyName("cuMemCreate", "cuMemAddressReserve", "cuMemMap", + "cuMemUnmap", "cuMemAddressFree", "cuMemRelease", + "cuMemSetAccess", "cuMemGetAllocationGranularity"); + }; + auto virtualmemoryType = [&]() { + return hasAnyName("CUmemAllocationProp", "CUmemGenericAllocationHandle", + "CUmemAccessDesc"); + }; + auto virtualmemoryEnum = [&]() { + return hasAnyName("CU_MEM_ALLOCATION_TYPE_PINNED", + "CU_MEM_LOCATION_TYPE_DEVICE", + "CU_MEM_ACCESS_FLAGS_PROT_READWRITE", + "CU_MEM_ALLOC_GRANULARITY_RECOMMENDED"); + }; + MF.addMatcher( + callExpr(callee(functionDecl(virtualmemoryAPI()))).bind("vmCall"), this); + MF.addMatcher( + typeLoc(loc(qualType(hasDeclaration(namedDecl(virtualmemoryType()))))) + .bind("vmType"), + this); + MF.addMatcher( + declRefExpr(to(enumConstantDecl(virtualmemoryEnum()))).bind("vmEnum"), + this); } void MemoryMigrationRule::runRule(const MatchFinder::MatchResult &Result) { @@ -11330,6 +11354,58 @@ void MemoryMigrationRule::runRule(const MatchFinder::MatchResult &Result) { getAssistNodeAsType(Result, "callExpr"), /* IsAssigned */ false, getAssistNodeAsType(Result, "unresolvedCall")); + + auto &SM = DpctGlobalInfo::getSourceManager(); + if (const CallExpr *CE = getNodeAsType(Result, "vmCall")) { + ExprAnalysis EA(CE); + emplaceTransformation(EA.getReplacement()); + EA.applyAllSubExprRepl(); + } + if (auto TL = getNodeAsType(Result, "vmType")) { + auto TypeStr = + DpctGlobalInfo::getTypeName(TL->getType().getUnqualifiedType()); + if (!DpctGlobalInfo::useExpVirtualMemory()) { + report(TL->getBeginLoc(), Diagnostics::TRY_EXPERIMENTAL_FEATURE, false, + TypeStr, "--use-experimental-features=virtual_memory"); + return; + } + if (!DpctGlobalInfo::isInAnalysisScope( + SM.getSpellingLoc(TL->getBeginLoc()))) { + return; + } + auto Range = getDefinitionRange(TL->getBeginLoc(), TL->getEndLoc()); + auto BeginLoc = Range.getBegin(); + auto EndLoc = Range.getEnd(); + + if (SM.isWrittenInScratchSpace(SM.getSpellingLoc(TL->getBeginLoc()))) { + BeginLoc = SM.getExpansionRange(TL->getBeginLoc()).getBegin(); + EndLoc = SM.getExpansionRange(TL->getBeginLoc()).getEnd(); + } + std::string Str = + MapNames::findReplacedName(MapNames::TypeNamesMap, TypeStr); + if (!Str.empty()) { + auto Len = Lexer::MeasureTokenLength( + EndLoc, SM, DpctGlobalInfo::getContext().getLangOpts()); + Len += SM.getDecomposedLoc(EndLoc).second - + SM.getDecomposedLoc(BeginLoc).second; + emplaceTransformation(new ReplaceText(BeginLoc, Len, std::move(Str))); + return; + } + } + if (auto *E = getNodeAsType(Result, "vmEnum")) { + std::string EnumName = E->getNameInfo().getName().getAsString(); + if (!DpctGlobalInfo::useExpVirtualMemory()) { + report(E->getBeginLoc(), Diagnostics::TRY_EXPERIMENTAL_FEATURE, false, + EnumName, "--use-experimental-features=virtual_memory"); + return; + } + auto Search = EnumConstantRule::EnumNamesMap.find(EnumName); + if (Search == EnumConstantRule::EnumNamesMap.end()) { + report(E->getBeginLoc(), Diagnostics::API_NOT_MIGRATED, false, EnumName); + return; + } + emplaceTransformation(new ReplaceStmt(E, Search->second->NewName)); + } } void MemoryMigrationRule::getSymbolAddressMigration( diff --git a/clang/lib/DPCT/AnalysisInfo.h b/clang/lib/DPCT/AnalysisInfo.h index f0c718ea8410..d8abb76affb5 100644 --- a/clang/lib/DPCT/AnalysisInfo.h +++ b/clang/lib/DPCT/AnalysisInfo.h @@ -1290,6 +1290,9 @@ class DpctGlobalInfo { static bool useExpDeviceGlobal() { return getUsingExperimental(); } + static bool useExpVirtualMemory() { + return getUsingExperimental(); + } static bool useNoQueueDevice() { return getHelperFuncPreference(HelperFuncPreference::NoQueueDevice); } diff --git a/clang/lib/DPCT/CallExprRewriterCommon.h b/clang/lib/DPCT/CallExprRewriterCommon.h index 374a19941dc6..d289a2d86caf 100644 --- a/clang/lib/DPCT/CallExprRewriterCommon.h +++ b/clang/lib/DPCT/CallExprRewriterCommon.h @@ -1713,6 +1713,10 @@ inline auto UseSYCLCompat = [](const CallExpr *C) -> bool { return DpctGlobalInfo::useSYCLCompat(); }; +inline auto UseExpVirtualMemory = [](const CallExpr *C) -> bool { + return DpctGlobalInfo::useExpVirtualMemory(); +}; + class CheckDerefedTypeBeforeCast { unsigned Idx; std::string TypeName; diff --git a/clang/lib/DPCT/MapNames.cpp b/clang/lib/DPCT/MapNames.cpp index b013d3c74cbd..aed97edb306c 100644 --- a/clang/lib/DPCT/MapNames.cpp +++ b/clang/lib/DPCT/MapNames.cpp @@ -832,6 +832,15 @@ void MapNames::setExplicitNamespaceMap( std::make_shared( getLibraryHelperNamespace() + "blas_gemm::experimental::transform_desc_ptr")}, + {"CUmemAllocationProp", + std::make_shared(getDpctNamespace() + + "experimental::mem_prop")}, + {"CUmemGenericAllocationHandle", + std::make_shared(getDpctNamespace() + + "experimental::mem_handle")}, + {"CUmemAccessDesc", + std::make_shared(getDpctNamespace() + + "experimental::mem_access_desc")}, {"cudaGraphicsMapFlags", std::make_shared("int")}, {"cudaGraphicsRegisterFlags", std::make_shared("int")}, // ... @@ -1439,6 +1448,54 @@ void MapNames::setExplicitNamespaceMap( std::make_shared("0")}, {"CU_MEM_ADVISE_SET_ACCESSED_BY", std::make_shared("0")}, {"CU_MEM_ADVISE_UNSET_ACCESSED_BY", std::make_shared("0")}, + {"CU_MEM_ALLOCATION_TYPE_PINNED", + std::make_shared( + getDpctNamespace() + + "experimental::mem_allocation_type::MEM_ALLOCATION_TYPE_DEFAULT")}, + {"CU_MEM_ALLOCATION_TYPE_INVALID", + std::make_shared( + getDpctNamespace() + + "experimental::mem_allocation_type::MEM_ALLOCATION_TYPE_INVALID")}, + {"CU_MEM_ALLOCATION_TYPE_MAX", + std::make_shared( + getDpctNamespace() + + "experimental::mem_allocation_type::MEM_ALLOCATION_TYPE_MAX")}, + {"CU_MEM_LOCATION_TYPE_DEVICE", + std::make_shared( + getDpctNamespace() + + "experimental::mem_location_type::MEM_LOCATION_TYPE_DEVICE")}, + {"CU_MEM_LOCATION_TYPE_INVALID", + std::make_shared( + getDpctNamespace() + + "experimental::mem_location_type::MEM_LOCATION_TYPE_INVALID")}, + {"CU_MEM_LOCATION_TYPE_MAX", + std::make_shared( + getDpctNamespace() + + "experimental::mem_location_type::MEM_LOCATION_TYPE_MAX")}, + {"CU_MEM_ACCESS_FLAGS_PROT_READWRITE", + std::make_shared(getDpctNamespace() + + "experimental::address_access_flags::" + "ADDRESS_ACCESS_FLAGS_READ_WRITE")}, + {"CU_MEM_ACCESS_FLAGS_PROT_NONE", + std::make_shared( + getDpctNamespace() + + "experimental::address_access_flags::ADDRESS_ACCESS_FLAGS_NONE")}, + {"CU_MEM_ACCESS_FLAGS_PROT_MAX", + std::make_shared( + getDpctNamespace() + + "experimental::address_access_flags::ADDRESS_ACCESS_FLAGS_MAX")}, + {"CU_MEM_ACCESS_FLAGS_PROT_READ", + std::make_shared( + getDpctNamespace() + + "experimental::address_access_flags::ADDRESS_ACCESS_FLAGS_READ")}, + {"CU_MEM_ALLOC_GRANULARITY_RECOMMENDED", + std::make_shared( + getDpctNamespace() + + "experimental::granularity_flags::GRANULARITY_FLAGS_RECOMMENDED")}, + {"CU_MEM_ALLOC_GRANULARITY_MINIMUM", + std::make_shared( + getDpctNamespace() + + "experimental::granularity_flags::GRANULARITY_FLAGS_MINIMUM")}, // enum Driver Device Attribute {"CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR", std::make_shared("get_major_version", diff --git a/clang/lib/DPCT/TypeNames.inc b/clang/lib/DPCT/TypeNames.inc index c15d66f1a4c8..fe0949cebfcd 100644 --- a/clang/lib/DPCT/TypeNames.inc +++ b/clang/lib/DPCT/TypeNames.inc @@ -137,11 +137,6 @@ ENTRY_TYPE(cooperative_groups::__v1::multi_grid_group, false, NO_FLAG, P4, "comm ENTRY_TYPE(cooperative_groups::__v1::thread_block_tile, true, NO_FLAG, P4, "Successful") ENTRY_TYPE(cooperative_groups::__v1::coalesced_group, true, NO_FLAG, P4, "Successful/DPCT1119") -//CUmem -ENTRY_TYPE(CUmemGenericAllocationHandle, false, NO_FLAG, P4, "comment") -ENTRY_TYPE(CUmemAllocationProp, false, NO_FLAG, P4, "comment") -ENTRY_TYPE(CUmemAccessDesc, false, NO_FLAG, P4, "comment") - // cuRand ENTRY_TYPE(curandStateScrambledSobol64_t, false, NO_FLAG, P4, "comment") ENTRY_TYPE(curandStateSobol64_t, false, NO_FLAG, P4, "comment") diff --git a/clang/lib/DPCT/ValidateArguments.h b/clang/lib/DPCT/ValidateArguments.h index 208d8a850e70..e91a4d1b2881 100644 --- a/clang/lib/DPCT/ValidateArguments.h +++ b/clang/lib/DPCT/ValidateArguments.h @@ -93,6 +93,7 @@ enum class ExperimentalFeatures : unsigned int { Exp_Graph, Exp_NonUniformGroups, Exp_DeviceGlobal, + Exp_VirtualMemory, Exp_ExperimentalFeaturesEnumSize, Exp_All }; diff --git a/clang/runtime/dpct-rt/include/dpct/memory.hpp b/clang/runtime/dpct-rt/include/dpct/memory.hpp index ce8aa699cc81..b262abe82411 100644 --- a/clang/runtime/dpct-rt/include/dpct/memory.hpp +++ b/clang/runtime/dpct-rt/include/dpct/memory.hpp @@ -1417,6 +1417,197 @@ static inline void async_dpct_memset(pitched_data pitch, int val, detail::dpct_memset(q, pitch, val, size); } +namespace experimental { +typedef sycl::ext::oneapi::experimental::physical_mem *mem_handle; + +enum mem_location_type { + MEM_LOCATION_TYPE_INVALID = 0x0, + MEM_LOCATION_TYPE_DEVICE = 0x1, + MEM_LOCATION_TYPE_MAX = 0xFFFFFFFF +}; + +enum mem_allocation_type { + MEM_ALLOCATION_TYPE_INVALID = 0x0, + MEM_ALLOCATION_TYPE_DEFAULT = 0x1, + MEM_ALLOCATION_TYPE_MAX = 0xFFFFFFFF +}; + +enum granularity_flags { + GRANULARITY_FLAGS_MINIMUM = 0x0, + GRANULARITY_FLAGS_RECOMMENDED = 0x1 +}; + +enum address_access_flags { + ADDRESS_ACCESS_FLAGS_NONE = 0x0, + ADDRESS_ACCESS_FLAGS_READ = 0x1, + ADDRESS_ACCESS_FLAGS_READ_WRITE = 0x3, + ADDRESS_ACCESS_FLAGS_MAX = 0xFFFFFFFF +}; + +struct mem_location { + int id; + mem_location_type type; +}; + +struct mem_prop { + mem_location location; + mem_allocation_type type; +}; + +struct mem_access_desc { + address_access_flags flags; + mem_location location; +}; + +/** + * @brief Reserves a virtual memory region of \p size bytes. The \p addr + * specifies the requested start of the new virtual memory range reservation. If + * the implementation is unable to reserve the virtual memory range at the + * specified address, the implementation will pick another suitable address. The + * \p size must be a multiple of the reserve granularity. The \p alignment + * parameter must be set to 0, which means the default alignment will be used. + * The \p addr must be aligned in accordance with the reserve granularity. + * @param [out] ptr Pointer to the reserved virtual memory. + * @param [in] size Size of the virtual memory region in bytes. + * @param [in] alignment Alignment of the virtual memory (must be 0). + * @param [in] addr Base address of the memory to be reserved. + * @param [in] flags Reserved for future use. + */ +void mem_address_reserve(device_ptr *ptr, size_t size, size_t alignment, + device_ptr addr, unsigned long long flags) { + *ptr = (device_ptr)sycl::ext::oneapi::experimental::reserve_virtual_mem( + (uintptr_t)addr, size, dpct::get_current_device().get_context()); +} + +/** + * @brief Frees a previously reserved virtual memory region pointed by \p ptr. + * @param [in] ptr Pointer to the device memory region to free. + * @param [in] size Size of the memory region in bytes. + */ +void mem_address_free(device_ptr ptr, size_t size) { + sycl::ext::oneapi::experimental::free_virtual_mem( + (uintptr_t)ptr, size, dpct::get_current_device().get_context()); +} + +/** + * @brief Creates a physical memory object for allocation. + * This will allocate \p size of physical memory on the device. The \p size must + * be a multiple of the allocation granularity, as returned by a call to + * get_mem_allocation_granularity. + * @param [out] handle Handle to the created memory object. + * @param [in] size Size of the physical memory in bytes. + * @param [in] prop Properties for the memory allocation. + * @param [in] flags Reserved for future use. + */ +void mem_create(mem_handle *handle, size_t size, const mem_prop *prop, + unsigned long long flags) { + auto &device = dpct::get_device(prop->location.id); + *handle = new sycl::ext::oneapi::experimental::physical_mem( + device, device.get_context(), size); +} + +/** + * @brief Releases the physical memory object. + * @param [in] handle Handle to the memory object to release. + */ +void mem_release(mem_handle handle) { + if (handle) { + delete handle; + } +} + +/** + * @brief Maps a virtual memory range, specified by \p ptr and \p size, to the + * physical memory specified by \p handle, starting at an offset of \p offset + * bytes. The \p ptr, \p size and \p offset are all need to be multiples of the + * allocation granularity. + * @param [in] ptr Pointer to the virtual memory region. + * @param [in] size Size of the memory region in bytes. + * @param [in] offset Offset into the physical memory. + * @param [in] handle Handle to the physical memory. + * @param [in] flags Reserved for future use. + */ +void mem_map(device_ptr ptr, size_t size, size_t offset, mem_handle handle, + unsigned long long flags) { + handle->map((uintptr_t)ptr, size, + sycl::ext::oneapi::experimental::address_access_mode::read_write, + offset); +} + +/** + * @brief Unmaps a previously mapped virtual memory region. + * @param [in] ptr Pointer to the virtual memory region to unmap. + * @param [in] size Size of the memory region in bytes. + */ +void mem_unmap(device_ptr ptr, size_t size) { + sycl::ext::oneapi::experimental::unmap( + ptr, size, dpct::get_current_device().get_context()); +} + +/** + * @brief Sets the access mode of a virtual memory region pointed by \p ptr. + * @param [in] ptr Pointer to the virtual memory region. + * @param [in] size Size of the memory region in bytes. + * @param [in] desc Array of access descriptors. + * @param [in] count Number of access descriptors. + */ +void mem_set_access(device_ptr ptr, size_t size, const mem_access_desc *desc, + size_t count) { + for (size_t index = 0; index < count; index++) { + sycl::ext::oneapi::experimental::address_access_mode mode; + switch (desc[index].flags) { + case address_access_flags::ADDRESS_ACCESS_FLAGS_NONE: + mode = sycl::ext::oneapi::experimental::address_access_mode::none; + break; + case address_access_flags::ADDRESS_ACCESS_FLAGS_READ: + mode = sycl::ext::oneapi::experimental::address_access_mode::read; + break; + case address_access_flags::ADDRESS_ACCESS_FLAGS_READ_WRITE: + mode = sycl::ext::oneapi::experimental::address_access_mode::read_write; + break; + default: + mode = sycl::ext::oneapi::experimental::address_access_mode::none; + break; + } + sycl::ext::oneapi::experimental::set_access_mode( + ptr, size, mode, + dpct::get_device(desc[index].location.id).get_context()); + } +} + +/** + * @brief Retrieves the granularity of the memory allocation. + * @param [out] granularity Pointer to store the retrieved granularity. + * @param [in] prop Properties for the memory allocation. + * @param [in] option Granularity option, either minimum or recommended. + */ +void mem_get_allocation_granularity(size_t *granularity, const mem_prop *prop, + granularity_flags option) { + auto &device = dpct::get_device(prop->location.id); + *granularity = sycl::ext::oneapi::experimental::get_mem_granularity( + device, device.get_context(), + option == granularity_flags::GRANULARITY_FLAGS_MINIMUM + ? sycl::ext::oneapi::experimental::granularity_mode::minimum + : sycl::ext::oneapi::experimental::granularity_mode::recommended); +} + +/** + * @brief Retrieves the granularity of the virtual memory region reserve. + * @param [out] granularity Pointer to store the retrieved granularity. + * @param [in] prop Properties for the virtual memory region reserve. + * @param [in] option Granularity option, either minimum or recommended. + */ +void mem_get_reserve_granularity(size_t *granularity, const mem_prop *prop, + granularity_flags option) { + auto &device = dpct::get_device(prop->location.id); + *granularity = sycl::ext::oneapi::experimental::get_mem_granularity( + device.get_context(), + option == granularity_flags::GRANULARITY_FLAGS_MINIMUM + ? sycl::ext::oneapi::experimental::granularity_mode::minimum + : sycl::ext::oneapi::experimental::granularity_mode::recommended); +} +} // namespace experimental + /// dpct accessor used as device function parameter. template class accessor; template class accessor { diff --git a/clang/test/dpct/virtual_memory.cu b/clang/test/dpct/virtual_memory.cu new file mode 100644 index 000000000000..deacfa98df74 --- /dev/null +++ b/clang/test/dpct/virtual_memory.cu @@ -0,0 +1,81 @@ +// RUN: dpct --format-range=none --use-experimental-features=virtual_memory -out-root %T/virtual_memory %s --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only +// RUN: FileCheck %s --match-full-lines --input-file %T/virtual_memory/virtual_memory.dp.cpp +#include +#include + +#define SIZE 100 +int main() { + // Initialize CUDA driver API + cuInit(0); + + CUdevice device; + cuDeviceGet(&device, 0); + + CUcontext context; + cuCtxCreate(&context, 0, device); + +// CHECK: dpct::experimental::mem_prop prop = {}; +// CHECK: prop.type = dpct::experimental::mem_allocation_type::MEM_ALLOCATION_TYPE_DEFAULT; +// CHECK: prop.location.type = dpct::experimental::mem_location_type::MEM_LOCATION_TYPE_DEVICE; +// CHECK: prop.location.id = device; + CUmemAllocationProp prop = {}; + prop.type = CU_MEM_ALLOCATION_TYPE_PINNED; + prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE; + prop.location.id = device; + size_t granularity; +// CHECK: dpct::experimental::mem_get_allocation_granularity(&granularity, &prop, dpct::experimental::granularity_flags::GRANULARITY_FLAGS_MINIMUM); + cuMemGetAllocationGranularity(&granularity, &prop, CU_MEM_ALLOC_GRANULARITY_MINIMUM); + size_t POOL_SIZE = granularity; + +// CHECK: dpct::device_ptr reserved_addr; +// CHECK: dpct::experimental::mem_handle allocHandle; +// CHECK: dpct::experimental::mem_address_reserve(&reserved_addr, POOL_SIZE, 0, 0, 0); +// CHECK: dpct::experimental::mem_create(&allocHandle, POOL_SIZE, &prop, 0); +// CHECK: dpct::experimental::mem_map(reserved_addr, POOL_SIZE, 0, allocHandle, 0); + CUdeviceptr reserved_addr; + CUmemGenericAllocationHandle allocHandle; + cuMemAddressReserve(&reserved_addr, POOL_SIZE, 0, 0, 0); + cuMemCreate(&allocHandle, POOL_SIZE, &prop, 0); + cuMemMap(reserved_addr, POOL_SIZE, 0, allocHandle, 0); + +// CHECK: dpct::experimental::mem_access_desc accessDesc = {}; +// CHECK: accessDesc.location.type = dpct::experimental::mem_location_type::MEM_LOCATION_TYPE_DEVICE; +// CHECK: accessDesc.location.id = device; +// CHECK: accessDesc.flags = dpct::experimental::address_access_flags::ADDRESS_ACCESS_FLAGS_READ_WRITE; +// CHECK: dpct::experimental::mem_set_access(reserved_addr, POOL_SIZE, &accessDesc, 1); + CUmemAccessDesc accessDesc = {}; + accessDesc.location.type = CU_MEM_LOCATION_TYPE_DEVICE; + accessDesc.location.id = device; + accessDesc.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE; + cuMemSetAccess(reserved_addr, POOL_SIZE, &accessDesc, 1); + int* host_data = new int[SIZE]; + int* host_data2 = new int[SIZE]; + for (int i = 0; i < SIZE; ++i) { + host_data[i] = i; + host_data2[i] = 0; + } + + cuMemcpyHtoD(reserved_addr, host_data, SIZE * sizeof(int)); + cuMemcpyDtoH(host_data2, reserved_addr, SIZE * sizeof(int)); + + for (int i = 0; i < SIZE; ++i) { + if(host_data[i] != host_data2[i]) { + std::cout << "test failed" << std::endl; + exit(-1); + } + } + std::cout << "test passed" << std::endl; + +// CHECK: dpct::experimental::mem_unmap(reserved_addr, POOL_SIZE); +// CHECK: dpct::experimental::mem_release(allocHandle); +// CHECK: dpct::experimental::mem_address_free(reserved_addr, POOL_SIZE); + cuMemUnmap(reserved_addr, POOL_SIZE); + cuMemRelease(allocHandle); + cuMemAddressFree(reserved_addr, POOL_SIZE); + + delete[] host_data; + delete[] host_data2; + + cuCtxDestroy(context); + return 0; +} From 74c4c34e8a58e9b70047a6adee54eceeab26d1ac Mon Sep 17 00:00:00 2001 From: intwanghao Date: Wed, 9 Oct 2024 09:20:55 +0800 Subject: [PATCH 02/10] fix Signed-off-by: intwanghao --- clang/lib/DPCT/MapNames.cpp | 12 +++++++----- 1 file changed, 7 insertions(+), 5 deletions(-) diff --git a/clang/lib/DPCT/MapNames.cpp b/clang/lib/DPCT/MapNames.cpp index 7985442756cb..b4265900619c 100644 --- a/clang/lib/DPCT/MapNames.cpp +++ b/clang/lib/DPCT/MapNames.cpp @@ -449,14 +449,16 @@ void MapNames::setExplicitNamespaceMap( getLibraryHelperNamespace() + "sparse::optimize_info>", HelperFeatureEnum::device_ext)}, - {"thrust::device_ptr", - std::make_shared(getLibraryHelperNamespace() + "device_pointer", - HelperFeatureEnum::device_ext)}, + {"thrust::device_ptr", std::make_shared( + getLibraryHelperNamespace() + "device_pointer", + HelperFeatureEnum::device_ext)}, {"thrust::device_reference", - std::make_shared(getLibraryHelperNamespace() + "device_reference", + std::make_shared(getLibraryHelperNamespace() + + "device_reference", HelperFeatureEnum::device_ext)}, {"thrust::device_vector", - std::make_shared(getLibraryHelperNamespace() + "device_vector", + std::make_shared(getLibraryHelperNamespace() + + "device_vector", HelperFeatureEnum::device_ext)}, {"thrust::device_malloc_allocator", std::make_shared(getDpctNamespace() + From 8811fe4530310b6b1dc388bdaf325b8de7b2bc14 Mon Sep 17 00:00:00 2001 From: intwanghao Date: Wed, 9 Oct 2024 09:37:58 +0800 Subject: [PATCH 03/10] fix Signed-off-by: intwanghao --- clang/lib/DPCT/APINames.inc | 4 ++++ clang/test/dpct/virtual_memory.cu | 4 +--- 2 files changed, 5 insertions(+), 3 deletions(-) diff --git a/clang/lib/DPCT/APINames.inc b/clang/lib/DPCT/APINames.inc index a1cd9cd4b537..f2fa3be225ca 100644 --- a/clang/lib/DPCT/APINames.inc +++ b/clang/lib/DPCT/APINames.inc @@ -32,6 +32,8 @@ * ****************************************************************************/ +// clang-format off + // API Names refer to "v12.2" // CUDA runtime API // Device management functions of runtime API @@ -2307,3 +2309,5 @@ ENTRY(__assert_fail, __assert_fail, true, NO_FLAG, P4, "Successful") ENTRY(__assertfail, __assertfail, true, NO_FLAG, P4, "Successful") ENTRY(cuGetExportTable, cuGetExportTable, true, NO_FLAG, P7, "Partial") + +// clang-format on diff --git a/clang/test/dpct/virtual_memory.cu b/clang/test/dpct/virtual_memory.cu index deacfa98df74..74532cb428fb 100644 --- a/clang/test/dpct/virtual_memory.cu +++ b/clang/test/dpct/virtual_memory.cu @@ -4,13 +4,11 @@ #include #define SIZE 100 + int main() { - // Initialize CUDA driver API cuInit(0); - CUdevice device; cuDeviceGet(&device, 0); - CUcontext context; cuCtxCreate(&context, 0, device); From dba0b1d07a92cc770e4b6fca40e554fd0268ee5d Mon Sep 17 00:00:00 2001 From: intwanghao Date: Wed, 9 Oct 2024 22:14:04 +0800 Subject: [PATCH 04/10] fix Signed-off-by: intwanghao --- clang/lib/DPCT/MapNames.cpp | 24 ++++++------ clang/runtime/dpct-rt/include/dpct/memory.hpp | 37 +++++++++---------- clang/test/dpct/virtual_memory.cu | 14 +++---- 3 files changed, 37 insertions(+), 38 deletions(-) diff --git a/clang/lib/DPCT/MapNames.cpp b/clang/lib/DPCT/MapNames.cpp index b4265900619c..bbbb04fd0003 100644 --- a/clang/lib/DPCT/MapNames.cpp +++ b/clang/lib/DPCT/MapNames.cpp @@ -1452,51 +1452,51 @@ void MapNames::setExplicitNamespaceMap( {"CU_MEM_ALLOCATION_TYPE_PINNED", std::make_shared( getDpctNamespace() + - "experimental::mem_allocation_type::MEM_ALLOCATION_TYPE_DEFAULT")}, + "experimental::mem_allocation_type::mem_allocation_type_default")}, {"CU_MEM_ALLOCATION_TYPE_INVALID", std::make_shared( getDpctNamespace() + - "experimental::mem_allocation_type::MEM_ALLOCATION_TYPE_INVALID")}, + "experimental::mem_allocation_type::mem_allocation_type_invalid")}, {"CU_MEM_ALLOCATION_TYPE_MAX", std::make_shared( getDpctNamespace() + - "experimental::mem_allocation_type::MEM_ALLOCATION_TYPE_MAX")}, + "experimental::mem_allocation_type::mem_allocation_type_max")}, {"CU_MEM_LOCATION_TYPE_DEVICE", std::make_shared( getDpctNamespace() + - "experimental::mem_location_type::MEM_LOCATION_TYPE_DEVICE")}, + "experimental::mem_location_type::mem_location_type_device")}, {"CU_MEM_LOCATION_TYPE_INVALID", std::make_shared( getDpctNamespace() + - "experimental::mem_location_type::MEM_LOCATION_TYPE_INVALID")}, + "experimental::mem_location_type::mem_location_type_invalid")}, {"CU_MEM_LOCATION_TYPE_MAX", std::make_shared( getDpctNamespace() + - "experimental::mem_location_type::MEM_LOCATION_TYPE_MAX")}, + "experimental::mem_location_type::mem_location_type_max")}, {"CU_MEM_ACCESS_FLAGS_PROT_READWRITE", std::make_shared(getDpctNamespace() + "experimental::address_access_flags::" - "ADDRESS_ACCESS_FLAGS_READ_WRITE")}, + "address_access_flags_read_write")}, {"CU_MEM_ACCESS_FLAGS_PROT_NONE", std::make_shared( getDpctNamespace() + - "experimental::address_access_flags::ADDRESS_ACCESS_FLAGS_NONE")}, + "experimental::address_access_flags::address_access_flags_none")}, {"CU_MEM_ACCESS_FLAGS_PROT_MAX", std::make_shared( getDpctNamespace() + - "experimental::address_access_flags::ADDRESS_ACCESS_FLAGS_MAX")}, + "experimental::address_access_flags::address_access_flags_max")}, {"CU_MEM_ACCESS_FLAGS_PROT_READ", std::make_shared( getDpctNamespace() + - "experimental::address_access_flags::ADDRESS_ACCESS_FLAGS_READ")}, + "experimental::address_access_flags::address_access_flags_read")}, {"CU_MEM_ALLOC_GRANULARITY_RECOMMENDED", std::make_shared( getDpctNamespace() + - "experimental::granularity_flags::GRANULARITY_FLAGS_RECOMMENDED")}, + "experimental::granularity_flags::granularity_flags_recommended")}, {"CU_MEM_ALLOC_GRANULARITY_MINIMUM", std::make_shared( getDpctNamespace() + - "experimental::granularity_flags::GRANULARITY_FLAGS_MINIMUM")}, + "experimental::granularity_flags::granularity_flags_minimum")}, // enum Driver Device Attribute {"CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR", std::make_shared("get_major_version", diff --git a/clang/runtime/dpct-rt/include/dpct/memory.hpp b/clang/runtime/dpct-rt/include/dpct/memory.hpp index b262abe82411..977046bdfd03 100644 --- a/clang/runtime/dpct-rt/include/dpct/memory.hpp +++ b/clang/runtime/dpct-rt/include/dpct/memory.hpp @@ -1421,27 +1421,27 @@ namespace experimental { typedef sycl::ext::oneapi::experimental::physical_mem *mem_handle; enum mem_location_type { - MEM_LOCATION_TYPE_INVALID = 0x0, - MEM_LOCATION_TYPE_DEVICE = 0x1, - MEM_LOCATION_TYPE_MAX = 0xFFFFFFFF + mem_location_type_invalid = 0x0, + mem_location_type_device = 0x1, + mem_location_type_max = 0xFFFFFFFF }; enum mem_allocation_type { - MEM_ALLOCATION_TYPE_INVALID = 0x0, - MEM_ALLOCATION_TYPE_DEFAULT = 0x1, - MEM_ALLOCATION_TYPE_MAX = 0xFFFFFFFF + mem_allocation_type_invalid = 0x0, + mem_allocation_type_default = 0x1, + mem_allocation_type_max = 0xFFFFFFFF }; enum granularity_flags { - GRANULARITY_FLAGS_MINIMUM = 0x0, - GRANULARITY_FLAGS_RECOMMENDED = 0x1 + granularity_flags_minimum = 0x0, + granularity_flags_recommended = 0x1 }; enum address_access_flags { - ADDRESS_ACCESS_FLAGS_NONE = 0x0, - ADDRESS_ACCESS_FLAGS_READ = 0x1, - ADDRESS_ACCESS_FLAGS_READ_WRITE = 0x3, - ADDRESS_ACCESS_FLAGS_MAX = 0xFFFFFFFF + address_access_flags_none = 0x0, + address_access_flags_read = 0x1, + address_access_flags_read_write = 0x3, + address_access_flags_max = 0xFFFFFFFF }; struct mem_location { @@ -1556,18 +1556,17 @@ void mem_set_access(device_ptr ptr, size_t size, const mem_access_desc *desc, for (size_t index = 0; index < count; index++) { sycl::ext::oneapi::experimental::address_access_mode mode; switch (desc[index].flags) { - case address_access_flags::ADDRESS_ACCESS_FLAGS_NONE: + case address_access_flags::address_access_flags_none: mode = sycl::ext::oneapi::experimental::address_access_mode::none; break; - case address_access_flags::ADDRESS_ACCESS_FLAGS_READ: + case address_access_flags::address_access_flags_read: mode = sycl::ext::oneapi::experimental::address_access_mode::read; break; - case address_access_flags::ADDRESS_ACCESS_FLAGS_READ_WRITE: + case address_access_flags::address_access_flags_read_write: mode = sycl::ext::oneapi::experimental::address_access_mode::read_write; break; default: - mode = sycl::ext::oneapi::experimental::address_access_mode::none; - break; + throw std::runtime_error("mem_set_access: invalid address access flags."); } sycl::ext::oneapi::experimental::set_access_mode( ptr, size, mode, @@ -1586,7 +1585,7 @@ void mem_get_allocation_granularity(size_t *granularity, const mem_prop *prop, auto &device = dpct::get_device(prop->location.id); *granularity = sycl::ext::oneapi::experimental::get_mem_granularity( device, device.get_context(), - option == granularity_flags::GRANULARITY_FLAGS_MINIMUM + option == granularity_flags::granularity_flags_minimum ? sycl::ext::oneapi::experimental::granularity_mode::minimum : sycl::ext::oneapi::experimental::granularity_mode::recommended); } @@ -1602,7 +1601,7 @@ void mem_get_reserve_granularity(size_t *granularity, const mem_prop *prop, auto &device = dpct::get_device(prop->location.id); *granularity = sycl::ext::oneapi::experimental::get_mem_granularity( device.get_context(), - option == granularity_flags::GRANULARITY_FLAGS_MINIMUM + option == granularity_flags::granularity_flags_minimum ? sycl::ext::oneapi::experimental::granularity_mode::minimum : sycl::ext::oneapi::experimental::granularity_mode::recommended); } diff --git a/clang/test/dpct/virtual_memory.cu b/clang/test/dpct/virtual_memory.cu index 74532cb428fb..6db44a3cd59d 100644 --- a/clang/test/dpct/virtual_memory.cu +++ b/clang/test/dpct/virtual_memory.cu @@ -13,15 +13,15 @@ int main() { cuCtxCreate(&context, 0, device); // CHECK: dpct::experimental::mem_prop prop = {}; -// CHECK: prop.type = dpct::experimental::mem_allocation_type::MEM_ALLOCATION_TYPE_DEFAULT; -// CHECK: prop.location.type = dpct::experimental::mem_location_type::MEM_LOCATION_TYPE_DEVICE; +// CHECK: prop.type = dpct::experimental::mem_allocation_type::mem_allocation_type_default; +// CHECK: prop.location.type = dpct::experimental::mem_location_type::mem_location_type_device; // CHECK: prop.location.id = device; CUmemAllocationProp prop = {}; prop.type = CU_MEM_ALLOCATION_TYPE_PINNED; - prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE; + prop.location.type = CU_mem_location_type_device; prop.location.id = device; size_t granularity; -// CHECK: dpct::experimental::mem_get_allocation_granularity(&granularity, &prop, dpct::experimental::granularity_flags::GRANULARITY_FLAGS_MINIMUM); +// CHECK: dpct::experimental::mem_get_allocation_granularity(&granularity, &prop, dpct::experimental::granularity_flags::granularity_flags_minimum); cuMemGetAllocationGranularity(&granularity, &prop, CU_MEM_ALLOC_GRANULARITY_MINIMUM); size_t POOL_SIZE = granularity; @@ -37,12 +37,12 @@ int main() { cuMemMap(reserved_addr, POOL_SIZE, 0, allocHandle, 0); // CHECK: dpct::experimental::mem_access_desc accessDesc = {}; -// CHECK: accessDesc.location.type = dpct::experimental::mem_location_type::MEM_LOCATION_TYPE_DEVICE; +// CHECK: accessDesc.location.type = dpct::experimental::mem_location_type::mem_location_type_device; // CHECK: accessDesc.location.id = device; -// CHECK: accessDesc.flags = dpct::experimental::address_access_flags::ADDRESS_ACCESS_FLAGS_READ_WRITE; +// CHECK: accessDesc.flags = dpct::experimental::address_access_flags::address_access_flags_read_write; // CHECK: dpct::experimental::mem_set_access(reserved_addr, POOL_SIZE, &accessDesc, 1); CUmemAccessDesc accessDesc = {}; - accessDesc.location.type = CU_MEM_LOCATION_TYPE_DEVICE; + accessDesc.location.type = CU_mem_location_type_device; accessDesc.location.id = device; accessDesc.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE; cuMemSetAccess(reserved_addr, POOL_SIZE, &accessDesc, 1); From b9b1e82ae991d30c263cde4c422f670247861b97 Mon Sep 17 00:00:00 2001 From: intwanghao Date: Thu, 10 Oct 2024 09:11:26 +0800 Subject: [PATCH 05/10] fix Signed-off-by: intwanghao --- clang/runtime/dpct-rt/include/dpct/memory.hpp | 33 ++++++++++--------- 1 file changed, 18 insertions(+), 15 deletions(-) diff --git a/clang/runtime/dpct-rt/include/dpct/memory.hpp b/clang/runtime/dpct-rt/include/dpct/memory.hpp index 977046bdfd03..cd33a7e29bbb 100644 --- a/clang/runtime/dpct-rt/include/dpct/memory.hpp +++ b/clang/runtime/dpct-rt/include/dpct/memory.hpp @@ -1473,8 +1473,9 @@ struct mem_access_desc { * @param [in] addr Base address of the memory to be reserved. * @param [in] flags Reserved for future use. */ -void mem_address_reserve(device_ptr *ptr, size_t size, size_t alignment, - device_ptr addr, unsigned long long flags) { +static inline void mem_address_reserve(device_ptr *ptr, size_t size, + size_t alignment, device_ptr addr, + unsigned long long flags) { *ptr = (device_ptr)sycl::ext::oneapi::experimental::reserve_virtual_mem( (uintptr_t)addr, size, dpct::get_current_device().get_context()); } @@ -1484,7 +1485,7 @@ void mem_address_reserve(device_ptr *ptr, size_t size, size_t alignment, * @param [in] ptr Pointer to the device memory region to free. * @param [in] size Size of the memory region in bytes. */ -void mem_address_free(device_ptr ptr, size_t size) { +static inline void mem_address_free(device_ptr ptr, size_t size) { sycl::ext::oneapi::experimental::free_virtual_mem( (uintptr_t)ptr, size, dpct::get_current_device().get_context()); } @@ -1499,8 +1500,8 @@ void mem_address_free(device_ptr ptr, size_t size) { * @param [in] prop Properties for the memory allocation. * @param [in] flags Reserved for future use. */ -void mem_create(mem_handle *handle, size_t size, const mem_prop *prop, - unsigned long long flags) { +static inline void mem_create(mem_handle *handle, size_t size, + const mem_prop *prop, unsigned long long flags) { auto &device = dpct::get_device(prop->location.id); *handle = new sycl::ext::oneapi::experimental::physical_mem( device, device.get_context(), size); @@ -1510,7 +1511,7 @@ void mem_create(mem_handle *handle, size_t size, const mem_prop *prop, * @brief Releases the physical memory object. * @param [in] handle Handle to the memory object to release. */ -void mem_release(mem_handle handle) { +static inline void mem_release(mem_handle handle) { if (handle) { delete handle; } @@ -1527,8 +1528,8 @@ void mem_release(mem_handle handle) { * @param [in] handle Handle to the physical memory. * @param [in] flags Reserved for future use. */ -void mem_map(device_ptr ptr, size_t size, size_t offset, mem_handle handle, - unsigned long long flags) { +static inline void mem_map(device_ptr ptr, size_t size, size_t offset, + mem_handle handle, unsigned long long flags) { handle->map((uintptr_t)ptr, size, sycl::ext::oneapi::experimental::address_access_mode::read_write, offset); @@ -1539,7 +1540,7 @@ void mem_map(device_ptr ptr, size_t size, size_t offset, mem_handle handle, * @param [in] ptr Pointer to the virtual memory region to unmap. * @param [in] size Size of the memory region in bytes. */ -void mem_unmap(device_ptr ptr, size_t size) { +static inline void mem_unmap(device_ptr ptr, size_t size) { sycl::ext::oneapi::experimental::unmap( ptr, size, dpct::get_current_device().get_context()); } @@ -1551,8 +1552,8 @@ void mem_unmap(device_ptr ptr, size_t size) { * @param [in] desc Array of access descriptors. * @param [in] count Number of access descriptors. */ -void mem_set_access(device_ptr ptr, size_t size, const mem_access_desc *desc, - size_t count) { +static inline void mem_set_access(device_ptr ptr, size_t size, + const mem_access_desc *desc, size_t count) { for (size_t index = 0; index < count; index++) { sycl::ext::oneapi::experimental::address_access_mode mode; switch (desc[index].flags) { @@ -1580,8 +1581,9 @@ void mem_set_access(device_ptr ptr, size_t size, const mem_access_desc *desc, * @param [in] prop Properties for the memory allocation. * @param [in] option Granularity option, either minimum or recommended. */ -void mem_get_allocation_granularity(size_t *granularity, const mem_prop *prop, - granularity_flags option) { +static inline void mem_get_allocation_granularity(size_t *granularity, + const mem_prop *prop, + granularity_flags option) { auto &device = dpct::get_device(prop->location.id); *granularity = sycl::ext::oneapi::experimental::get_mem_granularity( device, device.get_context(), @@ -1596,8 +1598,9 @@ void mem_get_allocation_granularity(size_t *granularity, const mem_prop *prop, * @param [in] prop Properties for the virtual memory region reserve. * @param [in] option Granularity option, either minimum or recommended. */ -void mem_get_reserve_granularity(size_t *granularity, const mem_prop *prop, - granularity_flags option) { +static inline void mem_get_reserve_granularity(size_t *granularity, + const mem_prop *prop, + granularity_flags option) { auto &device = dpct::get_device(prop->location.id); *granularity = sycl::ext::oneapi::experimental::get_mem_granularity( device.get_context(), From 332135679dc95f83197b3dee07d9151f338bcc92 Mon Sep 17 00:00:00 2001 From: intwanghao Date: Thu, 10 Oct 2024 10:41:32 +0800 Subject: [PATCH 06/10] fix Signed-off-by: intwanghao --- clang/test/dpct/virtual_memory.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/test/dpct/virtual_memory.cu b/clang/test/dpct/virtual_memory.cu index 6db44a3cd59d..735ae344bb7e 100644 --- a/clang/test/dpct/virtual_memory.cu +++ b/clang/test/dpct/virtual_memory.cu @@ -18,7 +18,7 @@ int main() { // CHECK: prop.location.id = device; CUmemAllocationProp prop = {}; prop.type = CU_MEM_ALLOCATION_TYPE_PINNED; - prop.location.type = CU_mem_location_type_device; + prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE; prop.location.id = device; size_t granularity; // CHECK: dpct::experimental::mem_get_allocation_granularity(&granularity, &prop, dpct::experimental::granularity_flags::granularity_flags_minimum); @@ -42,7 +42,7 @@ int main() { // CHECK: accessDesc.flags = dpct::experimental::address_access_flags::address_access_flags_read_write; // CHECK: dpct::experimental::mem_set_access(reserved_addr, POOL_SIZE, &accessDesc, 1); CUmemAccessDesc accessDesc = {}; - accessDesc.location.type = CU_mem_location_type_device; + accessDesc.location.type = CU_MEM_LOCATION_TYPE_DEVICE; accessDesc.location.id = device; accessDesc.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE; cuMemSetAccess(reserved_addr, POOL_SIZE, &accessDesc, 1); From 03939739bd5c1cfe9d734f5a0694618cc91d7a2a Mon Sep 17 00:00:00 2001 From: intwanghao Date: Thu, 10 Oct 2024 13:20:12 +0800 Subject: [PATCH 07/10] fix Signed-off-by: intwanghao --- clang/test/dpct/virtual_memory.cu | 2 ++ 1 file changed, 2 insertions(+) diff --git a/clang/test/dpct/virtual_memory.cu b/clang/test/dpct/virtual_memory.cu index 735ae344bb7e..8ec983bb127d 100644 --- a/clang/test/dpct/virtual_memory.cu +++ b/clang/test/dpct/virtual_memory.cu @@ -1,3 +1,5 @@ +// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2 +// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2 // RUN: dpct --format-range=none --use-experimental-features=virtual_memory -out-root %T/virtual_memory %s --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only // RUN: FileCheck %s --match-full-lines --input-file %T/virtual_memory/virtual_memory.dp.cpp #include From 627ed0d0cc8930b1b4dc7ee14e483a99983325bb Mon Sep 17 00:00:00 2001 From: intwanghao Date: Thu, 10 Oct 2024 14:35:25 +0800 Subject: [PATCH 08/10] fix Signed-off-by: intwanghao --- clang/lib/DPCT/ASTTraversal.cpp | 160 +++++++++++++++++--------------- clang/lib/DPCT/ASTTraversal.h | 6 ++ 2 files changed, 90 insertions(+), 76 deletions(-) diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index cc67e15c71d8..001d5fb53781 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/clang/lib/DPCT/ASTTraversal.cpp @@ -10851,30 +10851,6 @@ void MemoryMigrationRule::registerMatcher(MatchFinder &MF) { hasParent(callExpr(parentStmt()).bind("callExpr"))) .bind("unresolvedCall"), this); - auto virtualmemoryAPI = [&]() { - return hasAnyName("cuMemCreate", "cuMemAddressReserve", "cuMemMap", - "cuMemUnmap", "cuMemAddressFree", "cuMemRelease", - "cuMemSetAccess", "cuMemGetAllocationGranularity"); - }; - auto virtualmemoryType = [&]() { - return hasAnyName("CUmemAllocationProp", "CUmemGenericAllocationHandle", - "CUmemAccessDesc"); - }; - auto virtualmemoryEnum = [&]() { - return hasAnyName("CU_MEM_ALLOCATION_TYPE_PINNED", - "CU_MEM_LOCATION_TYPE_DEVICE", - "CU_MEM_ACCESS_FLAGS_PROT_READWRITE", - "CU_MEM_ALLOC_GRANULARITY_RECOMMENDED"); - }; - MF.addMatcher( - callExpr(callee(functionDecl(virtualmemoryAPI()))).bind("vmCall"), this); - MF.addMatcher( - typeLoc(loc(qualType(hasDeclaration(namedDecl(virtualmemoryType()))))) - .bind("vmType"), - this); - MF.addMatcher( - declRefExpr(to(enumConstantDecl(virtualmemoryEnum()))).bind("vmEnum"), - this); } void MemoryMigrationRule::runRule(const MatchFinder::MatchResult &Result) { @@ -10977,58 +10953,6 @@ void MemoryMigrationRule::runRule(const MatchFinder::MatchResult &Result) { getAssistNodeAsType(Result, "callExpr"), /* IsAssigned */ false, getAssistNodeAsType(Result, "unresolvedCall")); - - auto &SM = DpctGlobalInfo::getSourceManager(); - if (const CallExpr *CE = getNodeAsType(Result, "vmCall")) { - ExprAnalysis EA(CE); - emplaceTransformation(EA.getReplacement()); - EA.applyAllSubExprRepl(); - } - if (auto TL = getNodeAsType(Result, "vmType")) { - auto TypeStr = - DpctGlobalInfo::getTypeName(TL->getType().getUnqualifiedType()); - if (!DpctGlobalInfo::useExpVirtualMemory()) { - report(TL->getBeginLoc(), Diagnostics::TRY_EXPERIMENTAL_FEATURE, false, - TypeStr, "--use-experimental-features=virtual_memory"); - return; - } - if (!DpctGlobalInfo::isInAnalysisScope( - SM.getSpellingLoc(TL->getBeginLoc()))) { - return; - } - auto Range = getDefinitionRange(TL->getBeginLoc(), TL->getEndLoc()); - auto BeginLoc = Range.getBegin(); - auto EndLoc = Range.getEnd(); - - if (SM.isWrittenInScratchSpace(SM.getSpellingLoc(TL->getBeginLoc()))) { - BeginLoc = SM.getExpansionRange(TL->getBeginLoc()).getBegin(); - EndLoc = SM.getExpansionRange(TL->getBeginLoc()).getEnd(); - } - std::string Str = - MapNames::findReplacedName(MapNames::TypeNamesMap, TypeStr); - if (!Str.empty()) { - auto Len = Lexer::MeasureTokenLength( - EndLoc, SM, DpctGlobalInfo::getContext().getLangOpts()); - Len += SM.getDecomposedLoc(EndLoc).second - - SM.getDecomposedLoc(BeginLoc).second; - emplaceTransformation(new ReplaceText(BeginLoc, Len, std::move(Str))); - return; - } - } - if (auto *E = getNodeAsType(Result, "vmEnum")) { - std::string EnumName = E->getNameInfo().getName().getAsString(); - if (!DpctGlobalInfo::useExpVirtualMemory()) { - report(E->getBeginLoc(), Diagnostics::TRY_EXPERIMENTAL_FEATURE, false, - EnumName, "--use-experimental-features=virtual_memory"); - return; - } - auto Search = EnumConstantRule::EnumNamesMap.find(EnumName); - if (Search == EnumConstantRule::EnumNamesMap.end()) { - report(E->getBeginLoc(), Diagnostics::API_NOT_MIGRATED, false, EnumName); - return; - } - emplaceTransformation(new ReplaceStmt(E, Search->second->NewName)); - } } void MemoryMigrationRule::getSymbolAddressMigration( @@ -14021,6 +13945,90 @@ void FFTFunctionCallRule::runRule(const MatchFinder::MatchResult &Result) { REGISTER_RULE(FFTFunctionCallRule, PassKind::PK_Migration, RuleGroupKind::RK_FFT) +void VirtualMemoryRule::registerMatcher(ast_matchers::MatchFinder &MF) { + auto virtualmemoryAPI = [&]() { + return hasAnyName("cuMemCreate", "cuMemAddressReserve", "cuMemMap", + "cuMemUnmap", "cuMemAddressFree", "cuMemRelease", + "cuMemSetAccess", "cuMemGetAllocationGranularity"); + }; + auto virtualmemoryType = [&]() { + return hasAnyName("CUmemAllocationProp", "CUmemGenericAllocationHandle", + "CUmemAccessDesc"); + }; + auto virtualmemoryEnum = [&]() { + return hasAnyName("CU_MEM_ALLOCATION_TYPE_PINNED", + "CU_MEM_LOCATION_TYPE_DEVICE", + "CU_MEM_ACCESS_FLAGS_PROT_READWRITE", + "CU_MEM_ALLOC_GRANULARITY_RECOMMENDED"); + }; + MF.addMatcher( + callExpr(callee(functionDecl(virtualmemoryAPI()))).bind("vmCall"), this); + MF.addMatcher( + typeLoc(loc(qualType(hasDeclaration(namedDecl(virtualmemoryType()))))) + .bind("vmType"), + this); + MF.addMatcher( + declRefExpr(to(enumConstantDecl(virtualmemoryEnum()))).bind("vmEnum"), + this); +} + +void VirtualMemoryRule::runRule( + const ast_matchers::MatchFinder::MatchResult &Result) { + auto &SM = DpctGlobalInfo::getSourceManager(); + if (const CallExpr *CE = getNodeAsType(Result, "vmCall")) { + ExprAnalysis EA(CE); + emplaceTransformation(EA.getReplacement()); + EA.applyAllSubExprRepl(); + } + if (auto TL = getNodeAsType(Result, "vmType")) { + auto TypeStr = + DpctGlobalInfo::getTypeName(TL->getType().getUnqualifiedType()); + if (!DpctGlobalInfo::useExpVirtualMemory()) { + report(TL->getBeginLoc(), Diagnostics::TRY_EXPERIMENTAL_FEATURE, false, + TypeStr, "--use-experimental-features=virtual_memory"); + return; + } + if (!DpctGlobalInfo::isInAnalysisScope( + SM.getSpellingLoc(TL->getBeginLoc()))) { + return; + } + auto Range = getDefinitionRange(TL->getBeginLoc(), TL->getEndLoc()); + auto BeginLoc = Range.getBegin(); + auto EndLoc = Range.getEnd(); + + if (SM.isWrittenInScratchSpace(SM.getSpellingLoc(TL->getBeginLoc()))) { + BeginLoc = SM.getExpansionRange(TL->getBeginLoc()).getBegin(); + EndLoc = SM.getExpansionRange(TL->getBeginLoc()).getEnd(); + } + std::string Str = + MapNames::findReplacedName(MapNames::TypeNamesMap, TypeStr); + if (!Str.empty()) { + auto Len = Lexer::MeasureTokenLength( + EndLoc, SM, DpctGlobalInfo::getContext().getLangOpts()); + Len += SM.getDecomposedLoc(EndLoc).second - + SM.getDecomposedLoc(BeginLoc).second; + emplaceTransformation(new ReplaceText(BeginLoc, Len, std::move(Str))); + return; + } + } + if (auto *E = getNodeAsType(Result, "vmEnum")) { + std::string EnumName = E->getNameInfo().getName().getAsString(); + if (!DpctGlobalInfo::useExpVirtualMemory()) { + report(E->getBeginLoc(), Diagnostics::TRY_EXPERIMENTAL_FEATURE, false, + EnumName, "--use-experimental-features=virtual_memory"); + return; + } + auto Search = EnumConstantRule::EnumNamesMap.find(EnumName); + if (Search == EnumConstantRule::EnumNamesMap.end()) { + report(E->getBeginLoc(), Diagnostics::API_NOT_MIGRATED, false, EnumName); + return; + } + emplaceTransformation(new ReplaceStmt(E, Search->second->NewName)); + } +} + +REGISTER_RULE(VirtualMemoryRule, PassKind::PK_Migration) + void DriverModuleAPIRule::registerMatcher(ast_matchers::MatchFinder &MF) { auto DriverModuleAPI = [&]() { return hasAnyName("cuModuleLoad", "cuModuleLoadData", "cuModuleLoadDataEx", diff --git a/clang/lib/DPCT/ASTTraversal.h b/clang/lib/DPCT/ASTTraversal.h index 2424405f04fb..936eeb2a022b 100644 --- a/clang/lib/DPCT/ASTTraversal.h +++ b/clang/lib/DPCT/ASTTraversal.h @@ -1390,6 +1390,12 @@ class DriverModuleAPIRule : public NamedMigrationRule { void runRule(const ast_matchers::MatchFinder::MatchResult &Result); }; +class VirtualMemoryRule : public NamedMigrationRule { +public: + void registerMatcher(ast_matchers::MatchFinder &MF) override; + void runRule(const ast_matchers::MatchFinder::MatchResult &Result); +}; + class DriverDeviceAPIRule : public NamedMigrationRule { public: void registerMatcher(ast_matchers::MatchFinder &MF) override; From 4c470822603637dbd85e198771f3a9178a643c1f Mon Sep 17 00:00:00 2001 From: intwanghao Date: Mon, 14 Oct 2024 17:53:37 +0800 Subject: [PATCH 09/10] fix Signed-off-by: intwanghao --- clang/include/clang/DPCT/DPCTOptions.inc | 2 +- clang/lib/DPCT/APINamesMemory.inc | 96 ++++++---- clang/lib/DPCT/ASTTraversal.cpp | 22 ++- clang/lib/DPCT/ASTTraversal.h | 2 +- clang/lib/DPCT/MapNames.cpp | 63 +++--- clang/runtime/dpct-rt/include/dpct/memory.hpp | 180 +----------------- clang/test/dpct/virtual_memory.cu | 36 ++-- 7 files changed, 128 insertions(+), 273 deletions(-) diff --git a/clang/include/clang/DPCT/DPCTOptions.inc b/clang/include/clang/DPCT/DPCTOptions.inc index 897e40b6235c..aaf7d9231b42 100644 --- a/clang/include/clang/DPCT/DPCTOptions.inc +++ b/clang/include/clang/DPCT/DPCTOptions.inc @@ -788,7 +788,7 @@ DPCT_ENUM_OPTION( "variables.\n", false), DPCT_OPTION_ENUM_VALUE( - "virtual_memory", int(ExperimentalFeatures::Exp_VirtualMemory), + "virtual_mem", int(ExperimentalFeatures::Exp_VirtualMemory), "Experimental extension that allows map an address range onto " "multiple allocations of physical memory.", false), diff --git a/clang/lib/DPCT/APINamesMemory.inc b/clang/lib/DPCT/APINamesMemory.inc index 6a9e8003fa46..fce6b16aef6b 100644 --- a/clang/lib/DPCT/APINamesMemory.inc +++ b/clang/lib/DPCT/APINamesMemory.inc @@ -702,10 +702,17 @@ CONDITIONAL_FACTORY_ENTRY( UseExpVirtualMemory, ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY( HelperFeatureEnum::device_ext, - CALL_FACTORY_ENTRY("cuMemCreate", - CALL(MapNames::getDpctNamespace() + - "experimental::mem_create", - ARG(0), ARG(1), ARG(2), ARG(3))))), + ASSIGN_FACTORY_ENTRY( + "cuMemCreate", DEREF(0), + NEW(MapNames::getClNamespace() + + "ext::oneapi::experimental::physical_mem", + CALL(MapNames::getDpctNamespace() + "get_device", + MEMBER_EXPR(DEREF(2), false, LITERAL("location.id"))), + MEMBER_CALL(CALL(MapNames::getDpctNamespace() + "get_device", + MEMBER_EXPR(DEREF(2), false, + LITERAL("location.id"))), + false, "get_context"), + ARG(1))))), UNSUPPORT_FACTORY_ENTRY("cuMemCreate", Diagnostics::TRY_EXPERIMENTAL_FEATURE, ARG("cuMemCreate"), @@ -715,10 +722,14 @@ CONDITIONAL_FACTORY_ENTRY( UseExpVirtualMemory, ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY( HelperFeatureEnum::device_ext, - CALL_FACTORY_ENTRY("cuMemAddressReserve", - CALL(MapNames::getDpctNamespace() + - "experimental::mem_address_reserve", - ARG(0), ARG(1), ARG(2), ARG(3), ARG(4))))), + ASSIGN_FACTORY_ENTRY( + "cuMemAddressReserve", DEREF(0), + CALL("(" + MapNames::getDpctNamespace() + "device_ptr)" + + MapNames::getClNamespace() + + "ext::oneapi::experimental::reserve_virtual_mem", + CAST(LITERAL("uintptr_t"), ARG(3)), ARG(1), + LITERAL(MapNames::getDpctNamespace() + + "get_current_device().get_context()"))))), UNSUPPORT_FACTORY_ENTRY("cuMemAddressReserve", Diagnostics::TRY_EXPERIMENTAL_FEATURE, ARG("cuMemAddressReserve"), @@ -728,10 +739,13 @@ CONDITIONAL_FACTORY_ENTRY( UseExpVirtualMemory, ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY( HelperFeatureEnum::device_ext, - CALL_FACTORY_ENTRY("cuMemAddressFree", - CALL(MapNames::getDpctNamespace() + - "experimental::mem_address_free", - ARG(0), ARG(1))))), + CALL_FACTORY_ENTRY( + "cuMemAddressFree", + CALL(MapNames::getClNamespace() + + "ext::oneapi::experimental::free_virtual_mem", + CAST(LITERAL("uintptr_t"), ARG(0)), ARG(1), + LITERAL(MapNames::getDpctNamespace() + + "get_current_device().get_context()"))))), UNSUPPORT_FACTORY_ENTRY("cuMemAddressFree", Diagnostics::TRY_EXPERIMENTAL_FEATURE, ARG("cuMemAddressFree"), @@ -741,11 +755,17 @@ CONDITIONAL_FACTORY_ENTRY( UseExpVirtualMemory, ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY( HelperFeatureEnum::device_ext, - CALL_FACTORY_ENTRY( - "cuMemGetAllocationGranularity", - CALL(MapNames::getDpctNamespace() + - "experimental::mem_get_allocation_granularity", - ARG(0), ARG(1), ARG(2))))), + ASSIGN_FACTORY_ENTRY( + "cuMemGetAllocationGranularity", DEREF(0), + CALL(MapNames::getClNamespace() + + "ext::oneapi::experimental::get_mem_granularity", + CALL(MapNames::getDpctNamespace() + "get_device", + MEMBER_EXPR(DEREF(1), false, LITERAL("location.id"))), + MEMBER_CALL(CALL(MapNames::getDpctNamespace() + "get_device", + MEMBER_EXPR(DEREF(1), false, + LITERAL("location.id"))), + false, "get_context"), + ARG(2))))), UNSUPPORT_FACTORY_ENTRY("cuMemGetAllocationGranularity", Diagnostics::TRY_EXPERIMENTAL_FEATURE, ARG("cuMemGetAllocationGranularity"), @@ -753,11 +773,9 @@ CONDITIONAL_FACTORY_ENTRY( CONDITIONAL_FACTORY_ENTRY( UseExpVirtualMemory, - ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY( - HelperFeatureEnum::device_ext, - CALL_FACTORY_ENTRY("cuMemRelease", CALL(MapNames::getDpctNamespace() + - "experimental::mem_release", - ARG(0))))), + ASSIGNABLE_FACTORY( + FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, + DELETE_FACTORY_ENTRY("cuMemRelease", ARG(0)))), UNSUPPORT_FACTORY_ENTRY("cuMemRelease", Diagnostics::TRY_EXPERIMENTAL_FEATURE, ARG("cuMemRelease"), @@ -767,10 +785,13 @@ CONDITIONAL_FACTORY_ENTRY( UseExpVirtualMemory, ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY( HelperFeatureEnum::device_ext, - CALL_FACTORY_ENTRY("cuMemMap", - CALL(MapNames::getDpctNamespace() + - "experimental::mem_map", - ARG(0), ARG(1), ARG(2), ARG(3), ARG(4))))), + MEMBER_CALL_FACTORY_ENTRY( + "cuMemMap", ARG(3), true, "map", CAST(LITERAL("uintptr_t"), ARG(0)), + ARG(1), + LITERAL( + MapNames::getClNamespace() + + "ext::oneapi::experimental::address_access_mode::read_write"), + ARG(2)))), UNSUPPORT_FACTORY_ENTRY("cuMemMap", Diagnostics::TRY_EXPERIMENTAL_FEATURE, ARG("cuMemMap"), ARG("--use-experimental-features=virtual_memory"))) @@ -779,9 +800,13 @@ CONDITIONAL_FACTORY_ENTRY( UseExpVirtualMemory, ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY( HelperFeatureEnum::device_ext, - CALL_FACTORY_ENTRY("cuMemUnmap", CALL(MapNames::getDpctNamespace() + - "experimental::mem_unmap", - ARG(0), ARG(1))))), + CALL_FACTORY_ENTRY( + "cuMemUnmap", + CALL(MapNames::getClNamespace() + + "ext::oneapi::experimental::unmap", + ARG(0), ARG(1), + LITERAL(MapNames::getDpctNamespace() + + "get_current_device().get_context()"))))), UNSUPPORT_FACTORY_ENTRY("cuMemUnmap", Diagnostics::TRY_EXPERIMENTAL_FEATURE, ARG("cuMemUnmap"), ARG("--use-experimental-features=virtual_memory"))) @@ -790,10 +815,15 @@ CONDITIONAL_FACTORY_ENTRY( UseExpVirtualMemory, ASSIGNABLE_FACTORY(FEATURE_REQUEST_FACTORY( HelperFeatureEnum::device_ext, - CALL_FACTORY_ENTRY("cuMemSetAccess", - CALL(MapNames::getDpctNamespace() + - "experimental::mem_set_access", - ARG(0), ARG(1), ARG(2), ARG(3))))), + CALL_FACTORY_ENTRY( + "cuMemSetAccess", + CALL(MapNames::getClNamespace() + + "ext::oneapi::experimental::set_access_mode", + ARG(0), ARG(1), MEMBER_EXPR(DEREF(2), false, LITERAL("flags")), + MEMBER_CALL(CALL(MapNames::getDpctNamespace() + "get_device", + MEMBER_EXPR(DEREF(2), false, + LITERAL("location.id"))), + false, "get_context"))))), UNSUPPORT_FACTORY_ENTRY("cuMemSetAccess", Diagnostics::TRY_EXPERIMENTAL_FEATURE, ARG("cuMemSetAccess"), diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index 6d8523cacbf9..a15b4608af18 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/clang/lib/DPCT/ASTTraversal.cpp @@ -13946,7 +13946,7 @@ void FFTFunctionCallRule::runRule(const MatchFinder::MatchResult &Result) { REGISTER_RULE(FFTFunctionCallRule, PassKind::PK_Migration, RuleGroupKind::RK_FFT) -void VirtualMemoryRule::registerMatcher(ast_matchers::MatchFinder &MF) { +void VirtualMemRule::registerMatcher(ast_matchers::MatchFinder &MF) { auto virtualmemoryAPI = [&]() { return hasAnyName("cuMemCreate", "cuMemAddressReserve", "cuMemMap", "cuMemUnmap", "cuMemAddressFree", "cuMemRelease", @@ -13954,13 +13954,19 @@ void VirtualMemoryRule::registerMatcher(ast_matchers::MatchFinder &MF) { }; auto virtualmemoryType = [&]() { return hasAnyName("CUmemAllocationProp", "CUmemGenericAllocationHandle", - "CUmemAccessDesc"); + "CUmemAccessDesc", "CUmemLocationType", + "CUmemAllocationType", "CUmemAllocationGranularity_flags", + "CUmemAccess_flags"); }; auto virtualmemoryEnum = [&]() { - return hasAnyName("CU_MEM_ALLOCATION_TYPE_PINNED", - "CU_MEM_LOCATION_TYPE_DEVICE", - "CU_MEM_ACCESS_FLAGS_PROT_READWRITE", - "CU_MEM_ALLOC_GRANULARITY_RECOMMENDED"); + return hasAnyName( + "CU_MEM_ALLOCATION_TYPE_PINNED", "CU_MEM_ALLOCATION_TYPE_INVALID", + "CU_MEM_ALLOCATION_TYPE_MAX", "CU_MEM_LOCATION_TYPE_DEVICE", + "CU_MEM_LOCATION_TYPE_INVALID", "CU_MEM_LOCATION_TYPE_MAX", + "CU_MEM_ACCESS_FLAGS_PROT_NONE", "CU_MEM_ACCESS_FLAGS_PROT_READ", + "CU_MEM_ACCESS_FLAGS_PROT_READWRITE", + "CU_MEM_ALLOC_GRANULARITY_RECOMMENDED", + "CU_MEM_ALLOC_GRANULARITY_MINIMUM"); }; MF.addMatcher( callExpr(callee(functionDecl(virtualmemoryAPI()))).bind("vmCall"), this); @@ -13973,7 +13979,7 @@ void VirtualMemoryRule::registerMatcher(ast_matchers::MatchFinder &MF) { this); } -void VirtualMemoryRule::runRule( +void VirtualMemRule::runRule( const ast_matchers::MatchFinder::MatchResult &Result) { auto &SM = DpctGlobalInfo::getSourceManager(); if (const CallExpr *CE = getNodeAsType(Result, "vmCall")) { @@ -14028,7 +14034,7 @@ void VirtualMemoryRule::runRule( } } -REGISTER_RULE(VirtualMemoryRule, PassKind::PK_Migration) +REGISTER_RULE(VirtualMemRule, PassKind::PK_Migration) void DriverModuleAPIRule::registerMatcher(ast_matchers::MatchFinder &MF) { auto DriverModuleAPI = [&]() { diff --git a/clang/lib/DPCT/ASTTraversal.h b/clang/lib/DPCT/ASTTraversal.h index 936eeb2a022b..195c7ac360b1 100644 --- a/clang/lib/DPCT/ASTTraversal.h +++ b/clang/lib/DPCT/ASTTraversal.h @@ -1390,7 +1390,7 @@ class DriverModuleAPIRule : public NamedMigrationRule { void runRule(const ast_matchers::MatchFinder::MatchResult &Result); }; -class VirtualMemoryRule : public NamedMigrationRule { +class VirtualMemRule : public NamedMigrationRule { public: void registerMatcher(ast_matchers::MatchFinder &MF) override; void runRule(const ast_matchers::MatchFinder::MatchResult &Result); diff --git a/clang/lib/DPCT/MapNames.cpp b/clang/lib/DPCT/MapNames.cpp index bbbb04fd0003..aab9d64b5d87 100644 --- a/clang/lib/DPCT/MapNames.cpp +++ b/clang/lib/DPCT/MapNames.cpp @@ -842,6 +842,15 @@ void MapNames::setExplicitNamespaceMap( {"CUmemAccessDesc", std::make_shared(getDpctNamespace() + "experimental::mem_access_desc")}, + {"CUmemLocationType", std::make_shared("int")}, + {"CUmemAllocationType", std::make_shared("int")}, + {"CUmemAllocationGranularity_flags", + std::make_shared( + getClNamespace() + "ext::oneapi::experimental::granularity_mode")}, + {"CUmemAccess_flags", + std::make_shared( + getClNamespace() + + "ext::oneapi::experimental::address_access_mode")}, {"cudaGraphicsMapFlags", std::make_shared("int")}, {"cudaGraphicsRegisterFlags", std::make_shared("int")}, // ... @@ -1449,54 +1458,38 @@ void MapNames::setExplicitNamespaceMap( std::make_shared("0")}, {"CU_MEM_ADVISE_SET_ACCESSED_BY", std::make_shared("0")}, {"CU_MEM_ADVISE_UNSET_ACCESSED_BY", std::make_shared("0")}, - {"CU_MEM_ALLOCATION_TYPE_PINNED", - std::make_shared( - getDpctNamespace() + - "experimental::mem_allocation_type::mem_allocation_type_default")}, - {"CU_MEM_ALLOCATION_TYPE_INVALID", - std::make_shared( - getDpctNamespace() + - "experimental::mem_allocation_type::mem_allocation_type_invalid")}, + {"CU_MEM_ALLOCATION_TYPE_PINNED", std::make_shared("0")}, + {"CU_MEM_ALLOCATION_TYPE_INVALID", std::make_shared("1")}, {"CU_MEM_ALLOCATION_TYPE_MAX", - std::make_shared( - getDpctNamespace() + - "experimental::mem_allocation_type::mem_allocation_type_max")}, - {"CU_MEM_LOCATION_TYPE_DEVICE", - std::make_shared( - getDpctNamespace() + - "experimental::mem_location_type::mem_location_type_device")}, - {"CU_MEM_LOCATION_TYPE_INVALID", - std::make_shared( - getDpctNamespace() + - "experimental::mem_location_type::mem_location_type_invalid")}, + std::make_shared("0xFFFFFFFF")}, + {"CU_MEM_LOCATION_TYPE_DEVICE", std::make_shared("1")}, + {"CU_MEM_LOCATION_TYPE_INVALID", std::make_shared("0")}, {"CU_MEM_LOCATION_TYPE_MAX", - std::make_shared( - getDpctNamespace() + - "experimental::mem_location_type::mem_location_type_max")}, + std::make_shared("0xFFFFFFFF")}, {"CU_MEM_ACCESS_FLAGS_PROT_READWRITE", - std::make_shared(getDpctNamespace() + - "experimental::address_access_flags::" - "address_access_flags_read_write")}, + std::make_shared( + getClNamespace() + + "ext::oneapi::experimental::address_access_mode::read_write")}, {"CU_MEM_ACCESS_FLAGS_PROT_NONE", std::make_shared( - getDpctNamespace() + - "experimental::address_access_flags::address_access_flags_none")}, + getClNamespace() + + "ext::oneapi::experimental::address_access_mode::none")}, {"CU_MEM_ACCESS_FLAGS_PROT_MAX", std::make_shared( - getDpctNamespace() + - "experimental::address_access_flags::address_access_flags_max")}, + getClNamespace() + + "ext::oneapi::experimental::address_access_mode::none")}, {"CU_MEM_ACCESS_FLAGS_PROT_READ", std::make_shared( - getDpctNamespace() + - "experimental::address_access_flags::address_access_flags_read")}, + getClNamespace() + + "ext::oneapi::experimental::address_access_mode::read")}, {"CU_MEM_ALLOC_GRANULARITY_RECOMMENDED", std::make_shared( - getDpctNamespace() + - "experimental::granularity_flags::granularity_flags_recommended")}, + getClNamespace() + + "ext::oneapi::experimental::granularity_mode::recommended")}, {"CU_MEM_ALLOC_GRANULARITY_MINIMUM", std::make_shared( - getDpctNamespace() + - "experimental::granularity_flags::granularity_flags_minimum")}, + getClNamespace() + + "ext::oneapi::experimental::granularity_mode::minimum")}, // enum Driver Device Attribute {"CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR", std::make_shared("get_major_version", diff --git a/clang/runtime/dpct-rt/include/dpct/memory.hpp b/clang/runtime/dpct-rt/include/dpct/memory.hpp index cd33a7e29bbb..741fd90210f1 100644 --- a/clang/runtime/dpct-rt/include/dpct/memory.hpp +++ b/clang/runtime/dpct-rt/include/dpct/memory.hpp @@ -1420,194 +1420,20 @@ static inline void async_dpct_memset(pitched_data pitch, int val, namespace experimental { typedef sycl::ext::oneapi::experimental::physical_mem *mem_handle; -enum mem_location_type { - mem_location_type_invalid = 0x0, - mem_location_type_device = 0x1, - mem_location_type_max = 0xFFFFFFFF -}; - -enum mem_allocation_type { - mem_allocation_type_invalid = 0x0, - mem_allocation_type_default = 0x1, - mem_allocation_type_max = 0xFFFFFFFF -}; - -enum granularity_flags { - granularity_flags_minimum = 0x0, - granularity_flags_recommended = 0x1 -}; - -enum address_access_flags { - address_access_flags_none = 0x0, - address_access_flags_read = 0x1, - address_access_flags_read_write = 0x3, - address_access_flags_max = 0xFFFFFFFF -}; - struct mem_location { int id; - mem_location_type type; + int type; }; struct mem_prop { mem_location location; - mem_allocation_type type; + int type; }; struct mem_access_desc { - address_access_flags flags; + sycl::ext::oneapi::experimental::address_access_mode flags; mem_location location; }; - -/** - * @brief Reserves a virtual memory region of \p size bytes. The \p addr - * specifies the requested start of the new virtual memory range reservation. If - * the implementation is unable to reserve the virtual memory range at the - * specified address, the implementation will pick another suitable address. The - * \p size must be a multiple of the reserve granularity. The \p alignment - * parameter must be set to 0, which means the default alignment will be used. - * The \p addr must be aligned in accordance with the reserve granularity. - * @param [out] ptr Pointer to the reserved virtual memory. - * @param [in] size Size of the virtual memory region in bytes. - * @param [in] alignment Alignment of the virtual memory (must be 0). - * @param [in] addr Base address of the memory to be reserved. - * @param [in] flags Reserved for future use. - */ -static inline void mem_address_reserve(device_ptr *ptr, size_t size, - size_t alignment, device_ptr addr, - unsigned long long flags) { - *ptr = (device_ptr)sycl::ext::oneapi::experimental::reserve_virtual_mem( - (uintptr_t)addr, size, dpct::get_current_device().get_context()); -} - -/** - * @brief Frees a previously reserved virtual memory region pointed by \p ptr. - * @param [in] ptr Pointer to the device memory region to free. - * @param [in] size Size of the memory region in bytes. - */ -static inline void mem_address_free(device_ptr ptr, size_t size) { - sycl::ext::oneapi::experimental::free_virtual_mem( - (uintptr_t)ptr, size, dpct::get_current_device().get_context()); -} - -/** - * @brief Creates a physical memory object for allocation. - * This will allocate \p size of physical memory on the device. The \p size must - * be a multiple of the allocation granularity, as returned by a call to - * get_mem_allocation_granularity. - * @param [out] handle Handle to the created memory object. - * @param [in] size Size of the physical memory in bytes. - * @param [in] prop Properties for the memory allocation. - * @param [in] flags Reserved for future use. - */ -static inline void mem_create(mem_handle *handle, size_t size, - const mem_prop *prop, unsigned long long flags) { - auto &device = dpct::get_device(prop->location.id); - *handle = new sycl::ext::oneapi::experimental::physical_mem( - device, device.get_context(), size); -} - -/** - * @brief Releases the physical memory object. - * @param [in] handle Handle to the memory object to release. - */ -static inline void mem_release(mem_handle handle) { - if (handle) { - delete handle; - } -} - -/** - * @brief Maps a virtual memory range, specified by \p ptr and \p size, to the - * physical memory specified by \p handle, starting at an offset of \p offset - * bytes. The \p ptr, \p size and \p offset are all need to be multiples of the - * allocation granularity. - * @param [in] ptr Pointer to the virtual memory region. - * @param [in] size Size of the memory region in bytes. - * @param [in] offset Offset into the physical memory. - * @param [in] handle Handle to the physical memory. - * @param [in] flags Reserved for future use. - */ -static inline void mem_map(device_ptr ptr, size_t size, size_t offset, - mem_handle handle, unsigned long long flags) { - handle->map((uintptr_t)ptr, size, - sycl::ext::oneapi::experimental::address_access_mode::read_write, - offset); -} - -/** - * @brief Unmaps a previously mapped virtual memory region. - * @param [in] ptr Pointer to the virtual memory region to unmap. - * @param [in] size Size of the memory region in bytes. - */ -static inline void mem_unmap(device_ptr ptr, size_t size) { - sycl::ext::oneapi::experimental::unmap( - ptr, size, dpct::get_current_device().get_context()); -} - -/** - * @brief Sets the access mode of a virtual memory region pointed by \p ptr. - * @param [in] ptr Pointer to the virtual memory region. - * @param [in] size Size of the memory region in bytes. - * @param [in] desc Array of access descriptors. - * @param [in] count Number of access descriptors. - */ -static inline void mem_set_access(device_ptr ptr, size_t size, - const mem_access_desc *desc, size_t count) { - for (size_t index = 0; index < count; index++) { - sycl::ext::oneapi::experimental::address_access_mode mode; - switch (desc[index].flags) { - case address_access_flags::address_access_flags_none: - mode = sycl::ext::oneapi::experimental::address_access_mode::none; - break; - case address_access_flags::address_access_flags_read: - mode = sycl::ext::oneapi::experimental::address_access_mode::read; - break; - case address_access_flags::address_access_flags_read_write: - mode = sycl::ext::oneapi::experimental::address_access_mode::read_write; - break; - default: - throw std::runtime_error("mem_set_access: invalid address access flags."); - } - sycl::ext::oneapi::experimental::set_access_mode( - ptr, size, mode, - dpct::get_device(desc[index].location.id).get_context()); - } -} - -/** - * @brief Retrieves the granularity of the memory allocation. - * @param [out] granularity Pointer to store the retrieved granularity. - * @param [in] prop Properties for the memory allocation. - * @param [in] option Granularity option, either minimum or recommended. - */ -static inline void mem_get_allocation_granularity(size_t *granularity, - const mem_prop *prop, - granularity_flags option) { - auto &device = dpct::get_device(prop->location.id); - *granularity = sycl::ext::oneapi::experimental::get_mem_granularity( - device, device.get_context(), - option == granularity_flags::granularity_flags_minimum - ? sycl::ext::oneapi::experimental::granularity_mode::minimum - : sycl::ext::oneapi::experimental::granularity_mode::recommended); -} - -/** - * @brief Retrieves the granularity of the virtual memory region reserve. - * @param [out] granularity Pointer to store the retrieved granularity. - * @param [in] prop Properties for the virtual memory region reserve. - * @param [in] option Granularity option, either minimum or recommended. - */ -static inline void mem_get_reserve_granularity(size_t *granularity, - const mem_prop *prop, - granularity_flags option) { - auto &device = dpct::get_device(prop->location.id); - *granularity = sycl::ext::oneapi::experimental::get_mem_granularity( - device.get_context(), - option == granularity_flags::granularity_flags_minimum - ? sycl::ext::oneapi::experimental::granularity_mode::minimum - : sycl::ext::oneapi::experimental::granularity_mode::recommended); -} } // namespace experimental /// dpct accessor used as device function parameter. diff --git a/clang/test/dpct/virtual_memory.cu b/clang/test/dpct/virtual_memory.cu index 8ec983bb127d..5d16ce4ba901 100644 --- a/clang/test/dpct/virtual_memory.cu +++ b/clang/test/dpct/virtual_memory.cu @@ -1,6 +1,6 @@ // UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2 // UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2 -// RUN: dpct --format-range=none --use-experimental-features=virtual_memory -out-root %T/virtual_memory %s --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only +// RUN: dpct --format-range=none --use-experimental-features=virtual_mem -out-root %T/virtual_memory %s --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only // RUN: FileCheck %s --match-full-lines --input-file %T/virtual_memory/virtual_memory.dp.cpp #include #include @@ -14,35 +14,35 @@ int main() { CUcontext context; cuCtxCreate(&context, 0, device); -// CHECK: dpct::experimental::mem_prop prop = {}; -// CHECK: prop.type = dpct::experimental::mem_allocation_type::mem_allocation_type_default; -// CHECK: prop.location.type = dpct::experimental::mem_location_type::mem_location_type_device; -// CHECK: prop.location.id = device; +// CHECK: dpct::experimental::mem_prop prop = {}; +// CHECK: prop.type = 0; +// CHECK: prop.location.type = 1; +// CHECK: prop.location.id = device; CUmemAllocationProp prop = {}; prop.type = CU_MEM_ALLOCATION_TYPE_PINNED; prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE; prop.location.id = device; size_t granularity; -// CHECK: dpct::experimental::mem_get_allocation_granularity(&granularity, &prop, dpct::experimental::granularity_flags::granularity_flags_minimum); - cuMemGetAllocationGranularity(&granularity, &prop, CU_MEM_ALLOC_GRANULARITY_MINIMUM); +// CHECK: granularity = sycl::ext::oneapi::experimental::get_mem_granularity(dpct::get_device(prop.location.id), dpct::get_device(prop.location.id).get_context(), sycl::ext::oneapi::experimental::granularity_mode::minimum); + cuMemGetAllocationGranularity(&granularity, &prop, CU_MEM_ALLOC_GRANULARITY_MINIMUM); size_t POOL_SIZE = granularity; // CHECK: dpct::device_ptr reserved_addr; // CHECK: dpct::experimental::mem_handle allocHandle; -// CHECK: dpct::experimental::mem_address_reserve(&reserved_addr, POOL_SIZE, 0, 0, 0); -// CHECK: dpct::experimental::mem_create(&allocHandle, POOL_SIZE, &prop, 0); -// CHECK: dpct::experimental::mem_map(reserved_addr, POOL_SIZE, 0, allocHandle, 0); +// CHECK: reserved_addr = (dpct::device_ptr)sycl::ext::oneapi::experimental::reserve_virtual_mem((uintptr_t)0, POOL_SIZE, dpct::get_current_device().get_context()); +// CHECK: allocHandle = new sycl::ext::oneapi::experimental::physical_mem(dpct::get_device(prop.location.id), dpct::get_device(prop.location.id).get_context(), POOL_SIZE); +// CHECK: allocHandle->map((uintptr_t)reserved_addr, POOL_SIZE, sycl::ext::oneapi::experimental::address_access_mode::read_write, 0); CUdeviceptr reserved_addr; CUmemGenericAllocationHandle allocHandle; cuMemAddressReserve(&reserved_addr, POOL_SIZE, 0, 0, 0); cuMemCreate(&allocHandle, POOL_SIZE, &prop, 0); cuMemMap(reserved_addr, POOL_SIZE, 0, allocHandle, 0); -// CHECK: dpct::experimental::mem_access_desc accessDesc = {}; -// CHECK: accessDesc.location.type = dpct::experimental::mem_location_type::mem_location_type_device; -// CHECK: accessDesc.location.id = device; -// CHECK: accessDesc.flags = dpct::experimental::address_access_flags::address_access_flags_read_write; -// CHECK: dpct::experimental::mem_set_access(reserved_addr, POOL_SIZE, &accessDesc, 1); +// CHECK: dpct::experimental::mem_access_desc accessDesc = {}; +// CHECK: accessDesc.location.type = 1; +// CHECK: accessDesc.location.id = device; +// CHECK: accessDesc.flags = sycl::ext::oneapi::experimental::address_access_mode::read_write; +// CHECK: sycl::ext::oneapi::experimental::set_access_mode(reserved_addr, POOL_SIZE, accessDesc.flags, dpct::get_device(accessDesc.location.id).get_context()); CUmemAccessDesc accessDesc = {}; accessDesc.location.type = CU_MEM_LOCATION_TYPE_DEVICE; accessDesc.location.id = device; @@ -66,9 +66,9 @@ int main() { } std::cout << "test passed" << std::endl; -// CHECK: dpct::experimental::mem_unmap(reserved_addr, POOL_SIZE); -// CHECK: dpct::experimental::mem_release(allocHandle); -// CHECK: dpct::experimental::mem_address_free(reserved_addr, POOL_SIZE); + // CHECK: sycl::ext::oneapi::experimental::unmap(reserved_addr, POOL_SIZE, dpct::get_current_device().get_context()); + // CHECK: delete (allocHandle); + // CHECK: sycl::ext::oneapi::experimental::free_virtual_mem((uintptr_t)reserved_addr, POOL_SIZE, dpct::get_current_device().get_context()); cuMemUnmap(reserved_addr, POOL_SIZE); cuMemRelease(allocHandle); cuMemAddressFree(reserved_addr, POOL_SIZE); From 308173c6263f216c77ff9fdac45f055f3cde2c50 Mon Sep 17 00:00:00 2001 From: intwanghao Date: Tue, 15 Oct 2024 16:17:39 +0800 Subject: [PATCH 10/10] fix Signed-off-by: intwanghao --- clang/lib/DPCT/MapNames.cpp | 2 +- clang/runtime/dpct-rt/include/dpct/memory.hpp | 8 +++++--- clang/test/dpct/virtual_memory.cu | 2 +- 3 files changed, 7 insertions(+), 5 deletions(-) diff --git a/clang/lib/DPCT/MapNames.cpp b/clang/lib/DPCT/MapNames.cpp index aab9d64b5d87..6c7d23a811ea 100644 --- a/clang/lib/DPCT/MapNames.cpp +++ b/clang/lib/DPCT/MapNames.cpp @@ -838,7 +838,7 @@ void MapNames::setExplicitNamespaceMap( "experimental::mem_prop")}, {"CUmemGenericAllocationHandle", std::make_shared(getDpctNamespace() + - "experimental::mem_handle")}, + "experimental::physical_mem_ptr")}, {"CUmemAccessDesc", std::make_shared(getDpctNamespace() + "experimental::mem_access_desc")}, diff --git a/clang/runtime/dpct-rt/include/dpct/memory.hpp b/clang/runtime/dpct-rt/include/dpct/memory.hpp index bf98ab6732eb..a7e521e38034 100644 --- a/clang/runtime/dpct-rt/include/dpct/memory.hpp +++ b/clang/runtime/dpct-rt/include/dpct/memory.hpp @@ -1429,16 +1429,18 @@ static inline void async_dpct_memset(pitched_data pitch, int val, } namespace experimental { -typedef sycl::ext::oneapi::experimental::physical_mem *mem_handle; +typedef sycl::ext::oneapi::experimental::physical_mem *physical_mem_ptr; struct mem_location { int id; - int type; + int type; // Location type. Value 1 means device location, and thus, id is a + // device id. Other values are reserved for future use. }; struct mem_prop { mem_location location; - int type; + int type; // Memory type. Value 1 means default device memory. Other values + // are reserved for future use. }; struct mem_access_desc { diff --git a/clang/test/dpct/virtual_memory.cu b/clang/test/dpct/virtual_memory.cu index 5d16ce4ba901..a2a5a1bf47b4 100644 --- a/clang/test/dpct/virtual_memory.cu +++ b/clang/test/dpct/virtual_memory.cu @@ -28,7 +28,7 @@ int main() { size_t POOL_SIZE = granularity; // CHECK: dpct::device_ptr reserved_addr; -// CHECK: dpct::experimental::mem_handle allocHandle; +// CHECK: dpct::experimental::physical_mem_ptr allocHandle; // CHECK: reserved_addr = (dpct::device_ptr)sycl::ext::oneapi::experimental::reserve_virtual_mem((uintptr_t)0, POOL_SIZE, dpct::get_current_device().get_context()); // CHECK: allocHandle = new sycl::ext::oneapi::experimental::physical_mem(dpct::get_device(prop.location.id), dpct::get_device(prop.location.id).get_context(), POOL_SIZE); // CHECK: allocHandle->map((uintptr_t)reserved_addr, POOL_SIZE, sycl::ext::oneapi::experimental::address_access_mode::read_write, 0);