diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -6034,15 +6034,14 @@
   AllocatorTraitsLVal = CGF.MakeAddrLValue(Addr, CGF.getContext().VoidPtrTy,
                                            AllocatorTraitsLVal.getBaseInfo(),
                                            AllocatorTraitsLVal.getTBAAInfo());
-  llvm::Value *Traits =
-      CGF.EmitLoadOfScalar(AllocatorTraitsLVal, AllocatorTraits->getExprLoc());
+  llvm::Value *Traits = Addr.getPointer();
 
   llvm::Value *AllocatorVal =
       CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
                               CGM.getModule(), OMPRTL___kmpc_init_allocator),
                           {ThreadId, MemSpaceHandle, NumTraits, Traits});
   // Store to allocator.
-  CGF.EmitVarDecl(*cast<VarDecl>(
+  CGF.EmitAutoVarAlloca(*cast<VarDecl>(
       cast<DeclRefExpr>(Allocator->IgnoreParenImpCasts())->getDecl()));
   LValue AllocatorLVal = CGF.EmitLValue(Allocator->IgnoreParenImpCasts());
   AllocatorVal =
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -3103,8 +3103,13 @@
     return nullptr;
   SmallVector<Sema::UsesAllocatorsData, 4> Data;
   do {
+    CXXScopeSpec SS;
+    Token Replacement;
     ExprResult Allocator =
-        getLangOpts().CPlusPlus ? ParseCXXIdExpression() : ParseExpression();
+        getLangOpts().CPlusPlus
+            ? ParseCXXIdExpression()
+            : tryParseCXXIdExpression(SS, /*isAddressOfOperand=*/false,
+                                      Replacement);
     if (Allocator.isInvalid()) {
       SkipUntil(tok::comma, tok::r_paren, tok::annot_pragma_openmp_end,
                 StopBeforeMatch);
diff --git a/clang/test/OpenMP/target_parallel_for_simd_uses_allocators_codegen.cpp b/clang/test/OpenMP/target_parallel_for_simd_uses_allocators_codegen.cpp
--- a/clang/test/OpenMP/target_parallel_for_simd_uses_allocators_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_for_simd_uses_allocators_codegen.cpp
@@ -79,8 +79,7 @@
 // CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
 // CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
 // CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS_ADDR]])
 // CHECK: [[CONV:%.+]] = ptrtoint ptr [[ALLOCATOR]] to i64
 // CHECK: store i64 [[CONV]], ptr [[MY_ALLOCATOR_ADDR]],
 
diff --git a/clang/test/OpenMP/target_parallel_for_uses_allocators_codegen.cpp b/clang/test/OpenMP/target_parallel_for_uses_allocators_codegen.cpp
--- a/clang/test/OpenMP/target_parallel_for_uses_allocators_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_for_uses_allocators_codegen.cpp
@@ -79,8 +79,7 @@
 // CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
 // CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
 // CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS_ADDR]])
 // CHECK: [[CONV:%.+]] = ptrtoint ptr [[ALLOCATOR]] to i64
 // CHECK: store i64 [[CONV]], ptr [[MY_ALLOCATOR_ADDR]],
 
diff --git a/clang/test/OpenMP/target_parallel_uses_allocators_codegen.cpp b/clang/test/OpenMP/target_parallel_uses_allocators_codegen.cpp
--- a/clang/test/OpenMP/target_parallel_uses_allocators_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_uses_allocators_codegen.cpp
@@ -78,8 +78,7 @@
 // CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
 // CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
 // CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS_ADDR]])
 // CHECK: [[CONV:%.+]] = ptrtoint ptr [[ALLOCATOR]] to i64
 // CHECK: store i64 [[CONV]], ptr [[MY_ALLOCATOR_ADDR]],
 
diff --git a/clang/test/OpenMP/target_simd_uses_allocators_codegen.cpp b/clang/test/OpenMP/target_simd_uses_allocators_codegen.cpp
--- a/clang/test/OpenMP/target_simd_uses_allocators_codegen.cpp
+++ b/clang/test/OpenMP/target_simd_uses_allocators_codegen.cpp
@@ -79,8 +79,7 @@
 // CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
 // CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
 // CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS_ADDR]])
 // CHECK: [[CONV:%.+]] = ptrtoint ptr [[ALLOCATOR]] to i64
 // CHECK: store i64 [[CONV]], ptr [[MY_ALLOCATOR_ADDR]],
 
diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_uses_allocators_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_uses_allocators_codegen.cpp
--- a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_uses_allocators_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_uses_allocators_codegen.cpp
@@ -79,8 +79,7 @@
 // CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
 // CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
 // CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS_ADDR]])
 // CHECK: [[CONV:%.+]] = ptrtoint ptr [[ALLOCATOR]] to i64
 // CHECK: store i64 [[CONV]], ptr [[MY_ALLOCATOR_ADDR]],
 
diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_uses_allocators_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_uses_allocators_codegen.cpp
--- a/clang/test/OpenMP/target_teams_distribute_parallel_for_uses_allocators_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_uses_allocators_codegen.cpp
@@ -79,8 +79,7 @@
 // CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
 // CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
 // CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS_ADDR]])
 // CHECK: [[CONV:%.+]] = ptrtoint ptr [[ALLOCATOR]] to i64
 // CHECK: store i64 [[CONV]], ptr [[MY_ALLOCATOR_ADDR]],
 
diff --git a/clang/test/OpenMP/target_teams_distribute_simd_uses_allocators_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_simd_uses_allocators_codegen.cpp
--- a/clang/test/OpenMP/target_teams_distribute_simd_uses_allocators_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_simd_uses_allocators_codegen.cpp
@@ -79,8 +79,7 @@
 // CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
 // CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
 // CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS_ADDR]])
 // CHECK: [[CONV:%.+]] = ptrtoint ptr [[ALLOCATOR]] to i64
 // CHECK: store i64 [[CONV]], ptr [[MY_ALLOCATOR_ADDR]],
 
diff --git a/clang/test/OpenMP/target_teams_distribute_uses_allocators_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_uses_allocators_codegen.cpp
--- a/clang/test/OpenMP/target_teams_distribute_uses_allocators_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_uses_allocators_codegen.cpp
@@ -79,8 +79,7 @@
 // CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
 // CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
 // CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS_ADDR]])
 // CHECK: [[CONV:%.+]] = ptrtoint ptr [[ALLOCATOR]] to i64
 // CHECK: store i64 [[CONV]], ptr [[MY_ALLOCATOR_ADDR]],
 
diff --git a/clang/test/OpenMP/target_teams_uses_allocators_codegen.cpp b/clang/test/OpenMP/target_teams_uses_allocators_codegen.cpp
--- a/clang/test/OpenMP/target_teams_uses_allocators_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_uses_allocators_codegen.cpp
@@ -78,8 +78,7 @@
 // CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
 // CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
 // CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS_ADDR]])
 // CHECK: [[CONV:%.+]] = ptrtoint ptr [[ALLOCATOR]] to i64
 // CHECK: store i64 [[CONV]], ptr [[MY_ALLOCATOR_ADDR]],
 
diff --git a/clang/test/OpenMP/target_uses_allocators.c b/clang/test/OpenMP/target_uses_allocators.c
--- a/clang/test/OpenMP/target_uses_allocators.c
+++ b/clang/test/OpenMP/target_uses_allocators.c
@@ -2,6 +2,8 @@
 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50  -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -verify -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t %s -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -verify -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t %s -emit-llvm -o - | FileCheck %s
 
 #ifndef HEADER
 #define HEADER
@@ -19,9 +21,27 @@
   KMP_ALLOCATOR_MAX_HANDLE = __UINTPTR_MAX__
 } omp_allocator_handle_t;
 
+typedef enum omp_alloctrait_key_t { omp_atk_sync_hint = 1,
+                                    omp_atk_alignment = 2,
+                                    omp_atk_access = 3,
+                                    omp_atk_pool_size = 4,
+                                    omp_atk_fallback = 5,
+                                    omp_atk_fb_data = 6,
+                                    omp_atk_pinned = 7,
+                                    omp_atk_partition = 8
+} omp_alloctrait_key_t;
+
+typedef struct omp_alloctrait_t {
+  omp_alloctrait_key_t key;
+  __UINTPTR_TYPE__ value;
+} omp_alloctrait_t;
+
+
 // CHECK: define {{.*}}[[FIE:@.+]]()
 void fie(void) {
   int x;
+  omp_allocator_handle_t my_allocator;
+  omp_alloctrait_t traits[10];
   #pragma omp target uses_allocators(omp_null_allocator) allocate(omp_null_allocator: x) firstprivate(x)
   {}
   #pragma omp target uses_allocators(omp_default_mem_alloc) allocate(omp_default_mem_alloc: x) firstprivate(x)
@@ -40,6 +60,8 @@
   {}
   #pragma omp target uses_allocators(omp_thread_mem_alloc) allocate(omp_thread_mem_alloc: x) firstprivate(x) // expected-warning {{allocator with the 'thread' trait access has unspecified behavior on 'target' directive}}
   {}
+#pragma omp target uses_allocators(omp_null_allocator, omp_thread_mem_alloc, my_allocator(traits))
+  {}
 }
 
 #endif
@@ -106,3 +128,15 @@
 // CHECK-NEXT: %[[#R1:]] = load i32, ptr %x.addr, align 4
 // CHECK-NEXT: store i32 %[[#R1]], ptr %.x..void.addr, align 4
 // CHECK-NEXT: call void @__kmpc_free(i32 %[[#R0]], ptr %.x..void.addr, ptr inttoptr (i64 8 to ptr))
+
+// CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
+// CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
+// CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
+// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS_ADDR]])
+// CHECK: [[CONV:%.+]] = ptrtoint ptr [[ALLOCATOR]] to i64
+// CHECK: store i64 [[CONV]], ptr [[MY_ALLOCATOR_ADDR]],
+
+// Destroy allocator upon exit from the region.
+// CHECK: [[ALLOCATOR:%.+]] = load i64, ptr [[MY_ALLOCATOR_ADDR]],
+// CHECK: [[CONV:%.+]] = inttoptr i64 [[ALLOCATOR]] to ptr
+// CHECK: call void @__kmpc_destroy_allocator(i32 %{{.+}}, ptr [[CONV]])
diff --git a/clang/test/OpenMP/target_uses_allocators_codegen.cpp b/clang/test/OpenMP/target_uses_allocators_codegen.cpp
--- a/clang/test/OpenMP/target_uses_allocators_codegen.cpp
+++ b/clang/test/OpenMP/target_uses_allocators_codegen.cpp
@@ -78,8 +78,7 @@
 // CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
 // CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
 // CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS_ADDR]])
 // CHECK: [[CONV:%.+]] = ptrtoint ptr [[ALLOCATOR]] to i64
 // CHECK: store i64 [[CONV]], ptr [[MY_ALLOCATOR_ADDR]],
 
diff --git a/openmp/libomptarget/test/mapping/target_uses_allocator.c b/openmp/libomptarget/test/mapping/target_uses_allocator.c
new file mode 100755
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/target_uses_allocator.c
@@ -0,0 +1,56 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+#define N 1024
+
+int test_omp_aligned_alloc_on_device() {
+  int errors = 0;
+
+  omp_memspace_handle_t memspace = omp_default_mem_space;
+  omp_alloctrait_t traits[2] = {{omp_atk_alignment, 64}, {omp_atk_access, 64}};
+  omp_allocator_handle_t alloc =
+      omp_init_allocator(omp_default_mem_space, 1, traits);
+
+#pragma omp target map(tofrom : errors) uses_allocators(alloc(traits))
+  {
+    int *x;
+    int not_correct_array_values = 0;
+
+    x = (int *)omp_aligned_alloc(64, N * sizeof(int), alloc);
+    if (x == NULL) {
+      errors++;
+    } else {
+#pragma omp parallel for simd simdlen(16) aligned(x : 64)
+      for (int i = 0; i < N; i++) {
+        x[i] = i;
+      }
+
+#pragma omp parallel for simd simdlen(16) aligned(x : 64)
+      for (int i = 0; i < N; i++) {
+        if (x[i] != i) {
+#pragma omp atomic write
+          not_correct_array_values = 1;
+        }
+      }
+      if (not_correct_array_values) {
+        errors++;
+      }
+      omp_free(x, alloc);
+    }
+  }
+
+  omp_destroy_allocator(alloc);
+
+  return errors;
+}
+
+int main() {
+  int errors = 0;
+  if (test_omp_aligned_alloc_on_device())
+    printf("FAILE\n");
+  else
+    // CHECK: PASSED
+    printf("PASSED\n");
+}