diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index c05cdf1fa961..d6669230162c 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -326,7 +326,8 @@ mlir::Value CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
   case AMDGPU::BI__builtin_amdgcn_sqrtf:
   case AMDGPU::BI__builtin_amdgcn_sqrth:
   case AMDGPU::BI__builtin_amdgcn_sqrt_bf16: {
-    llvm_unreachable("sqrt_* NYI");
+    return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.sqrt")
+        .getScalarVal();
   }
   case AMDGPU::BI__builtin_amdgcn_rsq:
   case AMDGPU::BI__builtin_amdgcn_rsqf:
diff --git a/clang/test/CIR/CodeGen/HIP/builtins-amdgcn-gfx1250.hip b/clang/test/CIR/CodeGen/HIP/builtins-amdgcn-gfx1250.hip
index 4aa0d32521ba..2c2714220f2f 100644
--- a/clang/test/CIR/CodeGen/HIP/builtins-amdgcn-gfx1250.hip
+++ b/clang/test/CIR/CodeGen/HIP/builtins-amdgcn-gfx1250.hip
@@ -27,3 +27,14 @@ __device__ void test_rcp_bf16(__bf16* out, __bf16 a)
 {
   *out = __builtin_amdgcn_rcp_bf16(a);
 }
+
+// CIR-LABEL: @_Z14test_sqrt_bf16PDF16bDF16b
+// CIR: cir.llvm.intrinsic "amdgcn.sqrt" {{.*}} : (!cir.bf16) -> !cir.bf16
+// LLVM: define{{.*}} void @_Z14test_sqrt_bf16PDF16bDF16b
+// LLVM: call{{.*}} bfloat @llvm.amdgcn.sqrt.bf16(bfloat %{{.*}})
+// OGCG: define{{.*}} void @_Z14test_sqrt_bf16PDF16bDF16b
+// OGCG: call{{.*}} bfloat @llvm.amdgcn.sqrt.bf16(bfloat %{{.*}})
+__device__ void test_sqrt_bf16(__bf16* out, __bf16 a)
+{
+  *out = __builtin_amdgcn_sqrt_bf16(a);
+}
diff --git a/clang/test/CIR/CodeGen/HIP/builtins-amdgcn-vi.hip b/clang/test/CIR/CodeGen/HIP/builtins-amdgcn-vi.hip
index c49cefde8dff..1f8bd39c0e0d 100644
--- a/clang/test/CIR/CodeGen/HIP/builtins-amdgcn-vi.hip
+++ b/clang/test/CIR/CodeGen/HIP/builtins-amdgcn-vi.hip
@@ -76,3 +76,14 @@ __device__ void test_rcp_f16(_Float16* out, _Float16 a)
 {
   *out = __builtin_amdgcn_rcph(a);
 }
