diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index b3412663afa7..d4ac3251169b 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -148,9 +148,13 @@ mlir::Value CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
     return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.permlane64")
         .getScalarVal();
   }
-  case AMDGPU::BI__builtin_amdgcn_readlane:
+  case AMDGPU::BI__builtin_amdgcn_readlane: {
+    return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.readlane")
+        .getScalarVal();
+  }
   case AMDGPU::BI__builtin_amdgcn_readfirstlane: {
-    llvm_unreachable("readlane_* NYI");
+    return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.readfirstlane")
+        .getScalarVal();
   }
   case AMDGPU::BI__builtin_amdgcn_div_fixup:
   case AMDGPU::BI__builtin_amdgcn_div_fixupf:
diff --git a/clang/test/CIR/CodeGen/HIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGen/HIP/builtins-amdgcn.hip
index 9356373b13e1..b70f64e87c7d 100644
--- a/clang/test/CIR/CodeGen/HIP/builtins-amdgcn.hip
+++ b/clang/test/CIR/CodeGen/HIP/builtins-amdgcn.hip
@@ -285,3 +285,23 @@ __device__ void test_div_fmas_f64(double* out, double a, double b, double c, int
 __device__ void test_ds_swizzle_i32(int* out, int a) {
   *out = __builtin_amdgcn_ds_swizzle(a, 32);
 }
+
+// CIR-LABEL: @_Z13test_readlanePiii
+// CIR: cir.llvm.intrinsic "amdgcn.readlane" {{.*}} : (!s32i, !s32i) -> !s32i
+// LLVM: define{{.*}} void @_Z13test_readlanePiii
+// LLVM: call i32 @llvm.amdgcn.readlane.i32(i32 %{{.*}}, i32 %{{.*}})
+// OGCG: define{{.*}} void @_Z13test_readlanePiii
+// OGCG: call i32 @llvm.amdgcn.readlane.i32(i32 %{{.*}}, i32 %{{.*}})
+__device__ void test_readlane(int* out, int a, int b) {
+  *out = __builtin_amdgcn_readlane(a, b);
+}
+
+// CIR-LABEL: @_Z18test_readfirstlanePii
+// CIR: cir.llvm.intrinsic "amdgcn.readfirstlane" {{.*}} : (!s32i) -> !s32i
+// LLVM: define{{.*}} void @_Z18test_readfirstlanePii
+// LLVM: call i32 @llvm.amdgcn.readfirstlane.i32(i32 %{{.*}})
+// OGCG: define{{.*}} void @_Z18test_readfirstlanePii
+// OGCG: call i32 @llvm.amdgcn.readfirstlane.i32(i32 %{{.*}})
+__device__ void test_readfirstlane(int* out, int a) {
+  *out = __builtin_amdgcn_readfirstlane(a);
+}
diff --git a/clang/test/CIR/CodeGen/OpenCL/builtins_amdgcn.cl b/clang/test/CIR/CodeGen/OpenCL/builtins_amdgcn.cl
index b1d62d618b95..c602b7405155 100644
--- a/clang/test/CIR/CodeGen/OpenCL/builtins_amdgcn.cl
+++ b/clang/test/CIR/CodeGen/OpenCL/builtins_amdgcn.cl
@@ -298,3 +298,23 @@ void test_div_fmas_f64(global double* out, double a, double b, double c, int d)
 void test_ds_swizzle(global int* out, int a) {
   *out = __builtin_amdgcn_ds_swizzle(a, 32);
 }
+
+// CIR-LABEL: @test_readlane
+// CIR: cir.llvm.intrinsic "amdgcn.readlane" {{.*}} : (!s32i, !s32i) -> !s32i
+// LLVM: define{{.*}} void @test_readlane
+// LLVM: call i32 @llvm.amdgcn.readlane.i32(i32 %{{.*}}, i32 %{{.*}})
+// OGCG: define{{.*}} void @test_readlane
+// OGCG: call i32 @llvm.amdgcn.readlane.i32(i32 %{{.*}}, i32 %{{.*}})
+void test_readlane(global int* out, int a, int b) {
+  *out = __builtin_amdgcn_readlane(a, b);
+}
+
+// CIR-LABEL: @test_readfirstlane
+// CIR: cir.llvm.intrinsic "amdgcn.readfirstlane" {{.*}} : (!s32i) -> !s32i
+// LLVM: define{{.*}} void @test_readfirstlane
+// LLVM: call i32 @llvm.amdgcn.readfirstlane.i32(i32 %{{.*}})
+// OGCG: define{{.*}} void @test_readfirstlane
+// OGCG: call i32 @llvm.amdgcn.readfirstlane.i32(i32 %{{.*}})
+void test_readfirstlane(global int* out, int a) {
+  *out = __builtin_amdgcn_readfirstlane(a);
+}