diff --git a/src/dawn/native/Pipeline.cpp b/src/dawn/native/Pipeline.cpp
index 51adb45..b7b855b 100644
--- a/src/dawn/native/Pipeline.cpp
+++ b/src/dawn/native/Pipeline.cpp
@@ -94,6 +94,11 @@
         DAWN_TRY(ValidateCompatibilityWithPipelineLayout(device, metadata, layout));
     }
 
+    DAWN_INVALID_IF(device->IsCompatibilityMode() && metadata.usesTextureLoadWithDepthTexture,
+                    "textureLoad can not be used with depth textures in compatibility mode in "
+                    "stage (%s), entry point \"%s\"",
+                    metadata.stage, entryPoint.name);
+
     // Validate if overridable constants exist in shader module
     // pipelineBase is not yet constructed at this moment so iterate constants from descriptor
     size_t numUninitializedConstants = metadata.uninitializedOverrides.size();
diff --git a/src/dawn/native/ShaderModule.cpp b/src/dawn/native/ShaderModule.cpp
index f779fd8..7f96311 100644
--- a/src/dawn/native/ShaderModule.cpp
+++ b/src/dawn/native/ShaderModule.cpp
@@ -670,6 +670,8 @@
         metadata->usesNumWorkgroups = entryPoint.num_workgroups_used;
     }
 
+    metadata->usesTextureLoadWithDepthTexture = entryPoint.has_texture_load_with_depth_texture;
+
     const CombinedLimits& limits = device->GetLimits();
     const uint32_t maxVertexAttributes = limits.v1.maxVertexAttributes;
     const uint32_t maxInterStageShaderVariables = limits.v1.maxInterStageShaderVariables;
diff --git a/src/dawn/native/ShaderModule.h b/src/dawn/native/ShaderModule.h
index f62158f..8c4e797 100644
--- a/src/dawn/native/ShaderModule.h
+++ b/src/dawn/native/ShaderModule.h
@@ -279,6 +279,7 @@
     bool usesSampleMaskOutput = false;
     bool usesSampleIndex = false;
     bool usesVertexIndex = false;
+    bool usesTextureLoadWithDepthTexture = false;
 };
 
 class ShaderModuleBase : public RefCountedWithExternalCount<ApiObjectBase>,
diff --git a/src/dawn/tests/unittests/validation/CompatValidationTests.cpp b/src/dawn/tests/unittests/validation/CompatValidationTests.cpp
index ac6ed9e..f33574a 100644
--- a/src/dawn/tests/unittests/validation/CompatValidationTests.cpp
+++ b/src/dawn/tests/unittests/validation/CompatValidationTests.cpp
@@ -213,7 +213,48 @@
     ASSERT_DEVICE_ERROR(device.CreateRenderPipeline(&testDescriptor));
 }
 
