Portable UM prefetch/advise helper (CUDA-13 signature under nvcc)#

A repo portable-spelling helper (like LAUNCH_STREAM): call uvmPrefetch/uvmAdvise in a test and it works on both toolchains. On the sim these are no-ops.

Defined in code/simulator.h.

This page documents 2 symbols: 2 functions.

Functions#

Signature

Description

Availability

Location

Example

inline void uvmAdvise(const void* , size_t , cudaMemoryAdvise , int )

sim + nvcc

simulator.h:2483

View

inline void uvmPrefetch(const void* , size_t , int , cudaStream_t = nullptr)

A repo portable-spelling helper (like LAUNCH_STREAM): call uvmPrefetch/uvmAdvise in a test and it works on both toolchains.

sim + nvcc

simulator.h:2482

View

Examples#

Code from the repository’s example corpus (code\*.cu); the highlighted line is the call site. These blocks are what the View links above open.

uvmAdvise (test_web_cis_unified_prefetch.cu:64)
54    cudaMallocManaged(&a, bytes);
55    cudaMallocManaged(&b, bytes);
56    cudaMallocManaged(&c, bytes);
57
58    // Get the device ID for prefetching calls (source assigns the return value;
59    // on device 0 that resolves to 0 either way).
60    int id = cudaGetDevice(&id);
61
62    // Set some hints about the data and do some prefetching
63    uvmAdvise(a, bytes, cudaMemAdviseSetPreferredLocation, cudaCpuDeviceId);
64    uvmAdvise(b, bytes, cudaMemAdviseSetPreferredLocation, cudaCpuDeviceId);
65    uvmPrefetch(c, bytes, id);
66
67    // Initialize vectors
68    for (int i = 0; i < N; i++)
69    {
70        a[i] = rand() % 100;
71        b[i] = rand() % 100;
72    }
73
74    // Pre-fetch 'a' and 'b' arrays to the specified device (GPU)
uvmPrefetch (test_web_array_inc.cu:68)
58                  static_cast<int>(e), cudaGetErrorString(e));
59      return 0;
60  }
61  cudaDeviceProp prop{}; cudaGetDeviceProperties(&prop, 0);
62  std::printf("Device 0: %s (compute %d.%d)\n", prop.name, prop.major, prop.minor);
63
64  int *h_array;
65  alloc_bytes(h_array, ds*sizeof(h_array[0]));
66  cudaCheckErrors("cudaMallocManaged Error");
67  memset(h_array, 0, ds*sizeof(h_array[0]));
68  uvmPrefetch(h_array, ds*sizeof(h_array[0]), 0); // add in step 2c
69  LAUNCH(inc, 256, 256, h_array, ds);                       // [SIM] was inc<<<256,256>>>
70  cudaCheckErrors("kernel launch error");
71  uvmPrefetch(h_array, ds*sizeof(h_array[0]), cudaCpuDeviceId); // add in step 2c
72  cudaDeviceSynchronize();
73  cudaCheckErrors("kernel execution error");
74
75  // CPU-reference self-check (was the mismatch loop + "success!")
76  bool ok = true;
77  for (size_t i = 0; i < ds; i++)
78    if (h_array[i] != 1) { ok = false; break; }
79  std::printf("array_inc (managed, ds=%zu, grid-stride 256x256): %s\n",
On this page