diff --git a/rpcs3/Emu/Cell/lv2/sys_rsx.cpp b/rpcs3/Emu/Cell/lv2/sys_rsx.cpp
index 9b09d57af9..05ace4818a 100644
--- a/rpcs3/Emu/Cell/lv2/sys_rsx.cpp
+++ b/rpcs3/Emu/Cell/lv2/sys_rsx.cpp
@@ -343,8 +343,12 @@ s32 sys_rsx_context_attribute(s32 context_id, u32 package_id, u64 a3, u64 a4, u6
 		zcull.height = (((a4 & 0xFFFFFFFF) >> 6) & 0xFF) << 6;
 		zcull.cullStart = (a5 >> 32);
 		zcull.offset = (a5 & 0xFFFFFFFF);
+		zcull.zcullDir = ((a6 >> 32) >> 1) & 0x1;
+		zcull.zcullFormat = ((a6 >> 32) >> 2) & 0x3FF;
+		zcull.sFunc = ((a6 >> 32) >> 12) & 0xF;
+		zcull.sRef = ((a6 >> 32) >> 16) & 0xFF;
+		zcull.sMask = ((a6 >> 32) >> 24) & 0xFF;
 		zcull.binded = (a6 & 0xFFFFFFFF) != 0;
-		//TODO: Set zculldir, format, sfunc, sref, smask
 	}
 	break;
 
diff --git a/rpcs3/Emu/RSX/GL/GLGSRender.cpp b/rpcs3/Emu/RSX/GL/GLGSRender.cpp
index 34741ce43b..906b06c05a 100644
--- a/rpcs3/Emu/RSX/GL/GLGSRender.cpp
+++ b/rpcs3/Emu/RSX/GL/GLGSRender.cpp
@@ -857,6 +857,14 @@ bool GLGSRender::do_method(u32 cmd, u32 arg)
 
 		return true;
 	}
+	case NV4097_CLEAR_ZCULL_SURFACE:
+	{
+		//TODO
+		init_buffers(true);
+		clear_surface(0x3);
+
+		return true;
+	}
 	case NV4097_TEXTURE_READ_SEMAPHORE_RELEASE:
 	case NV4097_BACK_END_WRITE_SEMAPHORE_RELEASE:
 		flush_draw_buffers = true;
@@ -1283,6 +1291,7 @@ void GLGSRender::clear_zcull_stats(u32 type)
 
 			//re-enable cull stats if stats are enabled
 			check_zcull_status(false, false);
+			zcull_task_queue.active_query->num_draws = 0;
 		}
 
 		current_zcull_stats.clear();
@@ -1296,10 +1305,12 @@ u32 GLGSRender::get_zcull_stats(u32 type)
 
 	if (zcull_task_queue.active_query &&
 		zcull_task_queue.active_query->active &&
-		current_zcull_stats.zpass_pixel_cnt == 0)
+		current_zcull_stats.zpass_pixel_cnt == 0 &&
+		type == CELL_GCM_ZPASS_PIXEL_CNT)
 	{
 		//The zcull unit is still bound as the read is happening and there are no results ready
-		check_zcull_status(false, true);
+		check_zcull_status(false, true);  //close current query
+		check_zcull_status(false, false); //start new query since stat counting is still active
 	}
 
 	switch (type)
@@ -1307,17 +1318,16 @@ u32 GLGSRender::get_zcull_stats(u32 type)
 	case CELL_GCM_ZPASS_PIXEL_CNT:
 	{
 		if (current_zcull_stats.zpass_pixel_cnt > 0)
-			return UINT32_MAX;
+			return UINT16_MAX;
 
-		//If we have no results, we might as well synchronize here and wait for results to become available
 		synchronize_zcull_stats(true);
-		return (current_zcull_stats.zpass_pixel_cnt > 0)? UINT32_MAX: 0;
+		return (current_zcull_stats.zpass_pixel_cnt > 0)? UINT16_MAX : 0;
 	}
 	case CELL_GCM_ZCULL_STATS:
 	case CELL_GCM_ZCULL_STATS1:
 	case CELL_GCM_ZCULL_STATS2:
 		//TODO
-		return UINT32_MAX;
+		return UINT16_MAX;
 	case CELL_GCM_ZCULL_STATS3:
 	{
 		//Some kind of inverse value
@@ -1325,7 +1335,7 @@ u32 GLGSRender::get_zcull_stats(u32 type)
 			return 0;
 		
 		synchronize_zcull_stats(true);
-		return (current_zcull_stats.zpass_pixel_cnt > 0) ? 0 : UINT32_MAX;
+		return (current_zcull_stats.zpass_pixel_cnt > 0) ? 0 : UINT16_MAX;
 	}
 	default:
 		LOG_ERROR(RSX, "Unknown zcull stat type %d", type);
diff --git a/rpcs3/Emu/RSX/RSXThread.h b/rpcs3/Emu/RSX/RSXThread.h
index eb179ab0af..a50c049528 100644
--- a/rpcs3/Emu/RSX/RSXThread.h
+++ b/rpcs3/Emu/RSX/RSXThread.h
@@ -259,7 +259,7 @@ namespace rsx
 		//zcull
 		virtual void notify_zcull_info_changed() {}
 		virtual void clear_zcull_stats(u32 /*type*/) {}
-		virtual u32 get_zcull_stats(u32 /*type*/) { return UINT32_MAX; }
+		virtual u32 get_zcull_stats(u32 /*type*/) { return UINT16_MAX; }
 
 		gsl::span<const gsl::byte> get_raw_index_array(const std::vector<std::pair<u32, u32> >& draw_indexed_clause) const;
 		gsl::span<const gsl::byte> get_raw_vertex_buffer(const rsx::data_array_format_info&, u32 base_offset, const std::vector<std::pair<u32, u32>>& vertex_ranges) const;
diff --git a/rpcs3/Emu/RSX/rsx_methods.cpp b/rpcs3/Emu/RSX/rsx_methods.cpp
index 7f13bff959..bc7402e917 100644
--- a/rpcs3/Emu/RSX/rsx_methods.cpp
+++ b/rpcs3/Emu/RSX/rsx_methods.cpp
@@ -98,6 +98,16 @@ namespace rsx
 			}
 		}
 
+		void clear_zcull(thread* rsx, u32 _reg, u32 arg)
+		{
+			rsx->do_method(NV4097_CLEAR_ZCULL_SURFACE, arg);
+
+			if (rsx->capture_current_frame)
+			{
+				rsx->capture_frame("clear zcull memory");
+			}
+		}
+
 		void texture_read_semaphore_release(thread* rsx, u32 _reg, u32 arg)
 		{
 			const u32 index = method_registers.semaphore_offset_4097() >> 4;
@@ -1534,6 +1544,7 @@ namespace rsx
 		bind<NV4097_SET_ZCULL_EN, nv4097::set_zcull_render_enable>();
 		bind<NV4097_SET_ZCULL_STATS_ENABLE, nv4097::set_zcull_stats_enable>();
 		bind<NV4097_SET_ZPASS_PIXEL_COUNT_ENABLE, nv4097::set_zcull_pixel_count_enable>();
+		bind<NV4097_CLEAR_ZCULL_SURFACE, nv4097::clear_zcull>();
 
 		//NV308A
 		bind_range<NV308A_COLOR, 1, 256, nv308a::color>();