diff --git a/lib/CodeGen/CGOpenMPRuntime.cpp b/lib/CodeGen/CGOpenMPRuntime.cpp
index 1f1db83..5957e59 100644
--- a/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -6041,7 +6041,7 @@
                               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/test/OpenMP/target_uses_allocators.c b/test/OpenMP/target_uses_allocators.c
index eab2026..0352a58 100644
--- a/test/OpenMP/target_uses_allocators.c
+++ b/test/OpenMP/target_uses_allocators.c
@@ -64,6 +64,35 @@
   {}
 }
 
+typedef enum omp_memspace_handle_t {
+  omp_default_mem_space = 0,
+  omp_large_cap_mem_space = 1,
+  omp_const_mem_space = 2,
+  omp_high_bw_mem_space = 3,
+  omp_low_lat_mem_space = 4,
+  llvm_omp_target_host_mem_space = 100,
+  llvm_omp_target_shared_mem_space = 101,
+  llvm_omp_target_device_mem_space = 102,
+  KMP_MEMSPACE_MAX_HANDLE = __UINTPTR_MAX__
+} omp_memspace_handle_t;
+
+extern omp_allocator_handle_t
+omp_init_allocator(omp_memspace_handle_t memspace, int ntraits,
+                   const omp_alloctrait_t traits[]);
+
+void *omp_aligned_alloc(unsigned long alignment, unsigned long size,
+                        omp_allocator_handle_t allocator);
+extern void * omp_alloc(int size, omp_allocator_handle_t a);
+#define N 1024
+
+void foo() {
+  int errors = 0;
+  omp_memspace_handle_t  memspace = omp_default_mem_space;
+  omp_alloctrait_t       traits[1] = {{omp_atk_alignment, 64}};
+  omp_allocator_handle_t alloc = omp_init_allocator(memspace,1,traits);
+  #pragma omp target map(tofrom: errors) uses_allocators(alloc(traits))
+  { }
+}
 #endif
 
 // CHECK: %[[#R0:]] = call i32 @__kmpc_global_thread_num(ptr @1)
@@ -140,3 +169,15 @@
 // CHECK: [[ALLOCATOR:%.+]] = load i64, ptr [[MY_ALLOCATOR_ADDR]],
 // CHECK: [[CONV:%.+]] = inttoptr i64 [[ALLOCATOR]] to ptr
 // CHECK: call void @__kmpc_destroy_allocator(i32 %{{.+}}, ptr [[CONV]])
+
+// CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
+// CHECK: [[MY_ALLOCATOR_ADDR:%alloc]] = alloca i64,
+// CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
+// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 1, 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: [[CONV1:%.+]] = inttoptr i64 [[ALLOCATOR]] to ptr
+// CHECK: call void @__kmpc_destroy_allocator(i32 %{{.+}}, ptr [[CONV1]])