diff --git a/src/tint/lang/spirv/builtin_fn.cc b/src/tint/lang/spirv/builtin_fn.cc
index f70d75d..60cb6b9 100644
--- a/src/tint/lang/spirv/builtin_fn.cc
+++ b/src/tint/lang/spirv/builtin_fn.cc
@@ -168,6 +168,8 @@
             return "convertFToS";
         case BuiltinFn::kConvertSToF:
             return "convertSToF";
+        case BuiltinFn::kConvertUToF:
+            return "convertUToF";
         case BuiltinFn::kSdot:
             return "sdot";
         case BuiltinFn::kUdot:
@@ -260,6 +262,7 @@
         case BuiltinFn::kSMod:
         case BuiltinFn::kConvertFToS:
         case BuiltinFn::kConvertSToF:
+        case BuiltinFn::kConvertUToF:
             break;
     }
     return core::ir::Instruction::Accesses{};
diff --git a/src/tint/lang/spirv/builtin_fn.cc.tmpl b/src/tint/lang/spirv/builtin_fn.cc.tmpl
index 59c0565..8203aa9 100644
--- a/src/tint/lang/spirv/builtin_fn.cc.tmpl
+++ b/src/tint/lang/spirv/builtin_fn.cc.tmpl
@@ -105,6 +105,7 @@
         case BuiltinFn::kSMod:
         case BuiltinFn::kConvertFToS:
         case BuiltinFn::kConvertSToF:
+        case BuiltinFn::kConvertUToF:
             break;
     }
     return core::ir::Instruction::Accesses{};
diff --git a/src/tint/lang/spirv/builtin_fn.h b/src/tint/lang/spirv/builtin_fn.h
index 430e0a5..687eaa6 100644
--- a/src/tint/lang/spirv/builtin_fn.h
+++ b/src/tint/lang/spirv/builtin_fn.h
@@ -111,6 +111,7 @@
     kSMod,
     kConvertFToS,
     kConvertSToF,
+    kConvertUToF,
     kSdot,
     kUdot,
     kCooperativeMatrixLoad,
diff --git a/src/tint/lang/spirv/intrinsic/data.cc b/src/tint/lang/spirv/intrinsic/data.cc
index 39c200e..15f7fee 100644
--- a/src/tint/lang/spirv/intrinsic/data.cc
+++ b/src/tint/lang/spirv/intrinsic/data.cc
@@ -6711,30 +6711,37 @@
   },
   {
     /* [63] */
+    /* fn convertUToF<R : f32_f16>[T : iu32](T) -> R */
+    /* fn convertUToF<R : f32_f16>[T : iu32, N : num](vec<N, T>) -> vec<N, R> */
+    /* num overloads */ 2,
+    /* overloads */ OverloadIndex(180),
+  },
+  {
+    /* [64] */
     /* fn sdot(u32, u32, u32) -> i32 */
     /* num overloads */ 1,
     /* overloads */ OverloadIndex(193),
   },
   {
-    /* [64] */
+    /* [65] */
     /* fn udot(u32, u32, u32) -> u32 */
     /* num overloads */ 1,
     /* overloads */ OverloadIndex(194),
   },
   {
-    /* [65] */
+    /* [66] */
     /* fn cooperative_matrix_load<T : subgroup_matrix<K, S, C, R>>[K : subgroup_matrix_kind, S : fiu32_f16, C : num, R : num](ptr<workgroup_or_storage, S, readable>, u32, u32, u32) -> T */
     /* num overloads */ 1,
     /* overloads */ OverloadIndex(195),
   },
   {
-    /* [66] */
+    /* [67] */
     /* fn cooperative_matrix_store[K : subgroup_matrix_kind, S : fiu32_f16, C : num, R : num](ptr<workgroup_or_storage, S, writable>, subgroup_matrix<K, S, C, R>, u32, u32, u32) */
     /* num overloads */ 1,
     /* overloads */ OverloadIndex(196),
   },
   {
-    /* [67] */
+    /* [68] */
     /* fn cooperative_matrix_mul_add[T : subgroup_matrix_elements, TR : subgroup_matrix_elements, C : num, R : num, K : num](subgroup_matrix<subgroup_matrix_kind_left, T, K, R>, subgroup_matrix<subgroup_matrix_kind_right, T, C, K>, subgroup_matrix<subgroup_matrix_kind_result, TR, C, R>, u32) -> subgroup_matrix<subgroup_matrix_kind_result, TR, C, R> */
     /* num overloads */ 1,
     /* overloads */ OverloadIndex(197),
diff --git a/src/tint/lang/spirv/reader/lower/builtins.cc b/src/tint/lang/spirv/reader/lower/builtins.cc
index 5229a87..fc2b6cc 100644
--- a/src/tint/lang/spirv/reader/lower/builtins.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins.cc
@@ -152,6 +152,9 @@
                 case spirv::BuiltinFn::kConvertSToF:
                     ConvertSToF(builtin);
                     break;
+                case spirv::BuiltinFn::kConvertUToF:
+                    ConvertUToF(builtin);
+                    break;
                 default:
                     TINT_UNREACHABLE() << "unknown spirv builtin: " << builtin->Func();
             }
