- Recall:
the host memory and
device memory
are separate memory systems
- Data must be
transfered between
host memory and
device memory:
- Prior to the
release of CUDA 6:
- CUDA programmers must use the CUDA library function cudaMemcpy( ) to transfer data between the host memory and device memory !!!
- Unified memory introduced
in CUDA 6:
- Unified memory =
a memory model presented to
the CUDA programmer where
the CPU and the
GPU
can use the same program variable:
- Unified memory =
a memory model presented to
the CUDA programmer where
the CPU and the
GPU
can use the same program variable:
- How does
unified memory work:
- Data stored
in unified memory is
managed by
the CUDA system software:
- When code running on a
CPU
accesses data
allocated as
CUDA managed data,
the CUDA system software
takes care of
migrating (= transfering)
the data
to the host memory
- When code running on a
GPU
accesses data
allocated as
CUDA managed data,
the CUDA system software
takes care of
migrating (= transfering)
the data
to the device memory
- When code running on a
CPU
accesses data
allocated as
CUDA managed data,
the CUDA system software
takes care of
migrating (= transfering)
the data
to the host memory
- Data stored
in unified memory is
managed by
the CUDA system software:
- Effect of
the unified memory programming
model:
- A variable defined a "CUDA managed" is effectively shared by the CPU code and GPU code !!!
- Syntax to define
CUDA managed data:
__managed__ global-C-varible-definitionExample:
__managed__ int x; __managed__ float A[100];
Note:
- __manage__ can only used for defining shared global variables !!!
- Consider the following
failed
example CUDA program
that shares the
global variable int x:
__managed__ int x; // Defines shared variable !!! __global__ void GPU_func( ) { printf("GPU sees x = %d\n", x); x = 4444; } int main() { x = 1234; GPU_func<<< 1, 1 >>>( ); // Run 1 thread printf("CPU sees x = %d\n", x); cudaDeviceSynchronize( ); // Wait for GPU function to print return 0; }Purpose of the program:
- CPU main( ) function set x = 1234
- GPU GPU_func( ) function
can see that x = 1234
- GPU GPU_func( ) function will set x = 4444
- CPU main( ) function should see that x = 4444 !!!
- Example Program:
(Demo above code)
/home/cs355001/demo/CUDA/2-unified-mem/shared-global-fail Output: CPU sees x = 1234 GPU sees x = 1234 // What's wrong ???Expanation of the observed phenomenom:
- Recall that
launching a kernel is
asynchronous:
- The kernel launch returns immediately (without waiting for the kernel execution to finish)
- The following is
a time-progression of
the events in the
program
(on the CPU and GPU:
- Initialization:
CPU program starts
execution (GPU program is
inactive !!)
The shared data x resides in the host memory
- CPU statement
x = 1234 updates
the shared variable
x:
- CPU statement
GPU_func<<< 1,1 >>>( )
launches
the CPU function:
The call GPU_func<<< 1,1 >>>( ) will return immediately and the CPU continues with the next statement !!!
- CPU continues and
prints first (= before GPU):
x = 1234:
Now the CPU executes cudaDeviceSynchronize( ) and will wait for the GPU function to complete !!!
- The GPU function runs and
uses the
shared variable x:
- First, x is transfered to the device memory
- Then, the GPU code can use the variable x (to print)
Scehmatically:
- The GPU function then
updates
x = 4444:
- The GPU function completes and
cudaDeviceSynchronize( ) can
finally continue:
That's why you see:
CPU sees x = 1234 // CPU prints OLD value FIRST !! GPU sees x = 1234 // That's why !! - Initialization:
CPU program starts
execution (GPU program is
inactive !!)
- Recall that
launching a kernel is
asynchronous:
- Programming rule when
using unified memory:
- The CPU function
first
launches
the GPU kernel function
to perform
(parallel) computation
- If the
CPU function wants to
use the
result of the
computation
(computed by the GPU kernel function,
then:
- The CPU function must call cudaDeviceSynchronize( ) and wait until the GPU kernel function to complete before using the results !!!
Example:
__managed__ int x; // Defines shared variable !!! __global__ void GPU_func( ) { printf("GPU sees x = %d\n", x); x = 4444; } int main() { x = 1234; GPU_func<<< 1, 1 >>>( ); // Start GPU function cudaDeviceSynchronize(); // Wait until GPU kernel function finishes !! printf("CPU sees x = %d\n", x); // Now obtain the result !! return 0; }
- The CPU function
first
launches
the GPU kernel function
to perform
(parallel) computation
- Example Program:
(Demo above code)
/home/cs355001/demo/CUDA/2-unified-mem/shared-global Output: GPU sees x = 1234 // GPU prints FIRST !!! CPU sees x = 4444 // CPU waits until GPU is done !!!