# Advanced OpenMP presentation

README.md in `HPCTrainingExamples/Pragma_Examples/OpenMP/CXX/memory_pragmas` from Training Exercises repository

## Memory Pragmas

Setup your environment

```
module load amdclang
   or
export CXX=<C++ Compiler>
      Example: export CXX=amdclang++

export LIBOMPTARGET_INFO=-1
export LIBOMPTARGET_KERNEL_TRACE=[1,2]
export OMP_TARGET_OFFLOAD=MANDATORY
```

You can also be more selective in the output generated by using the individual bit masks

```
export LIBOMPTARGET_INFO=$((0x01 | 0x02 | 0x04 | 0x08 | 0x10 | 0x20))
```

The first example code uses just a single pragma with a map clause at the computational loop.

```
   mem1.cc:#pragma omp target teams distribute parallel for simd map(to: x[0:n], y[0:n]) map(from: z[0:n])
```

Examine this code and then compile and run. There is a map clause on pragma line just before computational loop

```
mkdir build && cd build
cmake ..
make 
./mem1
```

You should get some output like the following

```
Libomptarget device 0 info: Entering OpenMP kernel at mem1.cc:89:1 with 5 arguments:
Libomptarget device 0 info: firstprivate(n)[4] (implicit)
Libomptarget device 0 info: from(z[0:n])[80000]
Libomptarget device 0 info: firstprivate(a)[8] (implicit)
Libomptarget device 0 info: to(x[0:n])[80000]
Libomptarget device 0 info: to(y[0:n])[80000]
Libomptarget device 0 info: Creating new map entry with HstPtrBase=0x0000000001772200, ... 
Libomptarget device 0 info: Creating new map entry with HstPtrBase=0x000000000174b0e0, ... 
Libomptarget device 0 info: Copying data from host to device, HstPtr=0x000000000174b0e0, ...
Libomptarget device 0 info: Creating new map entry with HstPtrBase=0x000000000175e970, ...
Libomptarget device 0 info: Copying data from host to device, HstPtr=0x000000000175e970, ...
Libomptarget device 0 info: Mapping exists with HstPtrBegin=0x0000000001772200, ...
Libomptarget device 0 info: Mapping exists with HstPtrBegin=0x000000000174b0e0, ...
Libomptarget device 0 info: Mapping exists with HstPtrBegin=0x000000000175e970, ...
Libomptarget device 0 info: Mapping exists with HstPtrBegin=0x000000000175e970, ...
Libomptarget device 0 info: Mapping exists with HstPtrBegin=0x000000000174b0e0, ...
Libomptarget device 0 info: Mapping exists with HstPtrBegin=0x0000000001772200, ...
Libomptarget device 0 info: Copying data from device to host, TgtPtr=0x00007f617c420000, ...
Libomptarget device 0 info: Removing map entry with HstPtrBegin=0x000000000175e970, ...
Libomptarget device 0 info: Removing map entry with HstPtrBegin=0x000000000174b0e0, ...
Libomptarget device 0 info: Removing map entry with HstPtrBegin=0x0000000001772200, ...
-Timing in Seconds: min=0.010115, max=0.010115, avg=0.010115
-Overall time is 0.010505
Last Value: z[9999]=7.000000
```

Explore examples 2 through 5 and observe the output produced when the `LIBOMPTARGET_INFO` environment variable is set.

Mem2 pattern : Add enter/exit data alloc/delete when memory is created/freed

   After new
   `mem2.cc:#pragma omp target enter data map(alloc: x[0:n], y[0:n], z[0:n])`

   Loop around computational loop and keep map on computational loop. The map to/from should check if the
   data exists. If not, it will allocate/delete it. Then it will do the copies to and from. This will
   increment the Reference Counter and decrement it at end of loop.
   `mem2.cc:#pragma omp target teams distribute parallel for simd map(always to: x[0:n], y[0:n]) map(always from: z[0:n])`

   Before delete
   `mem2.cc:#pragma omp target exit data map(delete: x[0:n], y[0:n], z[0:n])`

