diff --git a/src/tint/writer/glsl/generator_impl.cc b/src/tint/writer/glsl/generator_impl.cc
index c4722f6beb..05fd7613b3 100644
--- a/src/tint/writer/glsl/generator_impl.cc
+++ b/src/tint/writer/glsl/generator_impl.cc
@@ -2488,7 +2488,8 @@ bool GeneratorImpl::EmitType(std::ostream& out,
       out << "out ";
       break;
     }
-    case ast::StorageClass::kUniform: {
+    case ast::StorageClass::kUniform:
+    case ast::StorageClass::kUniformConstant: {
       out << "uniform ";
       break;
     }
@@ -2552,7 +2553,7 @@ bool GeneratorImpl::EmitType(std::ostream& out,
     auto* depth_ms = tex->As<sem::DepthMultisampledTexture>();
     auto* sampled = tex->As<sem::SampledTexture>();
 
-    out << "uniform highp ";
+    out << "highp ";
 
     if (storage && storage->access() != ast::Access::kRead) {
       out << "writeonly ";
diff --git a/src/tint/writer/glsl/generator_impl_type_test.cc b/src/tint/writer/glsl/generator_impl_type_test.cc
index 1d29857fec..c6d7e86e74 100644
--- a/src/tint/writer/glsl/generator_impl_type_test.cc
+++ b/src/tint/writer/glsl/generator_impl_type_test.cc
@@ -472,7 +472,7 @@ TEST_F(GlslGeneratorImplTest_Type, EmitMultisampledTexture) {
   ASSERT_TRUE(gen.EmitType(out, s, ast::StorageClass::kNone,
                            ast::Access::kReadWrite, ""))
       << gen.error();
-  EXPECT_EQ(out.str(), "uniform highp sampler2DMS");
+  EXPECT_EQ(out.str(), "highp sampler2DMS");
 }
 
 struct GlslStorageTextureData {
diff --git a/test/tint/builtins/textureLoad/texture_param.wgsl b/test/tint/builtins/textureLoad/texture_param.wgsl
new file mode 100644
index 0000000000..81f359ea41
--- /dev/null
+++ b/test/tint/builtins/textureLoad/texture_param.wgsl
@@ -0,0 +1,25 @@
+@group(1) @binding(0) var arg_0: texture_2d<i32>;
+
+fn textureLoad2d(texture: texture_2d<i32>, coords: vec2<i32>, level: i32) -> vec4<i32> {
+  return textureLoad(texture, coords, level);
+}
+
+fn doTextureLoad() {
+  var res: vec4<i32> = textureLoad2d(arg_0, vec2<i32>(), 0);
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  doTextureLoad();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  doTextureLoad();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  doTextureLoad();
+}
diff --git a/test/tint/builtins/textureLoad/texture_param.wgsl.expected.glsl b/test/tint/builtins/textureLoad/texture_param.wgsl.expected.glsl
new file mode 100644
index 0000000000..862fcee33f
--- /dev/null
+++ b/test/tint/builtins/textureLoad/texture_param.wgsl.expected.glsl
@@ -0,0 +1,63 @@
+#version 310 es
+
+ivec4 textureLoad2d(highp isampler2D tint_symbol_1, ivec2 coords, int level) {
+  return texelFetch(tint_symbol_1, coords, level);
+}
+
+uniform highp isampler2D arg_0_1;
+void doTextureLoad() {
+  ivec4 res = textureLoad2d(arg_0_1, ivec2(0, 0), 0);
+}
+
+vec4 vertex_main() {
+  doTextureLoad();
+  return vec4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+void main() {
+  vec4 inner_result = vertex_main();
+  gl_Position = inner_result;
+  gl_Position.y = -(gl_Position.y);
+  gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+  return;
+}
+#version 310 es
+precision mediump float;
+
+ivec4 textureLoad2d(highp isampler2D tint_symbol_1, ivec2 coords, int level) {
+  return texelFetch(tint_symbol_1, coords, level);
+}
+
+uniform highp isampler2D arg_0_1;
+void doTextureLoad() {
+  ivec4 res = textureLoad2d(arg_0_1, ivec2(0, 0), 0);
+}
+
+void fragment_main() {
+  doTextureLoad();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+ivec4 textureLoad2d(highp isampler2D tint_symbol_1, ivec2 coords, int level) {
+  return texelFetch(tint_symbol_1, coords, level);
+}
+
+uniform highp isampler2D arg_0_1;
+void doTextureLoad() {
+  ivec4 res = textureLoad2d(arg_0_1, ivec2(0, 0), 0);
+}
+
+void compute_main() {
+  doTextureLoad();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}
diff --git a/test/tint/builtins/textureLoad/texture_param.wgsl.expected.hlsl b/test/tint/builtins/textureLoad/texture_param.wgsl.expected.hlsl
new file mode 100644
index 0000000000..1544f5ec7c
--- /dev/null
+++ b/test/tint/builtins/textureLoad/texture_param.wgsl.expected.hlsl
@@ -0,0 +1,36 @@
+Texture2D<int4> arg_0 : register(t0, space1);
+
+int4 textureLoad2d(Texture2D<int4> tint_symbol, int2 coords, int level) {
+  return tint_symbol.Load(int3(coords, level));
+}
+
+void doTextureLoad() {
+  int4 res = textureLoad2d(arg_0, int2(0, 0), 0);
+}
+
+struct tint_symbol_1 {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  doTextureLoad();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol_1 vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol_1 wrapper_result = (tint_symbol_1)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  doTextureLoad();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  doTextureLoad();
+  return;
+}
diff --git a/test/tint/builtins/textureLoad/texture_param.wgsl.expected.msl b/test/tint/builtins/textureLoad/texture_param.wgsl.expected.msl
new file mode 100644
index 0000000000..cee9f8ef51
--- /dev/null
+++ b/test/tint/builtins/textureLoad/texture_param.wgsl.expected.msl
@@ -0,0 +1,37 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int4 textureLoad2d(texture2d<int, access::sample> tint_symbol, int2 coords, int level) {
+  return tint_symbol.read(uint2(coords), level);
+}
+
+void doTextureLoad(texture2d<int, access::sample> tint_symbol_2) {
+  int4 res = textureLoad2d(tint_symbol_2, int2(), 0);
+}
+
+struct tint_symbol_1 {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner(texture2d<int, access::sample> tint_symbol_3) {
+  doTextureLoad(tint_symbol_3);
+  return float4();
+}
+
+vertex tint_symbol_1 vertex_main(texture2d<int, access::sample> tint_symbol_4 [[texture(0)]]) {
+  float4 const inner_result = vertex_main_inner(tint_symbol_4);
+  tint_symbol_1 wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main(texture2d<int, access::sample> tint_symbol_5 [[texture(0)]]) {
+  doTextureLoad(tint_symbol_5);
+  return;
+}
+
+kernel void compute_main(texture2d<int, access::sample> tint_symbol_6 [[texture(0)]]) {
+  doTextureLoad(tint_symbol_6);
+  return;
+}
+
diff --git a/test/tint/builtins/textureLoad/texture_param.wgsl.expected.spvasm b/test/tint/builtins/textureLoad/texture_param.wgsl.expected.spvasm
new file mode 100644
index 0000000000..b292e121d8
--- /dev/null
+++ b/test/tint/builtins/textureLoad/texture_param.wgsl.expected.spvasm
@@ -0,0 +1,90 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 47
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %arg_0 "arg_0"
+               OpName %textureLoad2d "textureLoad2d"
+               OpName %texture "texture"
+               OpName %coords "coords"
+               OpName %level "level"
+               OpName %doTextureLoad "doTextureLoad"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+               OpDecorate %arg_0 DescriptorSet 1
+               OpDecorate %arg_0 Binding 0
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+        %int = OpTypeInt 32 1
+         %11 = OpTypeImage %int 2D 0 0 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+      %arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+      %v4int = OpTypeVector %int 4
+      %v2int = OpTypeVector %int 2
+         %13 = OpTypeFunction %v4int %11 %v2int %int
+       %void = OpTypeVoid
+         %22 = OpTypeFunction %void
+         %28 = OpConstantNull %v2int
+      %int_0 = OpConstant %int 0
+%_ptr_Function_v4int = OpTypePointer Function %v4int
+         %32 = OpConstantNull %v4int
+         %33 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%textureLoad2d = OpFunction %v4int None %13
+    %texture = OpFunctionParameter %11
+     %coords = OpFunctionParameter %v2int
+      %level = OpFunctionParameter %int
+         %20 = OpLabel
+         %21 = OpImageFetch %v4int %texture %coords Lod %level
+               OpReturnValue %21
+               OpFunctionEnd
+%doTextureLoad = OpFunction %void None %22
+         %25 = OpLabel
+        %res = OpVariable %_ptr_Function_v4int Function %32
+         %27 = OpLoad %11 %arg_0
+         %26 = OpFunctionCall %v4int %textureLoad2d %27 %28 %int_0
+               OpStore %res %26
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %33
+         %35 = OpLabel
+         %36 = OpFunctionCall %void %doTextureLoad
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %22
+         %38 = OpLabel
+         %39 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %39
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %22
+         %42 = OpLabel
+         %43 = OpFunctionCall %void %doTextureLoad
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %22
+         %45 = OpLabel
+         %46 = OpFunctionCall %void %doTextureLoad
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/textureLoad/texture_param.wgsl.expected.wgsl b/test/tint/builtins/textureLoad/texture_param.wgsl.expected.wgsl
new file mode 100644
index 0000000000..13f6d82c09
--- /dev/null
+++ b/test/tint/builtins/textureLoad/texture_param.wgsl.expected.wgsl
@@ -0,0 +1,25 @@
+@group(1) @binding(0) var arg_0 : texture_2d<i32>;
+
+fn textureLoad2d(texture : texture_2d<i32>, coords : vec2<i32>, level : i32) -> vec4<i32> {
+  return textureLoad(texture, coords, level);
+}
+
+fn doTextureLoad() {
+  var res : vec4<i32> = textureLoad2d(arg_0, vec2<i32>(), 0);
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  doTextureLoad();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  doTextureLoad();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  doTextureLoad();
+}