update outdated cudaMemAdvise/cudaMemPrefetchAsync calls to the new cudaMemLocation API

Author: coketaste

modules/module2/content.md

@@ -470,19 +470,25 @@ __global__ void processData(float *data, size_t n) {
 ```cuda
 void optimizedUnifiedMemory(float *data, size_t n, int device) {
     // Prefetch data to GPU before kernel launch
-    cudaMemPrefetchAsync(data, n * sizeof(float), device);
+    cudaMemLocation loc{};
+    loc.type = cudaMemLocationTypeDevice;
+    loc.id = device;
+    cudaMemPrefetchAsync(data, n * sizeof(float), loc, /*stream=*/0);
     
     // Set memory usage hints
-    cudaMemAdvise(data, n * sizeof(float), cudaMemAdviseSetReadMostly, device);
-    cudaMemAdvise(data, n * sizeof(float), cudaMemAdviseSetPreferredLocation, device);
+    cudaMemAdvise(data, n * sizeof(float), cudaMemAdviseSetReadMostly, loc);
+    cudaMemAdvise(data, n * sizeof(float), cudaMemAdviseSetPreferredLocation, loc);
     
     // Launch kernel
     int blockSize = 256;
     int gridSize = (n + blockSize - 1) / blockSize;
     processData<<<gridSize, blockSize>>>(data, n);
     
     // Prefetch back to CPU if needed
-    cudaMemPrefetchAsync(data, n * sizeof(float), cudaCpuDeviceId);
+    cudaMemLocation hostLoc{};
+    hostLoc.type = cudaMemLocationTypeHost;
+    hostLoc.id = 0;
+    cudaMemPrefetchAsync(data, n * sizeof(float), hostLoc, /*stream=*/0);
 }
 ```
 

modules/module2/examples/04_unified_memory_cuda.cu

@@ -222,16 +222,23 @@ void demonstrateMemoryMigration() {
     }
 
     int device = 0;
+    // CUDA 13 updated UM APIs use cudaMemLocation instead of raw int device IDs
+    cudaMemLocation locDevice{};
+    locDevice.type = cudaMemLocationTypeDevice;
+    locDevice.id = device;
+    cudaMemLocation locHost{};
+    locHost.type = cudaMemLocationTypeHost;
+    locHost.id = 0; // host id is unused
 
     printf("Testing memory migration with prefetching and hints...\n");
 
-    // Set memory advice
-    CUDA_CHECK(cudaMemAdvise(data, bytes, cudaMemAdviseSetReadMostly, device));
-    CUDA_CHECK(cudaMemAdvise(data, bytes, cudaMemAdviseSetPreferredLocation, device));
+    // Set memory advice (location-aware in CUDA 13)
+    CUDA_CHECK(cudaMemAdvise(data, bytes, cudaMemAdviseSetReadMostly, locDevice));
+    CUDA_CHECK(cudaMemAdvise(data, bytes, cudaMemAdviseSetPreferredLocation, locDevice));
 
-    // Prefetch to GPU
+    // Prefetch to GPU (location-aware + explicit stream)
     printf("Prefetching to GPU...\n");
-    CUDA_CHECK(cudaMemPrefetchAsync(data, bytes, device));
+    CUDA_CHECK(cudaMemPrefetchAsync(data, bytes, locDevice, 0));
     CUDA_CHECK(cudaDeviceSynchronize());
 
     int blockSize = 256;
@@ -250,9 +257,9 @@ void demonstrateMemoryMigration() {
     float gpu_time;
     CUDA_CHECK(cudaEventElapsedTime(&gpu_time, start, stop));
 
-    // Prefetch to CPU
+    // Prefetch to CPU (location-aware + explicit stream)
     printf("Prefetching to CPU...\n");
-    CUDA_CHECK(cudaMemPrefetchAsync(data, bytes, cudaCpuDeviceId));
+    CUDA_CHECK(cudaMemPrefetchAsync(data, bytes, locHost, 0));
     CUDA_CHECK(cudaDeviceSynchronize());
 
     // CPU computation (data already on CPU)
@@ -274,9 +281,9 @@ void demonstrateMemoryMigration() {
     // Test without prefetching for comparison
     printf("\nTesting without prefetching...\n");
 
-    // Reset memory advice
-    CUDA_CHECK(cudaMemAdvise(data, bytes, cudaMemAdviseUnsetReadMostly, device));
-    CUDA_CHECK(cudaMemAdvise(data, bytes, cudaMemAdviseUnsetPreferredLocation, device));
+    // Reset memory advice (location-aware in CUDA 13)
+    CUDA_CHECK(cudaMemAdvise(data, bytes, cudaMemAdviseUnsetReadMostly, locDevice));
+    CUDA_CHECK(cudaMemAdvise(data, bytes, cudaMemAdviseUnsetPreferredLocation, locDevice));
 
     CUDA_CHECK(cudaEventRecord(start));
     computeIntensive<<<gridSize, blockSize>>>(data, n);

modules/module4/README.md

@@ -262,12 +262,15 @@ for (int chunk = 0; chunk < numChunks; chunk++) {
 
 **Memory Hints:**
 ```cuda
-// Guide data placement
-cudaMemAdvise(data, size, cudaMemAdviseSetReadMostly, deviceId);
-cudaMemAdvise(data, size, cudaMemAdviseSetPreferredLocation, deviceId);
-
-// Prefetch data proactively
-cudaMemPrefetchAsync(data, size, deviceId);
+// Guide data placement (CUDA 13+)
+cudaMemLocation loc{};
+loc.type = cudaMemLocationTypeDevice;
+loc.id = deviceId;
+cudaMemAdvise(data, size, cudaMemAdviseSetReadMostly, loc);
+cudaMemAdvise(data, size, cudaMemAdviseSetPreferredLocation, loc);
+
+// Prefetch data proactively (CUDA 13+)
+cudaMemPrefetchAsync(data, size, loc, /*stream=*/0);
 ```
 
 ### 4. P2P Communication Patterns

modules/module4/content.md

@@ -135,11 +135,14 @@ kernel<<<grid, block>>>(data, n);
 #### Memory Access Patterns
 
 ```cuda
-// Prefetch data to GPU
-cudaMemPrefetchAsync(data, size, deviceId);
-
-// Provide memory access hints
-cudaMemAdvise(data, size, cudaMemAdviseSetReadMostly, deviceId);
+// Prefetch data to GPU (CUDA 13+)
+cudaMemLocation loc{};
+loc.type = cudaMemLocationTypeDevice;
+loc.id = deviceId;
+cudaMemPrefetchAsync(data, size, loc, /*stream=*/0);
+
+// Provide memory access hints (CUDA 13+)
+cudaMemAdvise(data, size, cudaMemAdviseSetReadMostly, loc);
 ```
 
 #### Unified Memory Best Practices

modules/module4/examples/03_unified_memory.cu

@@ -204,23 +204,30 @@ double optimizedUnifiedMemory(int n) {
 
     auto start = std::chrono::high_resolution_clock::now();
 
-    // Provide memory hints
+    // Provide memory hints (CUDA 13: use cudaMemLocation)
     int deviceId = 0;
-    CUDA_CHECK(cudaMemAdvise(a, bytes, cudaMemAdviseSetReadMostly, deviceId));
-    CUDA_CHECK(cudaMemAdvise(b, bytes, cudaMemAdviseSetReadMostly, deviceId));
+    cudaMemLocation locDevice{};
+    locDevice.type = cudaMemLocationTypeDevice;
+    locDevice.id = deviceId;
+    cudaMemLocation locHost{};
+    locHost.type = cudaMemLocationTypeHost;
+    locHost.id = 0;
+
+    CUDA_CHECK(cudaMemAdvise(a, bytes, cudaMemAdviseSetReadMostly, locDevice));
+    CUDA_CHECK(cudaMemAdvise(b, bytes, cudaMemAdviseSetReadMostly, locDevice));
 
-    // Prefetch data to GPU
-    CUDA_CHECK(cudaMemPrefetchAsync(a, bytes, deviceId));
-    CUDA_CHECK(cudaMemPrefetchAsync(b, bytes, deviceId));
+    // Prefetch data to GPU (location-aware + explicit stream)
+    CUDA_CHECK(cudaMemPrefetchAsync(a, bytes, locDevice, 0));
+    CUDA_CHECK(cudaMemPrefetchAsync(b, bytes, locDevice, 0));
 
     // Launch kernel
     dim3 block(BLOCK_SIZE);
     dim3 grid((n + block.x - 1) / block.x);
     vectorAdd<<<grid, block>>>(a, b, c, n);
     CUDA_CHECK(cudaGetLastError());
 
-    // Prefetch result back to CPU
-    CUDA_CHECK(cudaMemPrefetchAsync(c, bytes, cudaCpuDeviceId));
+    // Prefetch result back to CPU (location-aware + explicit stream)
+    CUDA_CHECK(cudaMemPrefetchAsync(c, bytes, locHost, 0));
     CUDA_CHECK(cudaDeviceSynchronize());
 
     // Access result on CPU
@@ -376,9 +383,13 @@ void multiGPUUnifiedMemory(int n) {
         int offset = gpu * chunkSize;
         int currentChunkSize = (gpu == deviceCount - 1) ? n - offset : chunkSize;
 
-        // Prefetch chunk to current GPU
-        CUDA_CHECK(cudaMemPrefetchAsync(data + offset, 
-                                       currentChunkSize * sizeof(float), gpu));
+    // Prefetch chunk to current GPU (location-aware + explicit stream)
+    cudaMemLocation locGpu{};
+    locGpu.type = cudaMemLocationTypeDevice;
+    locGpu.id = gpu;
+    CUDA_CHECK(cudaMemPrefetchAsync(data + offset,
+                       currentChunkSize * sizeof(float),
+                       locGpu, 0));
 
         // Process on this GPU
         dim3 block(BLOCK_SIZE);