+
+// CIR-LABEL: @_Z13test_sqrt_f16PDF16_DF16_
+// CIR: cir.llvm.intrinsic "amdgcn.sqrt" {{.*}} : (!cir.f16) -> !cir.f16
+// LLVM: define{{.*}} void @_Z13test_sqrt_f16PDF16_DF16_
+// LLVM: call{{.*}} half @llvm.{{((amdgcn.){0,1})}}sqrt.f16(half %{{.*}})
+// OGCG: define{{.*}} void @_Z13test_sqrt_f16PDF16_DF16_
+// OGCG: call{{.*}} half @llvm.{{((amdgcn.){0,1})}}sqrt.f16(half %{{.*}})
+__device__ void test_sqrt_f16(_Float16* out, _Float16 a)
+{
+  *out = __builtin_amdgcn_sqrth(a);
+}
diff --git a/clang/test/CIR/CodeGen/HIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGen/HIP/builtins-amdgcn.hip
index b12e2aa301de..1745b8c1f9cf 100644
--- a/clang/test/CIR/CodeGen/HIP/builtins-amdgcn.hip
+++ b/clang/test/CIR/CodeGen/HIP/builtins-amdgcn.hip
@@ -365,3 +365,23 @@ __device__ void test_rcp_f32(float* out, float a) {
 __device__ void test_rcp_f64(double* out, double a) {
   *out = __builtin_amdgcn_rcp(a);
 }
+
+// CIR-LABEL: @_Z13test_sqrt_f32Pff
+// CIR: cir.llvm.intrinsic "amdgcn.sqrt" {{.*}} : (!cir.float) -> !cir.float
+// LLVM: define{{.*}} void @_Z13test_sqrt_f32Pff
+// LLVM: call{{.*}} float @llvm.amdgcn.sqrt.f32(float %{{.*}})
+// OGCG: define{{.*}} void @_Z13test_sqrt_f32Pff
+// OGCG: call{{.*}} float @llvm.amdgcn.sqrt.f32(float %{{.*}})
+__device__ void test_sqrt_f32(float* out, float a) {
+  *out = __builtin_amdgcn_sqrtf(a);
+}
+
+// CIR-LABEL: @_Z13test_sqrt_f64Pdd
+// CIR: cir.llvm.intrinsic "amdgcn.sqrt" {{.*}} : (!cir.double) -> !cir.double
+// LLVM: define{{.*}} void @_Z13test_sqrt_f64Pdd
+// LLVM: call{{.*}} double @llvm.amdgcn.sqrt.f64(double %{{.*}})
+// OGCG: define{{.*}} void @_Z13test_sqrt_f64Pdd
+// OGCG: call{{.*}} double @llvm.amdgcn.sqrt.f64(double %{{.*}})
+__device__ void test_sqrt_f64(double* out, double a) {
+  *out = __builtin_amdgcn_sqrt(a);
+}
diff --git a/clang/test/CIR/CodeGen/OpenCL/builtins-amdgcn-gfx1250.cl b/clang/test/CIR/CodeGen/OpenCL/builtins-amdgcn-gfx1250.cl
index aa272c230e1f..7d72125cafd0 100644
--- a/clang/test/CIR/CodeGen/OpenCL/builtins-amdgcn-gfx1250.cl
+++ b/clang/test/CIR/CodeGen/OpenCL/builtins-amdgcn-gfx1250.cl
@@ -27,3 +27,14 @@ void test_rcp_bf16(global __bf16* out, __bf16 a)
 {
   *out = __builtin_amdgcn_rcp_bf16(a);
 }
+
+// CIR-LABEL: @test_sqrt_bf16
+// CIR: cir.llvm.intrinsic "amdgcn.sqrt" {{.*}} : (!cir.bf16) -> !cir.bf16
+// LLVM: define{{.*}} void @test_sqrt_bf16
+// LLVM: call{{.*}} bfloat @llvm.amdgcn.sqrt.bf16(bfloat %{{.*}})
+// OGCG: define{{.*}} void @test_sqrt_bf16
+// OGCG: call{{.*}} bfloat @llvm.amdgcn.sqrt.bf16(bfloat %{{.*}})
+void test_sqrt_bf16(global __bf16* out, __bf16 a)
+{
+  *out = __builtin_amdgcn_sqrt_bf16(a);
+}
diff --git a/clang/test/CIR/CodeGen/OpenCL/builtins-amdgcn-vi.cl b/clang/test/CIR/CodeGen/OpenCL/builtins-amdgcn-vi.cl
index a7aa7192234a..021fe0050a11 100644
--- a/clang/test/CIR/CodeGen/OpenCL/builtins-amdgcn-vi.cl
+++ b/clang/test/CIR/CodeGen/OpenCL/builtins-amdgcn-vi.cl
@@ -76,3 +76,14 @@ void test_rcp_f16(global half* out, half a)
 {
   *out = __builtin_amdgcn_rcph(a);
 }
+
+// CIR-LABEL: @test_sqrt_f16
+// CIR: cir.llvm.intrinsic "amdgcn.sqrt" {{.*}} : (!cir.f16) -> !cir.f16
+// LLVM: define{{.*}} void @test_sqrt_f16
+// LLVM: call{{.*}} half @llvm.{{((amdgcn.){0,1})}}sqrt.f16(half %{{.*}})
+// OGCG: define{{.*}} void @test_sqrt_f16
+// OGCG: call{{.*}} half @llvm.{{((amdgcn.){0,1})}}sqrt.f16(half %{{.*}})
+void test_sqrt_f16(global half* out, half a)
+{
+  *out = __builtin_amdgcn_sqrth(a);
+}
diff --git a/clang/test/CIR/CodeGen/OpenCL/builtins_amdgcn.cl b/clang/test/CIR/CodeGen/OpenCL/builtins_amdgcn.cl
index 4fb408d92bf6..feda227849c5 100644
--- a/clang/test/CIR/CodeGen/OpenCL/builtins_amdgcn.cl
+++ b/clang/test/CIR/CodeGen/OpenCL/builtins_amdgcn.cl
@@ -380,3 +380,23 @@ void test_rcp_f32(global float* out, float a) {
 void test_rcp_f64(global double* out, double a) {
   *out = __builtin_amdgcn_rcp(a);
 }
+
+// CIR-LABEL: @test_sqrt_f32
+// CIR: cir.llvm.intrinsic "amdgcn.sqrt" {{.*}} : (!cir.float) -> !cir.float
+// LLVM: define{{.*}} void @test_sqrt_f32
+// LLVM: call{{.*}} float @llvm.amdgcn.sqrt.f32(float %{{.*}})
+// OGCG: define{{.*}} void @test_sqrt_f32
+// OGCG: call{{.*}} float @llvm.amdgcn.sqrt.f32(float %{{.*}})
+void test_sqrt_f32(global float* out, float a) {
+  *out = __builtin_amdgcn_sqrtf(a);
+}
+
+// CIR-LABEL: @test_sqrt_f64
+// CIR: cir.llvm.intrinsic "amdgcn.sqrt" {{.*}} : (!cir.double) -> !cir.double
+// LLVM: define{{.*}} void @test_sqrt_f64
+// LLVM: call{{.*}} double @llvm.amdgcn.sqrt.f64(double %{{.*}})
+// OGCG: define{{.*}} void @test_sqrt_f64
+// OGCG: call{{.*}} double @llvm.amdgcn.sqrt.f64(double %{{.*}})
+void test_sqrt_f64(global double* out, double a) {
+  *out = __builtin_amdgcn_sqrt(a);
+}