Index: clang/lib/Sema/SemaCUDA.cpp
===================================================================
--- clang/lib/Sema/SemaCUDA.cpp
+++ clang/lib/Sema/SemaCUDA.cpp
@@ -718,6 +718,11 @@
   FunctionDecl *CurFn = dyn_cast<FunctionDecl>(CurContext);
   if (!CurFn)
     return;
+  if (getLangOpts().HIP) {
+    Method->addAttr(CUDADeviceAttr::CreateImplicit(Context));
+    Method->addAttr(CUDAHostAttr::CreateImplicit(Context));
+    return;
+  }
   CUDAFunctionTarget Target = IdentifyCUDATarget(CurFn);
   if (Target == CFT_Global || Target == CFT_Device) {
     Method->addAttr(CUDADeviceAttr::CreateImplicit(Context));
Index: clang/test/CodeGenCUDA/lambda.cu
===================================================================
--- /dev/null
+++ clang/test/CodeGenCUDA/lambda.cu
@@ -0,0 +1,26 @@
+// RUN: %clang_cc1 -x hip -emit-llvm -std=c++11 %s -o - \
+// RUN:   -triple x86_64-linux-gnu | FileCheck -check-prefix=HOST %s
+// RUN: %clang_cc1 -x hip -emit-llvm -std=c++11 %s -o - \
+// RUN:   -triple amdgcn-amd-amdhsa -fcuda-is-device \
+// RUN:   | FileCheck -check-prefix=DEV %s
+
+#include "Inputs/cuda.h"
+
+// HOST: @[[KERN:[0-9]+]] = private unnamed_addr constant [22 x i8] c"_Z1gIZ4mainEUlvE_EvT_\00"
+// HOST: define internal void @_Z1hIZ4mainEUlvE_EvT_
+// HOST: define internal void @_Z16__device_stub__gIZ4mainEUlvE_EvT_
+// HOST: @__hipRegisterFunction(i8** %0, i8* bitcast ({{.*}}@[[KERN]]
+// HOST-NOT: define{{.*}}@_ZZ4mainENKUlvE_clEv
+// DEV: define amdgpu_kernel void @_Z1gIZ4mainEUlvE_EvT_
+// DEV: define internal void @_ZZ4mainENKUlvE_clEv
+template<class F>
+__global__ void g(F f) { f(); }
+
+template<class F>
+void h(F f) { g<<<1,1>>>(f); }
+
+__device__ int a;
+
+int main(void) {
+  h([&](){ a=1;});
+}