diff --git a/.github/workflows/sycl-linux-run-tests.yml b/.github/workflows/sycl-linux-run-tests.yml index 528772589c01d..5dad941e5a64e 100644 --- a/.github/workflows/sycl-linux-run-tests.yml +++ b/.github/workflows/sycl-linux-run-tests.yml @@ -330,6 +330,18 @@ jobs: cat /usr/local/lib/igc/IGCTAG.txt fi + # Best-effort reservation of a few HugeTLB pages so the explicit huge-page + # path of the register_host_memory E2E test is exercised rather than being + # UNSUPPORTED. Failing to set it just means the huge-page test is UNSUPPORTED + # there, so never fail the job. + - name: Reserve HugeTLB pages + if: inputs.tests_selector == 'e2e' + continue-on-error: true + run: | + grep -i Huge /proc/meminfo + echo 2 | sudo tee /proc/sys/vm/nr_hugepages || true + grep -i Huge /proc/meminfo + - name: Run E2E Tests if: inputs.tests_selector == 'e2e' uses: ./devops/actions/run-tests/linux/e2e diff --git a/sycl/test-e2e/USM/Inputs/register_host_memory_helpers.hpp b/sycl/test-e2e/USM/Inputs/register_host_memory_helpers.hpp new file mode 100644 index 0000000000000..59903423dcfe6 --- /dev/null +++ b/sycl/test-e2e/USM/Inputs/register_host_memory_helpers.hpp @@ -0,0 +1,32 @@ +// Shared helpers for the sycl_ext_oneapi_register_host_memory end-to-end tests. + +#pragma once + +#include +#include + +#if defined(_WIN32) +#include +#else +#include +#endif + +// Returns the host page size. The extension requires registered ranges to be +// aligned to (and a multiple of) the host page size. +inline size_t getHostPageSize() { +#if defined(_WIN32) + SYSTEM_INFO Info; + GetSystemInfo(&Info); + return static_cast(Info.dwPageSize); +#else + return static_cast(sysconf(_SC_PAGESIZE)); +#endif +} + +// Rounds NumBytes up to a whole number of host pages, as the extension requires +// the registered size to be a multiple of the host page size. +inline size_t roundUpToPage(size_t NumBytes, size_t PageSize) { + assert(PageSize != 0 && (PageSize & (PageSize - 1)) == 0 && + "PageSize must be a power of two"); + return (NumBytes + PageSize - 1) & ~(PageSize - 1); +} diff --git a/sycl/test-e2e/USM/register_host_memory.cpp b/sycl/test-e2e/USM/register_host_memory.cpp index 4f0ca89a16697..b411a61852464 100644 --- a/sycl/test-e2e/USM/register_host_memory.cpp +++ b/sycl/test-e2e/USM/register_host_memory.cpp @@ -14,6 +14,8 @@ // code (device writes to a read_only range are undefined behavior and are // therefore not exercised). +#include "Inputs/register_host_memory_helpers.hpp" + #include #include #include @@ -24,23 +26,10 @@ #if defined(_WIN32) #include -#include -#else -#include #endif namespace syclexp = sycl::ext::oneapi::experimental; -static size_t getHostPageSize() { -#if defined(_WIN32) - SYSTEM_INFO Info; - GetSystemInfo(&Info); - return static_cast(Info.dwPageSize); -#else - return static_cast(sysconf(_SC_PAGESIZE)); -#endif -} - static void *allocatePageAligned(size_t Alignment, size_t Size) { #if defined(_WIN32) return _aligned_malloc(Size, Alignment); @@ -64,8 +53,7 @@ int main() { const size_t PageSize = getHostPageSize(); const size_t NumElems = 1024; // Round the byte size up to a multiple of the page size as required. - size_t NumBytes = NumElems * sizeof(int); - NumBytes = (NumBytes + PageSize - 1) & ~(PageSize - 1); + const size_t NumBytes = roundUpToPage(NumElems * sizeof(int), PageSize); int *Data = static_cast(allocatePageAligned(PageSize, NumBytes)); assert(Data != nullptr && "host allocation failed"); diff --git a/sycl/test-e2e/USM/register_host_memory_huge_pages.cpp b/sycl/test-e2e/USM/register_host_memory_huge_pages.cpp new file mode 100644 index 0000000000000..3960bd1f16416 --- /dev/null +++ b/sycl/test-e2e/USM/register_host_memory_huge_pages.cpp @@ -0,0 +1,152 @@ +// REQUIRES: aspect-ext_oneapi_register_host_memory +// REQUIRES: level_zero_v2_adapter +// REQUIRES: linux +// REQUIRES: hugepages + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out | FileCheck %s + +// End-to-end test for sycl_ext_oneapi_register_host_memory with huge-page +// backed host memory. Two huge-page acquisition paths are exercised: +// - explicit huge pages via mmap(MAP_HUGETLB), and +// - transparent huge pages via mmap + madvise(MADV_HUGEPAGE). +// +// Explicit huge pages require the OS to have huge pages reserved (e.g. via +// /proc/sys/vm/nr_hugepages). The test therefore REQUIRES the "hugepages" +// feature (free HugeTLB pages present) so it is UNSUPPORTED rather than +// silently skipped where none are available. A huge page is itself a multiple +// of the host base page size, so a huge-page-aligned range trivially satisfies +// the extension's page-alignment and size requirements. + +#include "Inputs/register_host_memory_helpers.hpp" + +#include +#include +#include + +#include +#include +#include +#include +#include +#include + +#include +#include + +// Older glibc headers may not define MAP_HUGETLB; define it to its well-known +// value so the test still builds. The mmap call will simply fail at runtime if +// the running kernel does not support the flag, and that path is skipped. +#ifndef MAP_HUGETLB +#define MAP_HUGETLB 0x40000 +#endif + +namespace syclexp = sycl::ext::oneapi::experimental; + +// Returns the kernel's default HugeTLB page size in bytes, read from +// /proc/meminfo. Falls back to 2 MiB (the common x86-64 default) if the entry +// is missing. +static size_t getHugePageSize() { + std::ifstream Meminfo("/proc/meminfo"); + std::string Key; + while (Meminfo >> Key) { + if (Key == "Hugepagesize:") { + size_t KiB = 0; + Meminfo >> KiB; + if (KiB != 0) + return KiB * 1024; + break; + } + Meminfo.ignore(std::numeric_limits::max(), '\n'); + } + return 2 * 1024 * 1024; +} + +// Runs device + copy exercises over a registered range and verifies results. +// Reused by both huge-page paths. +static void exerciseRegisteredRange(sycl::queue &Q, sycl::context &Ctxt, + int *Data, size_t NumElems) { + syclexp::register_host_memory(Data, NumElems * sizeof(int), Ctxt); + + // The pointer behaves like a USM host allocation while registered. + assert(sycl::get_pointer_type(Data, Ctxt) == sycl::usm::alloc::host); + // Interior pointers are reported as host allocations too. + assert(sycl::get_pointer_type(Data + NumElems / 2, Ctxt) == + sycl::usm::alloc::host); + + // Use the registered pointer directly from device code. + Q.parallel_for(NumElems, [=](sycl::id<1> I) { + Data[I] = static_cast(I.get(0)) + 11; + }).wait(); + for (size_t I = 0; I < NumElems; ++I) + assert(Data[I] == static_cast(I) + 11); + + // Explicit copy out of the registered range. + std::vector HostDst(NumElems, 0); + Q.memcpy(HostDst.data(), Data, NumElems * sizeof(int)).wait(); + for (size_t I = 0; I < NumElems; ++I) + assert(HostDst[I] == static_cast(I) + 11); + + syclexp::unregister_host_memory(Data, Ctxt); +} + +// Path 1: explicit huge pages via MAP_HUGETLB. The "hugepages" REQUIRES feature +// guarantees free HugeTLB pages exist, so the mapping is expected to succeed +// when the mapping length matches the kernel's default HugeTLB page size. +static void testExplicitHugePages(sycl::queue &Q, sycl::context &Ctxt) { + // One whole huge page worth of memory, sized to the kernel's default HugeTLB + // page size so MAP_HUGETLB accepts the length regardless of platform. + const size_t NumBytes = getHugePageSize(); + const size_t NumElems = NumBytes / sizeof(int); + + void *Map = mmap(nullptr, NumBytes, PROT_READ | PROT_WRITE, + MAP_PRIVATE | MAP_ANONYMOUS | MAP_HUGETLB, -1, 0); + assert(Map != MAP_FAILED && "explicit MAP_HUGETLB mmap failed"); + int *Data = static_cast(Map); + + exerciseRegisteredRange(Q, Ctxt, Data, NumElems); + + assert(munmap(Map, NumBytes) == 0 && "munmap failed"); +} + +// Path 2: transparent huge pages. madvise(MADV_HUGEPAGE) is a best-effort hint +// and does not change the mapping's address or size, so registration of the +// (base-page-aligned) range is valid regardless of whether the kernel actually +// backs it with a transparent huge page. +static void testTransparentHugePages(sycl::queue &Q, sycl::context &Ctxt) { + const size_t PageSize = getHostPageSize(); + // Request a region whose length is a multiple of the kernel's HugeTLB page + // size, giving the kernel a chance to promote it to a transparent huge page. + // Note: mmap(nullptr, ...) only guarantees base-page alignment for the + // returned address, not huge-page alignment — this is a best-effort THP hint + // (size-only), so the promotion is not guaranteed. + const size_t NumBytes = getHugePageSize(); + const size_t NumElems = NumBytes / sizeof(int); + + void *Map = mmap(nullptr, NumBytes, PROT_READ | PROT_WRITE, + MAP_PRIVATE | MAP_ANONYMOUS, -1, 0); + assert(Map != MAP_FAILED && "anonymous mmap failed"); + assert((reinterpret_cast(Map) & (PageSize - 1)) == 0 && + "mmap result is not page aligned"); + + // Best-effort hint; ignore the result. Even if unsupported, the plain + // anonymous mapping below is still valid registrable host memory. + (void)madvise(Map, NumBytes, MADV_HUGEPAGE); + + int *Data = static_cast(Map); + exerciseRegisteredRange(Q, Ctxt, Data, NumElems); + + assert(munmap(Map, NumBytes) == 0 && "munmap failed"); +} + +int main() { + sycl::queue Q; + sycl::context Ctxt = Q.get_context(); + + testExplicitHugePages(Q, Ctxt); + testTransparentHugePages(Q, Ctxt); + + // CHECK: Done (explicit and transparent huge pages tested). + printf("Done (explicit and transparent huge pages tested).\n"); + return 0; +} diff --git a/sycl/test-e2e/USM/register_host_memory_mmap.cpp b/sycl/test-e2e/USM/register_host_memory_mmap.cpp new file mode 100644 index 0000000000000..b41b154c57993 --- /dev/null +++ b/sycl/test-e2e/USM/register_host_memory_mmap.cpp @@ -0,0 +1,60 @@ +// REQUIRES: aspect-ext_oneapi_register_host_memory +// REQUIRES: level_zero_v2_adapter +// REQUIRES: linux + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +// End-to-end test for sycl_ext_oneapi_register_host_memory with host memory +// obtained from an anonymous private read/write mmap mapping +// (MAP_PRIVATE | MAP_ANONYMOUS). mmap returns page-aligned memory, so the +// registration's page-alignment requirement is satisfied naturally. The +// registered pointer is used directly in device code and read back through an +// explicit copy. + +#include "Inputs/register_host_memory_helpers.hpp" + +#include +#include +#include + +#include +#include + +#include +#include + +namespace syclexp = sycl::ext::oneapi::experimental; + +int main() { + sycl::queue Q; + sycl::context Ctxt = Q.get_context(); + const size_t PageSize = getHostPageSize(); + + const size_t NumElems = 1024; + const size_t NumBytes = roundUpToPage(NumElems * sizeof(int), PageSize); + + void *Map = mmap(nullptr, NumBytes, PROT_READ | PROT_WRITE, + MAP_PRIVATE | MAP_ANONYMOUS, -1, 0); + assert(Map != MAP_FAILED && "anonymous mmap failed"); + int *Data = static_cast(Map); + + syclexp::register_host_memory(Data, NumBytes, Ctxt); + + // The pointer behaves like a USM host allocation while registered. + assert(sycl::get_pointer_type(Data, Ctxt) == sycl::usm::alloc::host); + + Q.parallel_for(NumElems, [=](sycl::id<1> I) { + Data[I] = static_cast(I.get(0)) + 1; + }).wait(); + + std::vector HostDst(NumElems, 0); + Q.memcpy(HostDst.data(), Data, NumElems * sizeof(int)).wait(); + for (size_t I = 0; I < NumElems; ++I) + assert(HostDst[I] == static_cast(I) + 1); + + syclexp::unregister_host_memory(Data, Ctxt); + assert(munmap(Map, NumBytes) == 0 && "munmap failed"); + + return 0; +} diff --git a/sycl/test-e2e/USM/register_host_memory_mmap_file_backed.cpp b/sycl/test-e2e/USM/register_host_memory_mmap_file_backed.cpp new file mode 100644 index 0000000000000..bc7f3f64fd3cf --- /dev/null +++ b/sycl/test-e2e/USM/register_host_memory_mmap_file_backed.cpp @@ -0,0 +1,73 @@ +// REQUIRES: aspect-ext_oneapi_register_host_memory +// REQUIRES: level_zero_v2_adapter +// REQUIRES: linux + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out %t + +// End-to-end test for sycl_ext_oneapi_register_host_memory with host memory +// obtained from a file-backed shared mmap mapping (MAP_SHARED) over a temporary +// file. mmap returns page-aligned memory, so the registration's page-alignment +// requirement is satisfied naturally. The temporary file is created next to +// TmpPrefix (a path within the test's output directory) rather than a global +// location such as /tmp. + +#include "Inputs/register_host_memory_helpers.hpp" + +#include +#include +#include + +#include +#include +#include + +#include +#include +#include + +namespace syclexp = sycl::ext::oneapi::experimental; + +int main(int argc, char **argv) { + // A path prefix inside the test's output directory is passed as the first + // argument and is used for the temporary file. + std::string TmpPrefix = argc > 1 ? argv[1] : "reghostmem"; + + sycl::queue Q; + sycl::context Ctxt = Q.get_context(); + const size_t PageSize = getHostPageSize(); + + const size_t NumElems = 1024; + const size_t NumBytes = roundUpToPage(NumElems * sizeof(int), PageSize); + + std::string Tmpl = TmpPrefix + "_reghostmem_XXXXXX"; + int Fd = mkstemp(Tmpl.data()); + assert(Fd >= 0 && "mkstemp failed"); + // Unlink immediately; the open fd keeps the file alive until close. + unlink(Tmpl.c_str()); + assert(ftruncate(Fd, static_cast(NumBytes)) == 0 && + "ftruncate failed"); + + void *Map = + mmap(nullptr, NumBytes, PROT_READ | PROT_WRITE, MAP_SHARED, Fd, 0); + assert(Map != MAP_FAILED && "file-backed mmap failed"); + int *Data = static_cast(Map); + + syclexp::register_host_memory(Data, NumBytes, Ctxt); + assert(sycl::get_pointer_type(Data, Ctxt) == sycl::usm::alloc::host); + + std::vector HostSrc(NumElems); + for (size_t I = 0; I < NumElems; ++I) + HostSrc[I] = static_cast(I) - 5; + Q.memcpy(Data, HostSrc.data(), NumElems * sizeof(int)).wait(); + + Q.parallel_for(NumElems, [=](sycl::id<1> I) { Data[I] *= 2; }).wait(); + for (size_t I = 0; I < NumElems; ++I) + assert(Data[I] == (static_cast(I) - 5) * 2); + + syclexp::unregister_host_memory(Data, Ctxt); + assert(munmap(Map, NumBytes) == 0 && "munmap failed"); + close(Fd); + + return 0; +} diff --git a/sycl/test-e2e/USM/register_host_memory_mmap_read_only.cpp b/sycl/test-e2e/USM/register_host_memory_mmap_read_only.cpp new file mode 100644 index 0000000000000..c4624cebcdb33 --- /dev/null +++ b/sycl/test-e2e/USM/register_host_memory_mmap_read_only.cpp @@ -0,0 +1,66 @@ +// REQUIRES: aspect-ext_oneapi_register_host_memory +// REQUIRES: level_zero_v2_adapter +// REQUIRES: linux + +// UNSUPPORTED: aspect-ext_oneapi_is_integrated_gpu +// UNSUPPORTED-TRACKER: GSD-13000 + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +// End-to-end test for sycl_ext_oneapi_register_host_memory with a read-only +// mmap'd host mapping. The host writes the contents while the mapping is still +// writable, then drops write permission via mprotect before registering, +// mirroring how an application might register a read-only mapping. Device code +// only reads it. + +#include "Inputs/register_host_memory_helpers.hpp" + +#include +#include +#include + +#include + +#include +#include + +namespace syclexp = sycl::ext::oneapi::experimental; + +int main() { + sycl::queue Q; + sycl::context Ctxt = Q.get_context(); + const size_t PageSize = getHostPageSize(); + + const size_t NumElems = 1024; + const size_t NumBytes = roundUpToPage(NumElems * sizeof(int), PageSize); + + void *Map = mmap(nullptr, NumBytes, PROT_READ | PROT_WRITE, + MAP_PRIVATE | MAP_ANONYMOUS, -1, 0); + assert(Map != MAP_FAILED && "anonymous mmap failed"); + int *Data = static_cast(Map); + + for (size_t I = 0; I < NumElems; ++I) + Data[I] = static_cast(I) * 3; + + // Make the mapping read-only on the host. The read_only property allows + // registering host memory that the application cannot write. + assert(mprotect(Map, NumBytes, PROT_READ) == 0 && "mprotect failed"); + + syclexp::register_host_memory(Data, NumBytes, Ctxt, + syclexp::properties{syclexp::read_only}); + assert(sycl::get_pointer_type(Data, Ctxt) == sycl::usm::alloc::host); + + // Device code reads the read_only range, writing results elsewhere. + int *Out = sycl::malloc_shared(NumElems, Q); + assert(Out != nullptr && "shared allocation failed"); + Q.parallel_for(NumElems, [=](sycl::id<1> I) { Out[I] = Data[I] + 1; }).wait(); + for (size_t I = 0; I < NumElems; ++I) + assert(Out[I] == static_cast(I) * 3 + 1); + + sycl::free(Out, Q); + syclexp::unregister_host_memory(Data, Ctxt); + assert(munmap(Map, NumBytes) == 0 && "munmap failed"); + + return 0; +} diff --git a/sycl/test-e2e/USM/register_host_memory_reregister.cpp b/sycl/test-e2e/USM/register_host_memory_reregister.cpp new file mode 100644 index 0000000000000..a1b465962789d --- /dev/null +++ b/sycl/test-e2e/USM/register_host_memory_reregister.cpp @@ -0,0 +1,78 @@ +// REQUIRES: aspect-ext_oneapi_register_host_memory +// REQUIRES: level_zero_v2_adapter +// REQUIRES: linux + +// UNSUPPORTED: true +// UNSUPPORTED-TRACKER: GSD-12994 + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +// End-to-end test for sycl_ext_oneapi_register_host_memory covering repeated +// registration of host memory in a single process. A page-aligned anonymous +// read/write mmap mapping is registered, used from device code, and +// unregistered; the same host range is then re-registered (writable) and +// written to again by the device. This exercises that unregistering a range +// fully releases it so a later registration of the same virtual address starts +// from clean device state. + +#include "Inputs/register_host_memory_helpers.hpp" + +#include +#include +#include + +#include +#include + +#include +#include + +namespace syclexp = sycl::ext::oneapi::experimental; + +// Registers NumBytes of Data as host memory (must be a multiple of the host +// page size), has the device write Base + I into element I over the first +// NumElems ints, verifies the result, then unregisters. All registrations are +// writable. +static void registerWriteVerifyUnregister(sycl::queue &Q, sycl::context &Ctxt, + int *Data, size_t NumBytes, + size_t NumElems, int Base) { + syclexp::register_host_memory(Data, NumBytes, Ctxt); + + // While registered, the pointer behaves like a USM host allocation. + assert(sycl::get_pointer_type(Data, Ctxt) == sycl::usm::alloc::host); + + Q.parallel_for(NumElems, [=](sycl::id<1> I) { + Data[I] = static_cast(I.get(0)) + Base; + }).wait(); + for (size_t I = 0; I < NumElems; ++I) + assert(Data[I] == static_cast(I) + Base); + + syclexp::unregister_host_memory(Data, Ctxt); +} + +int main() { + sycl::queue Q; + sycl::context Ctxt = Q.get_context(); + const size_t PageSize = getHostPageSize(); + + const size_t NumElems = 1024; + const size_t NumBytes = roundUpToPage(NumElems * sizeof(int), PageSize); + + void *Map = mmap(nullptr, NumBytes, PROT_READ | PROT_WRITE, + MAP_PRIVATE | MAP_ANONYMOUS, -1, 0); + assert(Map != MAP_FAILED && "anonymous mmap failed"); + int *Data = static_cast(Map); + + // First registration cycle over the range. + registerWriteVerifyUnregister(Q, Ctxt, Data, NumBytes, NumElems, /*Base=*/1); + + // Second registration cycle reusing the same virtual address. The device + // write here must land in the range after the first cycle was unregistered. + registerWriteVerifyUnregister(Q, Ctxt, Data, NumBytes, NumElems, + /*Base=*/100); + + assert(munmap(Map, NumBytes) == 0 && "munmap failed"); + + return 0; +} diff --git a/sycl/test-e2e/lit.cfg.py b/sycl/test-e2e/lit.cfg.py index 6d935fec57c13..6da64016d3c63 100644 --- a/sycl/test-e2e/lit.cfg.py +++ b/sycl/test-e2e/lit.cfg.py @@ -151,6 +151,19 @@ "LD_LIBRARY_PATH", config.sycl_libs_dir, append_path=True ) + # Expose a "hugepages" feature when the system has free HugeTLB pages, so + # tests that require explicit huge-page mappings (mmap(MAP_HUGETLB)) run + # only where they can actually get such pages and are UNSUPPORTED (rather + # than silently self-skipping) elsewhere. + try: + with open("/proc/meminfo") as meminfo: + for line in meminfo: + if line.startswith("HugePages_Free:") and int(line.split()[1]) > 0: + config.available_features.add("hugepages") + break + except OSError: + pass + elif platform.system() == "Windows": config.available_features.add("windows") llvm_config.with_system_environment(