diff --git a/reference/opt/shaders-msl/frag/ray-query-object-in-function.spv14.vk.msl24.frag b/reference/opt/shaders-msl/frag/ray-query-object-in-function.spv14.vk.msl24.frag
new file mode 100644
index 00000000..859ace2c
--- /dev/null
+++ b/reference/opt/shaders-msl/frag/ray-query-object-in-function.spv14.vk.msl24.frag
@@ -0,0 +1,45 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+#if __METAL_VERSION__ >= 230
+#include <metal_raytracing>
+using namespace metal::raytracing;
+#endif
+
+using namespace metal;
+
+struct main0_out
+{
+    float4 outColor [[color(0)]];
+};
+
+struct main0_in
+{
+    float4 inPos [[user(locn0)]];
+};
+
+fragment main0_out main0(main0_in in [[stage_in]], raytracing::acceleration_structure<raytracing::instancing> topLevelAS [[buffer(0)]])
+{
+    main0_out out = {};
+    raytracing::intersection_query<raytracing::instancing, raytracing::triangle_data> rayQuery;
+    rayQuery.reset(ray(float3((in.inPos.xy * 4.0) - float2(2.0), 1.0), float3(0.0, 0.0, -1.0), 0.001000000047497451305389404296875, 2.0), topLevelAS, intersection_params());
+    for (;;)
+    {
+        bool _88 = rayQuery.next();
+        if (_88)
+        {
+            continue;
+        }
+        else
+        {
+            break;
+        }
+    }
+    uint _92 = uint(rayQuery.get_committed_intersection_type());
+    if (_92 == 0u)
+    {
+        discard_fragment();
+    }
+    out.outColor = in.inPos;
+    return out;
+}
+
diff --git a/reference/shaders-msl/frag/ray-query-object-in-function.spv14.vk.msl24.frag b/reference/shaders-msl/frag/ray-query-object-in-function.spv14.vk.msl24.frag
new file mode 100644
index 00000000..3ab6a471
--- /dev/null
+++ b/reference/shaders-msl/frag/ray-query-object-in-function.spv14.vk.msl24.frag
@@ -0,0 +1,60 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#if __METAL_VERSION__ >= 230
+#include <metal_raytracing>
+using namespace metal::raytracing;
+#endif
+
+using namespace metal;
+
+struct main0_out
+{
+    float4 outColor [[color(0)]];
+};
+
+struct main0_in
+{
+    float4 inPos [[user(locn0)]];
+};
+
+static inline __attribute__((always_inline))
+uint doRay(thread const float3& rayOrigin, thread const float3& rayDirection, thread const float& rayDistance, thread raytracing::intersection_query<raytracing::instancing, raytracing::triangle_data>& rayQuery, thread const raytracing::acceleration_structure<raytracing::instancing>& topLevelAS)
+{
+    rayQuery.reset(ray(rayOrigin, rayDirection, 0.001000000047497451305389404296875, rayDistance), topLevelAS, intersection_params());
+    for (;;)
+    {
+        bool _36 = rayQuery.next();
+        if (_36)
+        {
+            continue;
+        }
+        else
+        {
+            break;
+        }
+    }
+    uint _40 = uint(rayQuery.get_committed_intersection_type());
+    return _40;
+}
+
+fragment main0_out main0(main0_in in [[stage_in]], raytracing::acceleration_structure<raytracing::instancing> topLevelAS [[buffer(0)]])
+{
+    main0_out out = {};
+    float3 rayOrigin = float3((in.inPos.xy * 4.0) - float2(2.0), 1.0);
+    float3 rayDirection = float3(0.0, 0.0, -1.0);
+    float rayDistance = 2.0;
+    float3 param = rayOrigin;
+    float3 param_1 = rayDirection;
+    float param_2 = rayDistance;
+    raytracing::intersection_query<raytracing::instancing, raytracing::triangle_data> rayQuery;
+    uint _70 = doRay(param, param_1, param_2, rayQuery, topLevelAS);
+    if (_70 == 0u)
+    {
+        discard_fragment();
+    }
+    out.outColor = in.inPos;
+    return out;
+}
+
diff --git a/shaders-msl/frag/ray-query-object-in-function.spv14.vk.msl24.frag b/shaders-msl/frag/ray-query-object-in-function.spv14.vk.msl24.frag
new file mode 100644
index 00000000..d2d94fcb
--- /dev/null
+++ b/shaders-msl/frag/ray-query-object-in-function.spv14.vk.msl24.frag
@@ -0,0 +1,31 @@
+#version 460
+
+#extension GL_ARB_separate_shader_objects : enable
+#extension GL_EXT_ray_query : enable
+
+layout(location = 0) in  vec4 inPos;
+layout(location = 0) out vec4 outColor;
+
+layout(binding = 0) uniform accelerationStructureEXT topLevelAS;
+
+uint doRay(vec3 rayOrigin, vec3 rayDirection, float rayDistance) {
+	rayQueryEXT rayQuery;
+	rayQueryInitializeEXT(rayQuery, topLevelAS, gl_RayFlagsTerminateOnFirstHitEXT, 0xFF,
+			rayOrigin, 0.001, rayDirection, rayDistance);
+
+	while(rayQueryProceedEXT(rayQuery))
+		;
+
+	return rayQueryGetIntersectionTypeEXT(rayQuery, true);
+}
+
+void main() {
+	vec3  rayOrigin    = vec3(inPos.xy*4.0-vec2(2.0),1.0);
+	vec3  rayDirection = vec3(0,0,-1);
+	float rayDistance  = 2.0;
+
+	if(doRay(rayOrigin,rayDirection,rayDistance) == gl_RayQueryCommittedIntersectionNoneEXT)
+		discard;
+
+	outColor = inPos;
+}
diff --git a/spirv_msl.cpp b/spirv_msl.cpp
index 04555f43..a9194fe1 100644
--- a/spirv_msl.cpp
+++ b/spirv_msl.cpp
@@ -1765,6 +1765,45 @@ void CompilerMSL::extract_global_variables_from_function(uint32_t func_id, std::
 				break;
 			}
 