-TEST_F(CompatValidationTest, CanNotUseFragmentShaderWithSampleMask) {
+TEST_F(CompatValidationTest, CanNotCreatePipelineWithTextureLoadOfDepthTexture) {
+    wgpu::ShaderModule module = utils::CreateShaderModule(device, R"(
+        @group(0) @binding(0) var<storage, read_write> dstBuf : array<vec4f>;
+        @group(0) @binding(1) var tex1 : texture_2d<f32>;
+        @group(0) @binding(2) var tex2 : texture_depth_2d;
+        @group(0) @binding(3) var tex3 : texture_depth_2d_array;
+        @group(0) @binding(4) var tex4 : texture_depth_multisampled_2d;
+
+        @compute @workgroup_size(1) fn main1() {
+            dstBuf[0] = textureLoad(tex1, vec2(0), 0);
+        }
+
+        @compute @workgroup_size(1) fn main2() {
+            dstBuf[0] = vec4f(textureLoad(tex2, vec2(0), 0));
+        }
+
+        @compute @workgroup_size(1) fn main3() {
+            dstBuf[0] = vec4f(textureLoad(tex3, vec2(0), 0, 0));
+        }
+
+        @compute @workgroup_size(1) fn main4() {
+            dstBuf[4] = vec4f(textureLoad(tex4, vec2(0), 0));
+        }
+    )");
+
+    const char* entryPoints[] = {"main1", "main2", "main3", "main4"};
+    for (auto entryPoint : entryPoints) {
+        wgpu::ComputePipelineDescriptor pDesc;
+        pDesc.compute.module = module;
+        pDesc.compute.entryPoint = entryPoint;
+        if (entryPoint == entryPoints[0]) {
+            device.CreateComputePipeline(&pDesc);
+        } else {
+            ASSERT_DEVICE_ERROR(
+                device.CreateComputePipeline(&pDesc),
+                testing::HasSubstr(
+                    "textureLoad can not be used with depth textures in compatibility mode"));
+        }
+    }
+}
+
+TEST_F(CompatValidationTest, CanNotUseSampleMask) {
     wgpu::ShaderModule moduleSampleMaskOutput = utils::CreateShaderModule(device, R"(
         @vertex fn vs() -> @builtin(position) vec4f {
             return vec4f(1);
diff --git a/src/tint/lang/wgsl/inspector/entry_point.h b/src/tint/lang/wgsl/inspector/entry_point.h
index f64d55c..4f7a605 100644
--- a/src/tint/lang/wgsl/inspector/entry_point.h
+++ b/src/tint/lang/wgsl/inspector/entry_point.h
@@ -207,6 +207,8 @@
     bool vertex_index_used = false;
     /// Does the entry point use the instance_index builtin
     bool instance_index_used = false;
+    /// Does the entry point have a textureLoad call with a texture_depth??? texture
+    bool has_texture_load_with_depth_texture = false;
     /// The array length of the clip_distances builtin. Holding no value means the clip_distances
     /// is not used.
     std::optional<uint32_t> clip_distances_size;
diff --git a/src/tint/lang/wgsl/inspector/inspector.cc b/src/tint/lang/wgsl/inspector/inspector.cc
index 83d045a..79e01a8 100644
--- a/src/tint/lang/wgsl/inspector/inspector.cc
+++ b/src/tint/lang/wgsl/inspector/inspector.cc
@@ -27,6 +27,7 @@
 
 #include "src/tint/lang/wgsl/inspector/inspector.h"
 
+#include <algorithm>
 #include <unordered_set>
 #include <utility>
 
@@ -246,6 +247,23 @@
         }
     }
 
+    {
+        auto filter = [](const tint::sem::Call* call,
+                         tint::wgsl::BuiltinFn builtin_fn) -> std::optional<TextureUsageType> {
+            if (builtin_fn == wgsl::BuiltinFn::kTextureLoad) {
+                if (call->Arguments()[0]
+                        ->Type()
+                        ->IsAnyOf<core::type::DepthTexture,
+                                  core::type::DepthMultisampledTexture>()) {
+                    return TextureUsageType::kTextureLoad;
+                }
+            }
+            return {};
+        };
+        entry_point.has_texture_load_with_depth_texture =
+            !GetTextureUsagesForEntryPoint(*func, filter).empty();
+    }
+
     return entry_point;
 }
 
@@ -1081,32 +1099,87 @@
 }
 
 std::vector<Inspector::LevelSampleInfo> Inspector::GetTextureQueries(const std::string& ep_name) {
-    std::vector<LevelSampleInfo> res;
-
-    std::unordered_set<BindingPoint> seen = {};
-
-    Hashmap<const sem::Function*, Hashmap<const ast::Parameter*, TextureQueryType, 4>, 8>
-        fn_to_data;
-
-    auto record_function_param = [&fn_to_data](const sem::Function* func,
-                                               const ast::Parameter* param, TextureQueryType type) {
-        fn_to_data.GetOrAddZero(func).Add(param, type);
-    };
-
-    auto save_if_needed = [&res, &seen](const sem::GlobalVariable* global, TextureQueryType type) {
-        auto binding = global->Attributes().binding_point.value();
-        if (seen.insert(binding).second) {
-            res.emplace_back(LevelSampleInfo{type, binding.group, binding.binding});
-        }
-    };
-
-    auto& sem = program_.Sem();
-
     const auto* ep = FindEntryPointByName(ep_name);
     if (!ep) {
         return {};
     }
 
+    auto filter = [&](const tint::sem::Call* call,
+                      tint::wgsl::BuiltinFn builtin_fn) -> std::optional<TextureUsageType> {
+        switch (builtin_fn) {
+            case wgsl::BuiltinFn::kTextureNumLevels: {
+                return TextureUsageType::kTextureNumLevels;
+            }
+            case wgsl::BuiltinFn::kTextureDimensions: {
+                if (call->Declaration()->args.Length() <= 1) {
+                    // When textureDimension only takes a texture as the input,
+                    // it doesn't require calls to textureNumLevels to clamp mip levels.
+                    return {};
+                }
+                return TextureUsageType::kTextureNumLevels;
+            }
+            case wgsl::BuiltinFn::kTextureLoad: {
+                if (call->Arguments()[0]
+                        ->Type()
+                        ->IsAnyOf<core::type::MultisampledTexture,
+                                  core::type::DepthMultisampledTexture>()) {
+                    // When textureLoad takes a multisampled texture as the input,
+                    // it doesn't require to query the mip level.
+                    return {};
+                }
+                return TextureUsageType::kTextureNumLevels;
+            }
+            case wgsl::BuiltinFn::kTextureNumSamples: {
+                return TextureUsageType::kTextureNumSamples;
+            }
+            default:
+                return {};
+        }
+    };
+
+    auto usages = GetTextureUsagesForEntryPoint(*ep, filter);
+
+    auto t = [](const TextureUsageInfo& info) -> LevelSampleInfo {
+        return {
+            info.type == TextureUsageType::kTextureNumSamples ? TextureQueryType::kTextureNumSamples
+                                                              : TextureQueryType::kTextureNumLevels,
+            info.group,
+            info.binding,
+        };
+    };
+
+    std::vector<LevelSampleInfo> res;
+    std::transform(usages.begin(), usages.end(), std::back_inserter(res), t);
+    return res;
+}
+
+std::vector<Inspector::TextureUsageInfo> Inspector::GetTextureUsagesForEntryPoint(
+    const tint::ast::Function& ep,
+    std::function<std::optional<TextureUsageType>(const tint::sem::Call* call,
+                                                  tint::wgsl::BuiltinFn builtin_fn)> filter) {
+    TINT_ASSERT(ep.IsEntryPoint());
+
+    std::vector<TextureUsageInfo> res;
+
+    std::unordered_set<BindingPoint> seen = {};
+
+    Hashmap<const sem::Function*, Hashmap<const ast::Parameter*, TextureUsageType, 4>, 8>
+        fn_to_data;
+
+    auto record_function_param = [&fn_to_data](const sem::Function* func,
+                                               const ast::Parameter* param, TextureUsageType type) {
+        fn_to_data.GetOrAddZero(func).Add(param, type);
+    };
+
+    auto save_if_needed = [&res, &seen](const sem::GlobalVariable* global, TextureUsageType type) {
+        auto binding = global->Attributes().binding_point.value();
+        if (seen.insert(binding).second) {
+            res.emplace_back(TextureUsageInfo{type, binding.group, binding.binding});
+        }
+    };
+
+    auto& sem = program_.Sem();
+
     // This works in dependency order such that we'll see the texture call first and can record
     // any function parameter information and then as we walk up the function chain we can look
     // the call data.
@@ -1118,16 +1191,32 @@
 
         // This is an entrypoint, make sure it's the requested entry point
         if (fn->Declaration()->IsEntryPoint()) {
-            if (fn->Declaration() != ep) {
+            if (fn->Declaration() != &ep) {
                 continue;
             }
         } else {
             // Not an entry point, make sure it was called from the requested entry point
-            if (!fn->HasAncestorEntryPoint(ep->name->symbol)) {
+            if (!fn->HasAncestorEntryPoint(ep.name->symbol)) {
                 continue;
             }
         }
 
+        auto queryTextureBuiltin = [&](TextureUsageType type, const sem::Call* builtin_call,
+                                       const sem::Variable* texture_sem = nullptr) {
+            TINT_ASSERT(builtin_call);
+            if (!texture_sem) {
+                auto* texture_expr = builtin_call->Declaration()->args[0];
+                texture_sem = sem.GetVal(texture_expr)->RootIdentifier();
+            }
+            tint::Switch(
+                texture_sem,  //
+                [&](const sem::GlobalVariable* global) { save_if_needed(global, type); },
+                [&](const sem::Parameter* param) {
+                    record_function_param(fn, param->Declaration(), type);
+                },
+                TINT_ICE_ON_NO_MATCH);
+        };
+
         for (auto* call : fn->DirectCalls()) {
             // Builtin function call, record the texture information. If the used texture maps
             // back up to a function parameter just store the type of the call and we'll track the
@@ -1135,61 +1224,9 @@
             tint::Switch(
                 call->Target(),
                 [&](const sem::BuiltinFn* builtin) {
-                    auto queryTextureBuiltin = [&](TextureQueryType type,
-                                                   const sem::Call* builtin_call,
-                                                   const sem::Variable* texture_sem = nullptr) {
-                        TINT_ASSERT(builtin_call);
-                        if (!texture_sem) {
-                            auto* texture_expr = builtin_call->Declaration()->args[0];
-                            texture_sem = sem.GetVal(texture_expr)->RootIdentifier();
-                        }
-                        tint::Switch(
-                            texture_sem,  //
-                            [&](const sem::GlobalVariable* global) {
-                                save_if_needed(global, type);
-                            },
-                            [&](const sem::Parameter* param) {
-                                record_function_param(fn, param->Declaration(), type);
-                            },
-                            TINT_ICE_ON_NO_MATCH);
-                    };
-
-                    switch (builtin->Fn()) {
-                        case wgsl::BuiltinFn::kTextureNumLevels: {
-                            queryTextureBuiltin(TextureQueryType::kTextureNumLevels, call);
-                            break;
-                        }
-                        case wgsl::BuiltinFn::kTextureDimensions: {
-                            if (call->Declaration()->args.Length() <= 1) {
-                                // When textureDimension only takes a texture as the input,
-                                // it doesn't require calls to textureNumLevels to clamp mip levels.
-                                return;
-                            }
-                            queryTextureBuiltin(TextureQueryType::kTextureNumLevels, call);
-                            break;
-                        }
-                        case wgsl::BuiltinFn::kTextureLoad: {
-                            auto* texture_expr = call->Declaration()->args[0];
-                            auto* texture_sem = sem.GetVal(texture_expr)->RootIdentifier();
-                            TINT_ASSERT(texture_sem);
-                            if (texture_sem->Type()
-                                    ->UnwrapRef()
-                                    ->IsAnyOf<core::type::MultisampledTexture,
-                                              core::type::DepthMultisampledTexture>()) {
-                                // When textureLoad takes a multisampled texture as the input,
-                                // it doesn't require to query the mip level.
-                                return;
-                            }
-                            queryTextureBuiltin(TextureQueryType::kTextureNumLevels, call,
-                                                texture_sem);
-                            break;
-                        }
-                        case wgsl::BuiltinFn::kTextureNumSamples: {
-                            queryTextureBuiltin(TextureQueryType::kTextureNumSamples, call);
-                            break;
-                        }
-                        default:
-                            return;
+                    auto type = filter(call, builtin->Fn());
+                    if (type) {
+                        queryTextureBuiltin(*type, call);
                     }
                 },
                 [&](const sem::Function* func) {
diff --git a/src/tint/lang/wgsl/inspector/inspector.h b/src/tint/lang/wgsl/inspector/inspector.h
index a55e1a7..5b90b17 100644
--- a/src/tint/lang/wgsl/inspector/inspector.h
+++ b/src/tint/lang/wgsl/inspector/inspector.h
@@ -43,6 +43,7 @@
 #include "src/tint/lang/wgsl/inspector/resource_binding.h"
 #include "src/tint/lang/wgsl/inspector/scalar.h"
 #include "src/tint/lang/wgsl/program/program.h"
+#include "src/tint/lang/wgsl/sem/call.h"
 #include "src/tint/lang/wgsl/sem/sampler_texture_pair.h"
 #include "src/tint/utils/containers/unique_vector.h"
 
@@ -180,7 +181,7 @@
         uint32_t binding = 0;
     };
 
-    /// @param ep the entry point ot get the information for
+    /// @param ep the entry point to get the information for
     /// @returns a vector of information for textures which call textureNumLevels and
     /// textureNumSamples for backends which require additional support for those methods. Each
     /// binding point will only be returned once regardless of the number of calls made. The
@@ -299,6 +300,30 @@
     /// @param func the function of the entry point. Must be non-nullptr and true for IsEntryPoint()
     /// @returns the entry point information
     EntryPoint GetEntryPoint(const tint::ast::Function* func);
+
+    /// The information needed to be supplied.
+    enum class TextureUsageType : uint8_t {
+        /// textureLoad
+        kTextureLoad,
+        /// textureNumLevels
+        kTextureNumLevels,
+        /// textureNumSamples
+        kTextureNumSamples,
+    };
+    /// Information on level and sample calls by a given texture binding point
+    struct TextureUsageInfo {
+        /// The type of function
+        TextureUsageType type = TextureUsageType::kTextureNumLevels;
+        /// The group number
+        uint32_t group = 0;
+        /// The binding number
+        uint32_t binding = 0;
+    };
+
+    std::vector<Inspector::TextureUsageInfo> GetTextureUsagesForEntryPoint(
+        const tint::ast::Function& ep,
+        std::function<std::optional<TextureUsageType>(const tint::sem::Call* call,
+                                                      tint::wgsl::BuiltinFn builtin_fn)> filter);
 };
 
 }  // namespace tint::inspector
diff --git a/src/tint/lang/wgsl/resolver/compatibility_mode_test.cc b/src/tint/lang/wgsl/resolver/compatibility_mode_test.cc
index d3644e8..34ed0fe 100644
--- a/src/tint/lang/wgsl/resolver/compatibility_mode_test.cc
+++ b/src/tint/lang/wgsl/resolver/compatibility_mode_test.cc
@@ -246,79 +246,5 @@
         R"(12:34 error: flat interpolation must use 'either' sampling parameter in compatibility mode)");
 }
 
-class ResolverCompatibilityModeTest_TextureLoad : public ResolverCompatibilityModeTest {
-  protected:
-    void add_call_param(std::string name, ast::Type type, ExpressionList* call_params) {
-        const std::string type_name = type->identifier->symbol.Name();
-        if (tint::HasPrefix(type_name, "texture")) {
-            GlobalVar(name, type, Binding(0_a), Group(0_a));
-        } else {
-            GlobalVar(name, type, core::AddressSpace::kPrivate);
-        }
-        call_params->Push(Expr(Source{{12, 34}}, name));
-    }
-};
-
-TEST_F(ResolverCompatibilityModeTest_TextureLoad, TextureDepth2D) {
-    // textureLoad(someDepthTexture2D, coords, level)
-    const ast::Type coords_type = ty.vec2(ty.i32());
-    auto texture_type = ty.depth_texture(core::type::TextureDimension::k2d);
-
-    ExpressionList call_params;
-
-    add_call_param("texture", texture_type, &call_params);
-    add_call_param("coords", coords_type, &call_params);
-    add_call_param("level", ty.i32(), &call_params);
-
-    auto* expr = Call("textureLoad", call_params);
-    WrapInFunction(expr);
-
-    EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(
-        r()->error(),
-        R"(12:34 error: use of texture_depth_2d with textureLoad is not allowed in compatibility mode)");
-}
-
-TEST_F(ResolverCompatibilityModeTest_TextureLoad, TextureDepth2DArray) {
-    // textureLoad(someDepthTexture2DArray, coords, layer, level)
-    const ast::Type coords_type = ty.vec2(ty.i32());
-    auto texture_type = ty.depth_texture(core::type::TextureDimension::k2dArray);
-
-    ExpressionList call_params;
-
-    add_call_param("texture", texture_type, &call_params);
-    add_call_param("coords", coords_type, &call_params);
-    add_call_param("array_index", ty.i32(), &call_params);
-    add_call_param("level", ty.i32(), &call_params);
-
-    auto* expr = Call("textureLoad", call_params);
-    WrapInFunction(expr);
-
-    EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(
-        r()->error(),
-        R"(12:34 error: use of texture_depth_2d_array with textureLoad is not allowed in compatibility mode)");
-}
-
-TEST_F(ResolverCompatibilityModeTest_TextureLoad, TextureDepthMultisampled2D) {
-    // textureLoad(someDepthTextureMultisampled2D, coords, sample_index)
-    const ast::Type coords_type = ty.vec2(ty.i32());
-    auto texture_type = ty.depth_multisampled_texture(core::type::TextureDimension::k2d);
-
-    ExpressionList call_params;
-
-    add_call_param("texture", texture_type, &call_params);
-    add_call_param("coords", coords_type, &call_params);
-    add_call_param("sample_index", ty.i32(), &call_params);
-
-    auto* expr = Call("textureLoad", call_params);
-    WrapInFunction(expr);
-
-    EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(
-        r()->error(),
-        R"(12:34 error: use of texture_depth_multisampled_2d with textureLoad is not allowed in compatibility mode)");
-}
-
 }  // namespace
 }  // namespace tint::resolver
