Update mempool example

Author: hokkanen

docs/04-memory.md

@@ -279,11 +279,11 @@ for (int i = 0; i < 100; i++) {
   hipMalloc(&ptr, size); 
   // Run GPU kernel
   kernel<<<..., stream>>>(ptr);
-  // Synchronize stream, does not influence memory allocations
-  hipStreamSynchronize(stream); 
   // Deallocate memory here
   hipFree(ptr); 
 }
+// Synchronize stream, does not influence memory allocations
+hipStreamSynchronize(stream); 
 ```
 * Allocating and deallocating memory in a loop is slow, and can have a significant impact on the performance
 </div>

memory/02-mempools/README.md

@@ -1,12 +1,10 @@
 # The stream-ordered memory allocator and memory pools
 
-The purpose of this exercise is to compare different memory allocation
-strategies within a loop and to understand the performance impact of using or not using a memory pool. The following timed functions are called at the end of the source file by the `main()` function:
+The purpose of this exercise is to compare different memory allocation strategies within a loop and to understand the performance impact of using or not using a memory pool. The following timed functions are called at the end of the source file by the `main()` function:
 
 * The function `noRecurringAlloc()` allocates memory outside loop only once
 * The function `recurringAllocNoMemPools()` allocates memory within a loop recurringly
-* The function `recurringAllocMemPoolNoSync()` obtains memory from a pool within a loop recurringly (no synchronization within loop)
-* The function `recurringAllocMemPoolSync()` obtains memory from a pool within a loop recurringly and synchronizes (synchronization within loop)
+* The function `recurringAllocMemPool()` obtains memory from a pool within a loop recurringly
 
 The task is to fill the missing function calls in the code indicated by lines beginning with `#error`, and followed by a descriptive instruction.
 

memory/02-mempools/mempools.cpp

@@ -61,9 +61,9 @@ void noRecurringAlloc(int nSteps, int size)
   {    
     // Launch GPU kernel
     hipKernel<<<gridsize, blocksize, 0, 0>>>(d_A, size);
-    // Synchronization
-    #error synchronize the default stream here
   }
+  // Synchronization
+  #error synchronize the default stream here
   // Check results and print timings
   checkTiming("noRecurringAlloc", (double)(clock() - tStart) / CLOCKS_PER_SEC);
 
@@ -87,17 +87,17 @@ void recurringAllocNoMemPools(int nSteps, int size)
     #error allocate memory with hipMalloc for d_A of size
     // Launch GPU kernel
     hipKernel<<<gridsize, blocksize, 0, 0>>>(d_A, size);
-    // Synchronization
-    #error synchronize the default stream here
     // Free allocation
     #error free d_A allocation using hipFree
   }
+  // Synchronization
+  #error synchronize the default stream here
   // Check results and print timings
   checkTiming("recurringAllocNoMemPools", (double)(clock() - tStart) / CLOCKS_PER_SEC);
 }
 
-/* Run using memory pooling but no recurring syncs */
-void recurringAllocMemPoolNoSync(int nSteps, int size)
+/* Do recurring allocation with memory pooling */
+void recurringAllocMemPool(int nSteps, int size)
 {
   // Create HIP stream
   hipStream_t stream;
@@ -128,38 +128,6 @@ void recurringAllocMemPoolNoSync(int nSteps, int size)
   hipStreamDestroy(stream);
 }
 
