The kinds of variables that can be managed by the Unified Memory subsystem

The Unified Memory technique can manage the following kinds of variables:

  1. Global variables qualified with __managed__

    Examples: 

        __managed__ int x;            // In last example

  __managed__ float A[100];     // one-dim array

  __managed__ double B[10][10]; // 2-dim array

  2. (Dynamically) allocate variables

      • The CUDA library function cudaMallocManaged( ) can (dynamically) allocate variables that are managed by the Unified Memory method

Using global variables as shared variables between CPU and GPU
 

  • Global variables are shared between the CPU and GPU when they are qualified with the CUDA keyword __managed__.

    Examples: 

        __managed__ int x;            // In last example

  __managed__ float A[100];     // one-dim array

  __managed__ double B[10][10]; // 2-dim array

Since CUDA programs use arrays (1-dim and 2-dim) frequently, I will show you an example of sharing a global array variable next

Example: sharing a global array between CPU and GPU 

__managed__ int x[10]; // Defines global shared array variable

__global__ void GPU_func( )
{
   printf("++ GPU sees x: ");
   for (int i = 0; i < 10; i++ ) 
      printf("%d ", x[i]);
   printf("\n");

   for (int i = 0; i < 10; i++ ) 
      x[i]+=i; // Update x[i]
}

int main()
{
   for (int i = 0; i < 10; i++ ) 
      x[i] = 1000+i;

   GPU_func<<< 1, 1 >>>( );  // Launch kernel on grid

   cudaDeviceSynchronize();

   printf("** CPU sees x: ");
   for (int i = 0; i < 10; i++ ) 
      printf("%d ", x[i]);
   printf("\n");
}

DEMO: /home/cs355001/demo/CUDA/2-unified-mem/shared-global-array.cu

Dynamically allocated variables
 

  • Global variables are fixed size

    I.e.: array variables must be defined with constant values: 

        __managed__ float  A[100]   
  __managed__ double B[10][10]

  • Dynamically allocated variables can have variable size

    Advantage of dynamically allocated variables

      • You can create an array of the right size for your problem

 

Before I show you how to create managed dynamically allocated variables, I will review some relevant C material first.