-
+### Unified shared memory
+Unified Shared Memory (USM) provides a pointer-based approach to memory
+management. To implement USM, fulfill the following system requirements along
+with Xnack capability.
-
+ return 0;
+}
+$ clang++ -O2 -target x86_64-pc-linux-gnu -fopenmp --offload-arch=gfx90a:xnack+ parallel_for.cpp
+$ HSA_XNACK=1 ./a.out
+```
+
+In the above code example, pointer “a” is not mapped in the target region, while
+pointer “b” is. Both are valid pointers on the GPU device and passed by-value to
+the kernel implementing the target region. This means the pointer values on the
+host and the device are the same.
+
+The difference between the memory pages pointed to by these two variables is
+that the pages pointed by “a” are in fine-grain memory, while the pages pointed
+to by “b” are in coarse-grain memory during and after the execution of the
+target region. This is accomplished in the OpenMP runtime library with calls to
+the ROCr runtime to set the pages pointed by “b” as coarse grain.
+
+### OMPT target support
+
+The OpenMP runtime in ROCm implements a subset of the OMPT device APIs, as
+described in the OpenMP specification document. These APIs allow first-party
+tools to examine the profile and kernel traces that execute on a device. A tool
+can register callbacks for data transfer and kernel dispatch entry points or use
+APIs to start and stop tracing for device-related activities such as data
+transfer and kernel dispatch timings and associated metadata. If device tracing
+is enabled, trace records for device activities are collected during program
+execution and returned to the tool using the APIs described in the
+specification.
+
+The following example demonstrates how a tool uses the supported OMPT target
+APIs. The `README` in `/opt/rocm/llvm/examples/tools/ompt` outlines the steps to
+be followed, and the provided example can be run as shown below:
+
+```bash
+cd $ROCM_PATH/share/openmp-extras/examples/tools/ompt/veccopy-ompt-target-tracing
+sudo make run
+```
+
+The file `veccopy-ompt-target-tracing.c` simulates how a tool initiates device
+activity tracing. The file `callbacks.h` shows the callbacks registered and
+implemented by the tool.
+
+### Floating point atomic operations
+
+The MI200-series GPUs support the generation of hardware floating-point atomics
+using the OpenMP atomic pragma. The support includes single- and
+double-precision floating-point atomic operations. The programmer must ensure
+that the memory subjected to the atomic operation is in coarse-grain memory by
+mapping it explicitly with the help of map clauses when not implicitly mapped by
+the compiler as per the [OpenMP
+specifications](https://www.openmp.org/specifications/). This makes these
+hardware floating-point atomic instructions “fast,” as they are faster than
+using a default compare-and-swap loop scheme, but at the same time “unsafe,” as
+they are not supported on fine-grain memory. The operation in
+`unified_shared_memory` mode also requires programmers to map the memory
+explicitly when not implicitly mapped by the compiler.
+
+To request fast floating-point atomic instructions at the file level, use
+compiler flag `-munsafe-fp-atomics` or a hint clause on a specific pragma:
+
+```bash
+double a = 0.0;
+#pragma omp atomic hint(AMD_fast_fp_atomics)
+a = a + 1.0;
+```
+
+:::{note}
+`AMD_unsafe_fp_atomics` is an alias for `AMD_fast_fp_atomics`, and
+`AMD_safe_fp_atomics` is implemented with a compare-and-swap loop.
+:::
+
+To disable the generation of fast floating-point atomic instructions at the file
+level, build using the option `-msafe-fp-atomics` or use a hint clause on a
+specific pragma:
+
+```bash
+double a = 0.0;
+#pragma omp atomic hint(AMD_safe_fp_atomics)
+a = a + 1.0;
+```
+
+The hint clause value always has a precedence over the compiler flag, which
+allows programmers to create atomic constructs with a different behavior than
+the rest of the file.
+
+See the example below, where the user builds the program using
+`-msafe-fp-atomics` to select a file-wide “safe atomic” compilation. However,
+the fast atomics hint clause over variable “a” takes precedence and operates on
+“a” using a fast/unsafe floating-point atomic, while the variable “b” in the
+absence of a hint clause is operated upon using safe floating-point atomics as
+per the compiler flag.
+
+```bash
+double a = 0.0;.
+#pragma omp atomic hint(AMD_fast_fp_atomics)
+a = a + 1.0;
+
+double b = 0.0;
+#pragma omp atomic
+b = b + 1.0;
+```
+
+### AddressSanitizer tool
+
+AddressSanitizer (ASan) is a memory error detector tool utilized by applications to
+detect various errors ranging from spatial issues such as out-of-bound access to
+temporal issues such as use-after-free. The AOMP compiler supports ASan for AMD
+GPUs with applications written in both HIP and OpenMP.
+
+**Features supported on host platform (Target x86_64):**
+
+* Use-after-free
+* Buffer overflows
+* Heap buffer overflow
+* Stack buffer overflow
+* Global buffer overflow
+* Use-after-return
+* Use-after-scope
+* Initialization order bugs
+
+**Features supported on AMDGPU platform (`amdgcn-amd-amdhsa`):**
+
+* Heap buffer overflow
+* Global buffer overflow
-
+**Software (kernel/OS) requirements:** Unified Shared Memory support with Xnack
+capability. See the section on [Unified Shared Memory](#unified-shared-memory)
+for prerequisites and details on Xnack.
+**Example:**
+* Heap buffer overflow
+```bash
+void main() {
+....... // Some program statements
+....... // Some program statements
+#pragma omp target map(to : A[0:N], B[0:N]) map(from: C[0:N])
+{
+#pragma omp parallel for
+ for(int i =0 ; i < N; i++){
+ C[i+10] = A[i] + B[i];
+ } // end of for loop
+}
+....... // Some program statements
+}// end of main
+```
+See the complete sample code for heap buffer overflow
+[here](https://github.com/ROCm/aomp/blob/aomp-dev/examples/tools/asan/heap_buffer_overflow/openmp/vecadd-HBO.cpp).
+* Global buffer overflow
+```bash
+#pragma omp declare target
+ int A[N],B[N],C[N];
+#pragma omp end declare target
+void main(){
+...... // some program statements
+...... // some program statements
+#pragma omp target data map(to:A[0:N],B[0:N]) map(from: C[0:N])
+{
+#pragma omp target update to(A,B)
+#pragma omp target parallel for
+for(int i=0; i
-
-
-
-
-
-
-
-
-
-