@@ -172,6 +175,20 @@
         call->Destroy();
     }
 
+    void ConvertUToF(spirv::ir::BuiltinCall* call) {
+        b.InsertBefore(call, [&] {
+            auto* result_ty = call->Result(0)->Type();
+
+            auto* arg = call->Args()[0];
+            if (arg->Type()->IsSignedIntegerScalarOrVector()) {
+                arg = b.Bitcast(ty.MatchWidth(ty.u32(), result_ty), arg)->Result(0);
+            }
+
+            b.ConvertWithResult(call->DetachResult(), arg);
+        });
+        call->Destroy();
+    }
+
     void ConvertFToS(spirv::ir::BuiltinCall* call) {
         b.InsertBefore(call, [&] {
             auto* res_ty = call->Result(0)->Type();
diff --git a/src/tint/lang/spirv/reader/lower/builtins_test.cc b/src/tint/lang/spirv/reader/lower/builtins_test.cc
index 7d9de47..71ba45e 100644
--- a/src/tint/lang/spirv/reader/lower/builtins_test.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins_test.cc
@@ -5341,5 +5341,131 @@
     EXPECT_EQ(expect, str());
 }
 
+TEST_F(SpirvParser_BuiltinsTest, ConvertUToF_ScalarSigned) {
+    auto* ep = b.ComputeFunction("foo");
+
+    b.Append(ep->Block(), [&] {  //
+        b.CallExplicit<spirv::ir::BuiltinCall>(ty.f32(), spirv::BuiltinFn::kConvertUToF,
+                                               Vector{ty.f32()}, 10_i);
+        b.Return(ep);
+    });
+
+    auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:f32 = spirv.convertUToF<f32> 10i
+    ret
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+    Run(Builtins);
+
+    auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:u32 = bitcast 10i
+    %3:f32 = convert %2
+    ret
+  }
+}
+)";
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, ConvertUToF_ScalarUnsigned) {
+    auto* ep = b.ComputeFunction("foo");
+
+    b.Append(ep->Block(), [&] {  //
+        b.CallExplicit<spirv::ir::BuiltinCall>(ty.f32(), spirv::BuiltinFn::kConvertUToF,
+                                               Vector{ty.f32()}, 10_u);
+        b.Return(ep);
+    });
+
+    auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:f32 = spirv.convertUToF<f32> 10u
+    ret
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+    Run(Builtins);
+
+    auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:f32 = convert 10u
+    ret
+  }
+}
+)";
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, ConvertUToF_VectorSigned) {
+    auto* ep = b.ComputeFunction("foo");
+
+    b.Append(ep->Block(), [&] {  //
+        b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<f32>(), spirv::BuiltinFn::kConvertUToF,
+                                               Vector{ty.f32()}, b.Splat<vec2<i32>>(10_i));
+        b.Return(ep);
+    });
+
+    auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:vec2<f32> = spirv.convertUToF<f32> vec2<i32>(10i)
+    ret
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+    Run(Builtins);
+
+    auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:vec2<u32> = bitcast vec2<i32>(10i)
+    %3:vec2<f32> = convert %2
+    ret
+  }
+}
+)";
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, ConvertUToF_VectorUnsigned) {
+    auto* ep = b.ComputeFunction("foo");
+
+    b.Append(ep->Block(), [&] {  //
+        b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<f32>(), spirv::BuiltinFn::kConvertUToF,
+                                               Vector{ty.f32()}, b.Splat<vec2<u32>>(10_u));
+        b.Return(ep);
+    });
+
+    auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:vec2<f32> = spirv.convertUToF<f32> vec2<u32>(10u)
+    ret
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+    Run(Builtins);
+
+    auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:vec2<f32> = convert vec2<u32>(10u)
+    ret
+  }
+}
+)";
+    EXPECT_EQ(expect, str());
+}
+
 }  // namespace
 }  // namespace tint::spirv::reader::lower