-/* Run using memory pooling and recurring syncs */
-void recurringAllocMemPoolSync(int nSteps, int size)
-{
-  // Create HIP stream
-  hipStream_t stream;
-  hipStreamCreate(&stream);
-
-  // Determine grid and block size
-  const int blocksize = BLOCKSIZE;
-  const int gridsize = (size - 1 + blocksize) / blocksize;
-
-  // Start timer and begin stepping loop
-  clock_t tStart = clock();
-  for(unsigned int i = 0; i < nSteps; i++)
-  {
-    int *d_A;
-    // Allocate pinned device memory
-    #error allocate memory with cudaMallocAsync for d_A of size in stream
-    // Launch GPU kernel
-    hipKernel<<<gridsize, blocksize, 0, stream>>>(d_A, size);
-    // Free allocation
-    #error free d_A allocation using cudaFreeAsync in stream
-    // Synchronization
-    #error synchronize the stream here
-  }
-  // Check results and print timings
-  checkTiming("recurringAllocMemPoolSync", (double)(clock() - tStart) / CLOCKS_PER_SEC);
-
-  // Destroy the stream
-  hipStreamDestroy(stream);
-}
-
 /* The main function */
 int main(int argc, char* argv[])
 {
@@ -172,6 +140,5 @@ int main(int argc, char* argv[])
   // Run with different memory allocatins strategies
   noRecurringAlloc(nSteps, size);
   recurringAllocNoMemPools(nSteps, size);
-  recurringAllocMemPoolNoSync(nSteps, size);
-  recurringAllocMemPoolSync(nSteps, size);
+  recurringAllocMemPool(nSteps, size);
 }

memory/02-mempools/solution/mempools.cpp

@@ -61,9 +61,9 @@ void noRecurringAlloc(int nSteps, int size)
   {    
     // Launch GPU kernel
     hipKernel<<<gridsize, blocksize, 0, 0>>>(d_A, size);
-    // Synchronization
-    hipStreamSynchronize(0);
   }
+  // Synchronization
+  hipStreamSynchronize(0);
   // Check results and print timings
   checkTiming("noRecurringAlloc", (double)(clock() - tStart) / CLOCKS_PER_SEC);
 
@@ -87,17 +87,17 @@ void recurringAllocNoMemPools(int nSteps, int size)
     hipMalloc((void**)&d_A, size);
     // Launch GPU kernel
     hipKernel<<<gridsize, blocksize, 0, 0>>>(d_A, size);
-    // Synchronization
-    hipStreamSynchronize(0);
     // Free allocation
     hipFree(d_A);
   }
+  // Synchronization
+  hipStreamSynchronize(0);
   // Check results and print timings
   checkTiming("recurringAllocNoMemPools", (double)(clock() - tStart) / CLOCKS_PER_SEC);
 }
 
-/* Run using memory pooling but no recurring syncs */
-void recurringAllocMemPoolNoSync(int nSteps, int size)
+/* Do recurring allocation with memory pooling */
+void recurringAllocMemPool(int nSteps, int size)
 {
   // Create HIP stream
   hipStream_t stream;
@@ -128,38 +128,6 @@ void recurringAllocMemPoolNoSync(int nSteps, int size)
   hipStreamDestroy(stream);
 }
 
-/* Run using memory pooling and recurring syncs */
-void recurringAllocMemPoolSync(int nSteps, int size)
-{
-  // Create HIP stream
-  hipStream_t stream;
-  hipStreamCreate(&stream);
-
-  // Determine grid and block size
-  const int blocksize = BLOCKSIZE;
-  const int gridsize = (size - 1 + blocksize) / blocksize;
-
-  // Start timer and begin stepping loop
-  clock_t tStart = clock();
-  for(unsigned int i = 0; i < nSteps; i++)
-  {
-    int *d_A;
-    // Allocate pinned device memory
-    cudaMallocAsync((void**)&d_A, size, stream);
-    // Launch GPU kernel
-    hipKernel<<<gridsize, blocksize, 0, stream>>>(d_A, size);
-    // Free allocation
-    cudaFreeAsync(d_A, stream);
-    // Synchronization
-    hipStreamSynchronize(stream);
-  }
-  // Check results and print timings
-  checkTiming("recurringAllocMemPoolSync", (double)(clock() - tStart) / CLOCKS_PER_SEC);
-
-  // Destroy the stream
-  hipStreamDestroy(stream);
-}
-
 /* The main function */
 int main(int argc, char* argv[])
 {
@@ -172,6 +140,5 @@ int main(int argc, char* argv[])
   // Run with different memory allocatins strategies
   noRecurringAlloc(nSteps, size);
   recurringAllocNoMemPools(nSteps, size);
-  recurringAllocMemPoolNoSync(nSteps, size);
-  recurringAllocMemPoolSync(nSteps, size);
+  recurringAllocMemPool(nSteps, size);
 }