Mem3 pattern: Replacing map to/from with updates to bypass unneeded device memory check

   After new
   `mem3.cc:#pragma omp target enter data map(alloc: x[0:n], y[0:n], z[0:n])`

   Before computational loop. Data should be copied. Reference counter should not change.
   `mem3.cc:#pragma omp target update to (x[0:n], y[0:n])`
   `mem3.cc:#pragma omp target teams distribute parallel for simd`

   After computational loop
   `mem3.cc:#pragma omp target update from (z[0:n])`

   Before delete
   `mem3.cc:#pragma omp target exit data map(delete: x[0:n], y[0:n], z[0:n])`

Mem4 pattern: Replacing delete with release to use Reference Counting

```
   mem4.cc:#pragma omp target enter data map(alloc: x[0:n], y[0:n], z[0:n])
   mem4.cc:#pragma omp target exit data map(release: x[0:n], y[0:n], z[0:n])
   mem4.cc:#pragma omp target teams distribute parallel for simd map(always to: x[0:n], y[0:n]) map(always from: z[0:n])
```

Mem5 pattern: Using enter data map to/from alloc/delete to reduce memory copies

```
   mem5.cc:#pragma omp target enter data map(to: x[0:n], y[0:n]) map(alloc: z[0:n])
   mem5.cc:#pragma omp target exit data map(from: z[0:n]) map(delete: x[0:n], y[0:n])
   mem5.cc:#pragma omp target teams distribute parallel for simd map(to:x[0:n], y[0:n]) map(from: z[0:n])
```

### One solution that miminizes data transfer

Mem6 pattern: Using enter data alloc/delete with update clause at end

```
   mem6.cc:#pragma omp target enter data map(alloc: x[0:n], y[0:n], z[0:n])
   mem6.cc:#pragma omp target teams distribute parallel for simd
   mem6.cc:#pragma omp target update from(z[0])
   mem6.cc:#pragma omp target exit data map(delete: x[0:n], y[0:n], z[0:n])
   mem6.cc:#pragma omp target teams distribute parallel for simd
```

### Unified Shared Memory

Mem7 pattern: Using Unified Shared Memory to automatically move data

```
   mem7.cc:#pragma omp requires unified_shared_memory
   mem7.cc:#pragma omp target teams distribute parallel for simd
   mem7.cc:#pragma omp target teams distribute parallel for simd
```

For this example, `HSA_XNACK=1` needs to be set

```
export HSA_XNACK=1
make mem7
./mem7
```


### Unified Shared Memory with backwards compatibility

Mem8 pattern: Demonstrating Unified Shared Memory with maps for backward compatibility

```
   set HSA_XNACK=1 at runtime

   mem8.cc:#pragma omp requires unified_shared_memory
   mem8.cc:#pragma omp target enter data map(alloc: x[0:n], y[0:n], z[0:n])
   mem8.cc:#pragma omp target teams distribute parallel for simd
   mem8.cc:#pragma omp target update from(z[0])
   mem8.cc:#pragma omp target exit data map(delete: x[0:n], y[0:n], z[0:n])
   mem8.cc:#pragma omp target teams distribute parallel for simd
```

### APU Code -- Unified Address in OpenMP

We now switch to how unified address programming would look in other languages. The language we will work 
with the most will be OpenMP. We'll start by looking at the unified address code shown in slide 31 and 32. 
It is also in the mem11.cc file in the directory given below. You should also compare
it to the original GPU code using explicit memory management in mem1.cc through mem6.cc.

We'll now run the unified address example if we have access to an MI300A GPU. If you don't have access to an MI300A, we'll
also run nearly the same code in mem7.cc with managed memory on the MI200 series GPUs. We'll be looking at all of the versions of this code in the Advanced OpenMP presentation.

```
cd ~/HPCTrainingExamples/Pragma_Examples/OpenMP/CXX/memory_pragmas
module load amdclang
mkdir build && cd build
cmake ..
make mem11
./mem11
```