diff --git a/src/tint/lang/spirv/reader/parser/convert_test.cc b/src/tint/lang/spirv/reader/parser/convert_test.cc
index ac9bcd5..c4bcacb 100644
--- a/src/tint/lang/spirv/reader/parser/convert_test.cc
+++ b/src/tint/lang/spirv/reader/parser/convert_test.cc
@@ -325,5 +325,123 @@
 )");
 }
 
+TEST_F(SpirvParserTest, ConvertUToF_ScalarSigned) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+        %int = OpTypeInt 32 1
+      %float = OpTypeFloat 32
+        %two = OpConstant %int 2
+    %void_fn = OpTypeFunction %void
+
+       %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+               %1 = OpConvertUToF %float %two
+               OpReturn
+               OpFunctionEnd
+)",
+              R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:f32 = spirv.convertUToF<f32> 2i
+    ret
+  }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, ConvertUToF_ScalarUnsigned) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+       %uint = OpTypeInt 32 0
+      %float = OpTypeFloat 32
+        %two = OpConstant %uint 2
+    %void_fn = OpTypeFunction %void
+
+       %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+               %1 = OpConvertUToF %float %two
+               OpReturn
+               OpFunctionEnd
+)",
+              R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:f32 = spirv.convertUToF<f32> 2u
+    ret
+  }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, ConvertUToF_VectorSigned) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+      %float = OpTypeFloat 32
+    %v2float = OpTypeVector %float 2
+        %two = OpConstant %int 2
+     %v2_two = OpConstantComposite %v2int %two %two
+    %void_fn = OpTypeFunction %void
+
+       %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+               %1 = OpConvertUToF %v2float %v2_two
+               OpReturn
+               OpFunctionEnd
+)",
+              R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:vec2<f32> = spirv.convertUToF<f32> vec2<i32>(2i)
+    ret
+  }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, ConvertUToF_VectorUnsigned) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+       %uint = OpTypeInt 32 0
+     %v2uint = OpTypeVector %uint 2
+      %float = OpTypeFloat 32
+    %v2float = OpTypeVector %float 2
+        %two = OpConstant %uint 2
+     %v2_two = OpConstantComposite %v2uint %two %two
+    %void_fn = OpTypeFunction %void
+
+       %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+               %1 = OpConvertUToF %v2float %v2_two
+               OpReturn
+               OpFunctionEnd
+)",
+              R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:vec2<f32> = spirv.convertUToF<f32> vec2<u32>(2u)
+    ret
+  }
+}
+)");
+}
+
 }  // namespace
 }  // namespace tint::spirv::reader
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index ed0c379..736ca0e 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -665,6 +665,9 @@
                 case spv::Op::OpConvertSToF:
                     EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kConvertSToF);
                     break;
+                case spv::Op::OpConvertUToF:
+                    EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kConvertUToF);
+                    break;
                 case spv::Op::OpAccessChain:
                 case spv::Op::OpInBoundsAccessChain:
                     EmitAccess(inst);
diff --git a/src/tint/lang/spirv/spirv.def b/src/tint/lang/spirv/spirv.def
index c74b5ce..3cf4ad5 100644
--- a/src/tint/lang/spirv/spirv.def
+++ b/src/tint/lang/spirv/spirv.def
@@ -418,6 +418,8 @@
 implicit(T: f32_f16, N: num) fn convertFToS<R: iu32>(vec<N, T>) -> vec<N, R>
 implicit(T: iu32) fn convertSToF<R: f32_f16>(T) -> R
 implicit(T: iu32, N: num) fn convertSToF<R: f32_f16>(vec<N, T>) -> vec<N, R>
+implicit(T: iu32) fn convertUToF<R: f32_f16>(T) -> R
+implicit(T: iu32, N: num) fn convertUToF<R: f32_f16>(vec<N, T>) -> vec<N, R>
 
 ////////////////////////////////////////////////////////////////////////////////
 // SPV_KHR_integer_dot_product instructions
diff --git a/src/tint/lang/spirv/writer/printer/printer.cc b/src/tint/lang/spirv/writer/printer/printer.cc
index 3d9fff1..eaed98a 100644
--- a/src/tint/lang/spirv/writer/printer/printer.cc
+++ b/src/tint/lang/spirv/writer/printer/printer.cc
@@ -1530,6 +1530,9 @@
             case BuiltinFn::kConvertSToF:
                 op = spv::Op::OpConvertSToF;
                 break;
+            case BuiltinFn::kConvertUToF:
+                op = spv::Op::OpConvertUToF;
+                break;
             case spirv::BuiltinFn::kNone:
                 TINT_ICE() << "undefined spirv ir function";
         }