diff --git a/src/tint/lang/wgsl/resolver/validator.cc b/src/tint/lang/wgsl/resolver/validator.cc
index 92004bb..6ec74c9 100644
--- a/src/tint/lang/wgsl/resolver/validator.cc
+++ b/src/tint/lang/wgsl/resolver/validator.cc
@@ -35,7 +35,6 @@
 #include <utility>
 
 #include "src/tint/lang/core/fluent_types.h"
-#include "src/tint/lang/core/parameter_usage.h"
 #include "src/tint/lang/core/type/abstract_numeric.h"
 #include "src/tint/lang/core/type/atomic.h"
 #include "src/tint/lang/core/type/depth_multisampled_texture.h"
@@ -68,7 +67,6 @@
 #include "src/tint/lang/wgsl/ast/unary_op_expression.h"
 #include "src/tint/lang/wgsl/ast/variable_decl_statement.h"
 #include "src/tint/lang/wgsl/ast/workgroup_attribute.h"
-#include "src/tint/lang/wgsl/builtin_fn.h"
 #include "src/tint/lang/wgsl/sem/array.h"
 #include "src/tint/lang/wgsl/sem/break_if_statement.h"
 #include "src/tint/lang/wgsl/sem/call.h"
@@ -1912,19 +1910,6 @@
     std::string func_name = builtin->str();
     auto& signature = builtin->Signature();
 
-    if (mode_ == wgsl::ValidationMode::kCompat) {
-        if (builtin->Fn() == wgsl::BuiltinFn::kTextureLoad) {
-            auto* arg = call->Arguments()[0];
-            if (arg->Type()
-                    ->IsAnyOf<core::type::DepthTexture, core::type::DepthMultisampledTexture>()) {
-                AddError(arg->Declaration()->source)
-                    << "use of " << arg->Type()->FriendlyName()
-                    << " with textureLoad is not allowed in compatibility mode";
-                return false;
-            }
-        }
-    }
-
     auto check_arg_is_constexpr = [&](core::ParameterUsage usage, int min, int max) {
         auto signed_index = signature.IndexOf(usage);
         if (signed_index < 0) {