+			case OpRayQueryInitializeKHR:
+			case OpRayQueryProceedKHR:
+			case OpRayQueryTerminateKHR:
+			case OpRayQueryGenerateIntersectionKHR:
+			case OpRayQueryConfirmIntersectionKHR:
+			{
+				// Ray query accesses memory directly, need check pass down object if using Private storage class.
+				uint32_t base_id = ops[0];
+				if (global_var_ids.find(base_id) != global_var_ids.end())
+					added_arg_ids.insert(base_id);
+				break;
+			}
+
+			case OpRayQueryGetRayTMinKHR:
+			case OpRayQueryGetRayFlagsKHR:
+			case OpRayQueryGetWorldRayOriginKHR:
+			case OpRayQueryGetWorldRayDirectionKHR:
+			case OpRayQueryGetIntersectionCandidateAABBOpaqueKHR:
+			case OpRayQueryGetIntersectionTypeKHR:
+			case OpRayQueryGetIntersectionTKHR:
+			case OpRayQueryGetIntersectionInstanceCustomIndexKHR:
+			case OpRayQueryGetIntersectionInstanceIdKHR:
+			case OpRayQueryGetIntersectionInstanceShaderBindingTableRecordOffsetKHR:
+			case OpRayQueryGetIntersectionGeometryIndexKHR:
+			case OpRayQueryGetIntersectionPrimitiveIndexKHR:
+			case OpRayQueryGetIntersectionBarycentricsKHR:
+			case OpRayQueryGetIntersectionFrontFaceKHR:
+			case OpRayQueryGetIntersectionObjectRayDirectionKHR:
+			case OpRayQueryGetIntersectionObjectRayOriginKHR:
+			case OpRayQueryGetIntersectionObjectToWorldKHR:
+			case OpRayQueryGetIntersectionWorldToObjectKHR:
+			{
+				// Ray query accesses memory directly, need check pass down object if using Private storage class.
+				uint32_t base_id = ops[2];
+				if (global_var_ids.find(base_id) != global_var_ids.end())
+					added_arg_ids.insert(base_id);
+				break;
+			}
+
 			default:
 				break;
 			}