zeVirtualMemQueryPageSize probe

README.md

Probes zeVirtualMemQueryPageSize across a range of allocation sizes to see how Battlemage L0 actually scales the page size and compares that with what public API returns.

Device: Intel(R) Arc(TM) Pro B60 Graphics
request      | page size   
-------------+-------------
1 B          | 64.00 KiB   
4.00 KiB     | 64.00 KiB   
64.00 KiB    | 64.00 KiB   
1.00 MiB     | 64.00 KiB   
2.00 MiB     | 2.00 MiB    
4.00 MiB     | 2.00 MiB    
16.00 MiB    | 2.00 MiB    
64.00 MiB    | 2.00 MiB    
256.00 MiB   | 2.00 MiB    
1.00 GiB     | 2.00 MiB    
4.00 GiB     | 2.00 MiB    
16.00 GiB    | 2.00 MiB    
32.00 GiB    | 2.00 MiB    

SYCL get_mem_granularity (size-agnostic, AllocationSize=1):
  minimum     = 64.00 KiB
  recommended = 64.00 KiB

Building

source /opt/intel/oneapi/setvars.sh --force
icpx -fsycl probe_page_size.cpp -lze_loader -o probe_page_size

probe_page_size.cpp

// Probe zeVirtualMemQueryPageSize across a range of allocation sizes
// to see how Battlemage L0 actually scales the page size.
//
// Build: icpx -fsycl probe_page_size.cpp -lze_loader -o probe_page_size
// Run:   ./probe_page_size

#include <level_zero/ze_api.h>
#include <sycl/sycl.hpp>
#include <sycl/ext/oneapi/backend/level_zero.hpp>
#include <cstdio>
#include <cstdint>

static const char * fmt_size(size_t b, char * buf, size_t buflen) {
    if (b >= (1ull << 30)) snprintf(buf, buflen, "%.2f GiB", b / double(1ull << 30));
    else if (b >= (1ull << 20)) snprintf(buf, buflen, "%.2f MiB", b / double(1ull << 20));
    else if (b >= (1ull << 10)) snprintf(buf, buflen, "%.2f KiB", b / double(1ull << 10));
    else snprintf(buf, buflen, "%zu B", b);
    return buf;
}

int main() {
    sycl::queue q;
    auto dev = q.get_device();
    auto ctx = q.get_context();

    printf("Device: %s\n", dev.get_info<sycl::info::device::name>().c_str());

    auto ze_ctx = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(ctx);
    auto ze_dev = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(dev);

    // Sweep across small / medium / large / huge.
    const size_t sizes[] = {
        1ull,
        4ull << 10,         // 4 KiB
        64ull << 10,        // 64 KiB
        1ull << 20,         // 1 MiB
        2ull << 20,         // 2 MiB
        4ull << 20,         // 4 MiB
        16ull << 20,        // 16 MiB
        64ull << 20,        // 64 MiB
        256ull << 20,       // 256 MiB
        1ull << 30,         // 1 GiB
        4ull << 30,         // 4 GiB
        16ull << 30,        // 16 GiB
        32ull << 30,        // 32 GiB
    };

    char rbuf[32], pbuf[32];
    printf("%-12s | %-12s\n", "request", "page size");
    printf("%-12s-+-%-12s\n", "------------", "------------");
    for (size_t s : sizes) {
        size_t page = 0;
        auto r = zeVirtualMemQueryPageSize(ze_ctx, ze_dev, s, &page);
        if (r != ZE_RESULT_SUCCESS) {
            printf("%-12s | error 0x%x\n", fmt_size(s, rbuf, sizeof(rbuf)), r);
            continue;
        }
        printf("%-12s | %-12s\n", fmt_size(s, rbuf, sizeof(rbuf)),
               fmt_size(page, pbuf, sizeof(pbuf)));
    }

    // Also show what SYCL's public API returns, for comparison.
    using namespace sycl::ext::oneapi::experimental;
    size_t sycl_min  = get_mem_granularity(dev, ctx, granularity_mode::minimum);
    size_t sycl_rec  = get_mem_granularity(dev, ctx, granularity_mode::recommended);
    printf("\nSYCL get_mem_granularity (size-agnostic, AllocationSize=1):\n");
    printf("  minimum     = %s\n", fmt_size(sycl_min, pbuf, sizeof(pbuf)));
    printf("  recommended = %s\n", fmt_size(sycl_rec, pbuf, sizeof(pbuf)));
    return 0;
}