summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--configs/tested-cfgs/SM2_GTX480/gpgpusim.config7
-rw-r--r--configs/tested-cfgs/SM6_TITANX/gpgpusim.config10
-rw-r--r--configs/tested-cfgs/SM7_TITANV/gpgpusim.config11
-rw-r--r--libcuda/cuda_runtime_api.cc68
-rw-r--r--src/abstract_hardware_model.cc14
-rw-r--r--src/abstract_hardware_model.h35
-rw-r--r--src/cuda-sim/cuda-sim.cc16
-rw-r--r--src/cuda-sim/cuda_device_runtime.cc3
-rw-r--r--src/cuda-sim/instructions.cc13
-rw-r--r--src/gpgpu-sim/gpu-sim.cc24
-rw-r--r--src/gpgpu-sim/gpu-sim.h16
11 files changed, 186 insertions, 31 deletions
diff --git a/configs/tested-cfgs/SM2_GTX480/gpgpusim.config b/configs/tested-cfgs/SM2_GTX480/gpgpusim.config
index d71b2fd..cf3627b 100644
--- a/configs/tested-cfgs/SM2_GTX480/gpgpusim.config
+++ b/configs/tested-cfgs/SM2_GTX480/gpgpusim.config
@@ -3,6 +3,13 @@
-gpgpu_ptx_sim_mode 0
-gpgpu_ptx_force_max_capability 20
+# Device Limits
+-gpgpu_stack_size_limit 1024
+-gpgpu_heap_size_limit 8388608
+
+# Compute Capability
+-gpgpu_compute_capability_major 2
+-gpgpu_compute_capability_minor 0
# SASS execution (only supported with CUDA >= 4.0)
-gpgpu_ptx_convert_to_ptxplus 0
diff --git a/configs/tested-cfgs/SM6_TITANX/gpgpusim.config b/configs/tested-cfgs/SM6_TITANX/gpgpusim.config
index cb23ab3..23d044c 100644
--- a/configs/tested-cfgs/SM6_TITANX/gpgpusim.config
+++ b/configs/tested-cfgs/SM6_TITANX/gpgpusim.config
@@ -8,6 +8,16 @@
-gpgpu_ptx_force_max_capability 61
-gpgpu_ignore_resources_limitation 1
+# Device Limits
+-gpgpu_stack_size_limit 1024
+-gpgpu_heap_size_limit 8388608
+-gpgpu_runtime_sync_depth_limit 2
+-gpgpu_runtime_pending_launch_count_limit 2048
+
+# Compute Capability
+-gpgpu_compute_capability_major 6
+-gpgpu_compute_capability_minor 1
+
# SASS execution (only supported with CUDA >= 4.0)
-gpgpu_ptx_convert_to_ptxplus 0
-gpgpu_ptx_save_converted_ptxplus 0
diff --git a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config
index 8ed4cd0..9de1050 100644
--- a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config
+++ b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config
@@ -13,6 +13,17 @@
-gpgpu_ptx_sim_mode 0
-gpgpu_ptx_force_max_capability 70
+
+# Device Limits
+-gpgpu_stack_size_limit 1024
+-gpgpu_heap_size_limit 8388608
+-gpgpu_runtime_sync_depth_limit 2
+-gpgpu_runtime_pending_launch_count_limit 2048
+
+# Compute Capability
+-gpgpu_compute_capability_major 7
+-gpgpu_compute_capability_minor 0
+
# SASS execution (only supported with CUDA >= 4.0)
-gpgpu_ptx_convert_to_ptxplus 0
-gpgpu_ptx_save_converted_ptxplus 0
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index cf16d9b..3a9d613 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -354,8 +354,8 @@ struct _cuda_device_id *GPGPUSim_Init()
cudaDeviceProp *prop = (cudaDeviceProp *) calloc(sizeof(cudaDeviceProp),1);
snprintf(prop->name,256,"GPGPU-Sim_v%s", g_gpgpusim_version_string );
- prop->major = 5;
- prop->minor = 2;
+ prop->major = the_gpu->compute_capability_major();
+ prop->minor = the_gpu->compute_capability_minor();
prop->totalGlobalMem = 0x80000000 /* 2 GB */;
prop->memPitch = 0;
if(prop->major >= 2) {
@@ -1120,11 +1120,11 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic
case 41:
*value= 0;
break;
- case 75:
- *value= 9 ;
+ case 75://cudaDevAttrComputeCapabilityMajor
+ *value= prop->major ;
break;
- case 76:
- *value= 3 ;
+ case 76://cudaDevAttrComputeCapabilityMinor
+ *value= prop->minor ;
break;
case 78:
*value= 0 ; //TODO: as of now, we dont support stream priorities.
@@ -1199,17 +1199,50 @@ __host__ cudaError_t CUDARTAPI cudaGetDevice(int *device)
*device = g_active_device;
return g_last_cudaError = cudaSuccess;
}
+
__host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit limit )
{
if(g_debug_execution >= 3){
announce_call(__my_func__);
}
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaSuccess;
+ _cuda_device_id *dev = GPGPUSim_Init();
+ const struct cudaDeviceProp *prop = dev->get_prop();
+ const gpgpu_sim_config& config=dev->get_gpgpu()->get_config();
+ switch(limit) {
+ case 0: // cudaLimitStackSize
+ *pValue=config.stack_limit();
+ break;
+ case 2: // cudaLimitMallocHeapSize
+ *pValue=config.heap_limit();
+ break;
+#if (CUDART_VERSION > 5050)
+ case 3: // cudaLimitDevRuntimeSyncDepth
+ if(prop->major > 2){
+ *pValue=config.sync_depth_limit();
+ break;
+ }
+ else{
+ printf("ERROR:Limit %s is not supported on this architecture \n",limit);
+ abort();
+ }
+ case 4: // cudaLimitDevRuntimePendingLaunchCount
+ if(prop->major > 2){
+ *pValue=config.pending_launch_count_limit();
+ break;
+ }
+ else{
+ printf("ERROR:Limit %s is not supported on this architecture \n",limit);
+ abort();
+ }
+#endif
+ default:
+ printf("ERROR:Limit %s unimplemented \n",limit);
+ abort();
+ }
+ return g_last_cudaError = cudaSuccess;
}
-
__host__ cudaError_t CUDARTAPI cudaStreamGetPriority ( cudaStream_t hStream, int* priority )
{
if(g_debug_execution >= 3){
@@ -1569,9 +1602,9 @@ __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream)
announce_call(__my_func__);
}
#if (CUDART_VERSION >= 3000)
- //synchronization required for application using external libraries without explicit synchronization in the code to
+ //per-stream synchronization required for application using external libraries without explicit synchronization in the code to
//avoid the stream_manager from spinning forever to destroy non-empty streams without making any forward progress.
- synchronize();
+ stream->synchronize();
g_stream_manager->destroy_stream(stream);
#endif
return g_last_cudaError = cudaSuccess;
@@ -2698,14 +2731,14 @@ cudaError_t cudaDeviceReset ( void ) {
return g_last_cudaError = cudaSuccess;
}
cudaError_t CUDARTAPI cudaDeviceSynchronize(void){
- // I don't know what this should do
if(g_debug_execution >= 3){
announce_call(__my_func__);
}
+ //Blocks until the device has completed all preceding requested tasks
+ synchronize();
return g_last_cudaError = cudaSuccess;
}
-
void CUDARTAPI __cudaRegisterFunction(
void **fatCubinHandle,
const char *hostFun,
@@ -3135,6 +3168,8 @@ __host__ cudaError_t CUDARTAPI cudaDeviceSetLimit(enum cudaLimit limit, size_t v
}
return g_last_cudaError = cudaSuccess;
}
+
+
#endif
#endif
@@ -3313,7 +3348,12 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun,
announce_call(__my_func__);
}
function_info *entry = context->get_kernel(hostFun);
- kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry);
+ gpgpu_t* gpu= context->get_device()->get_gpgpu();
+ /*
+ Passing a snapshot of the GPU's current texture mapping to the kernel's info
+ as kernels should use texture bindings present at the time of their launch.
+ */
+ kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry,gpu->getNameArrayMapping(),gpu->getNameInfoMapping());
if( entry == NULL ) {
printf("GPGPU-Sim PTX: ERROR launching kernel -- no PTX implementation found for %p\n", hostFun);
abort();
diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc
index 307df40..cebdb25 100644
--- a/src/abstract_hardware_model.cc
+++ b/src/abstract_hardware_model.cc
@@ -88,8 +88,6 @@ void checkpoint::load_global_mem(class memory_space *temp_mem, char * f1name)
fclose ( fp2 );
}
-
-
void checkpoint::store_global_mem(class memory_space * mem, char *fname, char * format)
{
@@ -192,6 +190,11 @@ gpgpu_t::gpgpu_t( const gpgpu_functional_sim_config &config )
checkpoint_CTA_t = m_function_model_config.get_checkpoint_CTA_t();
checkpoint_insn_Y = m_function_model_config.get_checkpoint_insn_Y();
+ // initialize texture mappings to empty
+ m_NameToTextureInfo.clear();
+ m_NameToCudaArray.clear();
+ m_TextureRefToName.clear();
+ m_NameToAttribute.clear();
if(m_function_model_config.get_ptx_inst_debug_to_file() != 0)
ptx_inst_debug_file = fopen(m_function_model_config.get_ptx_inst_debug_file(), "w");
@@ -688,7 +691,10 @@ unsigned g_kernel_launch_latency;
unsigned kernel_info_t::m_next_uid = 1;
-kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry )
+/*A snapshot of the texture mappings needs to be stored in the kernel's info as
+kernels should use the texture bindings seen at the time of launch and textures
+ can be bound/unbound asynchronously with respect to streams. */
+kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map<std::string, const struct cudaArray*> nameToCudaArray, std::map<std::string, const struct textureInfo*> nameToTextureInfo)
{
m_kernel_entry=entry;
m_grid_dim=gridDim;
@@ -708,6 +714,8 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *
m_launch_latency = g_kernel_launch_latency;
volta_cache_config_set=false;
+ m_NameToCudaArray = nameToCudaArray;
+ m_NameToTextureInfo = nameToTextureInfo;
}
kernel_info_t::~kernel_info_t()
diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h
index 28a22e0..e8716ab 100644
--- a/src/abstract_hardware_model.h
+++ b/src/abstract_hardware_model.h
@@ -212,7 +212,7 @@ public:
// m_num_cores_running=0;
// m_param_mem=NULL;
// }
- kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry );
+ kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map<std::string, const struct cudaArray*> nameToCudaArray, std::map<std::string, const struct textureInfo*> nameToTextureInfo);
~kernel_info_t();
void inc_running() { m_num_cores_running++; }
@@ -275,6 +275,23 @@ public:
std::list<class ptx_thread_info *> &active_threads() { return m_active_threads; }
class memory_space *get_param_memory() { return m_param_mem; }
+
+ //The following functions access texture bindings present at the kernel's launch
+
+ const struct cudaArray* get_texarray( const std::string &texname ) const
+ {
+ std::map<std::string,const struct cudaArray*>::const_iterator t=m_NameToCudaArray.find(texname);
+ assert(t != m_NameToCudaArray.end());
+ return t->second;
+ }
+
+ const struct textureInfo* get_texinfo( const std::string &texname ) const
+ {
+ std::map<std::string, const struct textureInfo*>::const_iterator t=m_NameToTextureInfo.find(texname);
+ assert(t != m_NameToTextureInfo.end());
+ return t->second;
+ }
+
private:
kernel_info_t( const kernel_info_t & ); // disable copy constructor
void operator=( const kernel_info_t & ); // disable copy operator
@@ -283,6 +300,10 @@ private:
unsigned m_uid;
static unsigned m_next_uid;
+
+ //These maps contain the snapshot of the texture mappings at kernel launch
+ std::map<std::string, const struct cudaArray*> m_NameToCudaArray;
+ std::map<std::string, const struct textureInfo*> m_NameToTextureInfo;
dim3 m_grid_dim;
dim3 m_block_dim;
@@ -580,8 +601,8 @@ public:
const struct textureInfo* get_texinfo( const std::string &texname ) const
{
- std::map<std::string, const struct textureInfo*>::const_iterator t=m_NameToTexureInfo.find(texname);
- assert(t != m_NameToTexureInfo.end());
+ std::map<std::string, const struct textureInfo*>::const_iterator t=m_NameToTextureInfo.find(texname);
+ assert(t != m_NameToTextureInfo.end());
return t->second;
}
@@ -594,6 +615,10 @@ public:
const gpgpu_functional_sim_config &get_config() const { return m_function_model_config; }
FILE* get_ptx_inst_debug_file() { return ptx_inst_debug_file; }
+
+ // These maps return the current texture mappings for the GPU at any given time.
+ std::map<std::string, const struct cudaArray*> getNameArrayMapping() {return m_NameToCudaArray;}
+ std::map<std::string, const struct textureInfo*> getNameInfoMapping() {return m_NameToTextureInfo;}
protected:
const gpgpu_functional_sim_config &m_function_model_config;
@@ -604,11 +629,11 @@ protected:
class memory_space *m_surf_mem;
unsigned long long m_dev_malloc;
-
+ // These maps contain the current texture mappings for the GPU at any given time.
std::map<std::string, std::set<const struct textureReference*> > m_NameToTextureRef;
std::map<const struct textureReference*, std::string> m_TextureRefToName;
std::map<std::string, const struct cudaArray*> m_NameToCudaArray;
- std::map<std::string, const struct textureInfo*> m_NameToTexureInfo;
+ std::map<std::string, const struct textureInfo*> m_NameToTextureInfo;
std::map<std::string, const struct textureReferenceAttr*> m_NameToAttribute;
};
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index ec68b5b..f7bb9cc 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -218,7 +218,7 @@ void gpgpu_t::gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* te
texInfo->Ty_numbits = intLOGB2(Ty);
texInfo->texel_size = texel_size;
texInfo->texel_size_numbits = intLOGB2(texel_size);
- m_NameToTexureInfo[texname] = texInfo;
+ m_NameToTextureInfo[texname] = texInfo;
}
void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref)
@@ -226,7 +226,7 @@ void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref)
//assumes bind-use-unbind-bind-use-unbind pattern
std::string texname = gpgpu_ptx_sim_findNamefromTexture(texref);
m_NameToCudaArray.erase(texname);
- m_NameToTexureInfo.erase(texname);
+ m_NameToTextureInfo.erase(texname);
}
unsigned g_assemble_code_next_pc=0;
@@ -1507,7 +1507,15 @@ static unsigned get_tex_datasize( const ptx_instruction *pI, ptx_thread_info *th
std::string texname = src1.name();
gpgpu_t *gpu = thread->get_gpu();
- const struct textureInfo* texInfo = gpu->get_texinfo(texname);
+ /*
+ For programs with many streams, textures can be bound and unbound
+ asynchronously. This means we need to use the kernel's "snapshot" of
+ the state of the texture mappings when it was launched (so that we
+ don't try to access the incorrect texture mapping if it's been updated,
+ or that we don't access a mapping that has been unbound).
+ */
+ kernel_info_t& k = thread->get_kernel();
+ const struct textureInfo* texInfo = k.get_texinfo(texname);
unsigned data_size = texInfo->texel_size;
return data_size;
@@ -1918,7 +1926,7 @@ kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,
struct dim3 blockDim,
gpgpu_t *gpu )
{
- kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry);
+ kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry,gpu->getNameArrayMapping(),gpu->getNameInfoMapping());
unsigned argcount=args.size();
unsigned argn=1;
for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) {
diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc
index b399133..917e7a8 100644
--- a/src/cuda-sim/cuda_device_runtime.cc
+++ b/src/cuda-sim/cuda_device_runtime.cc
@@ -198,7 +198,8 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info *
memory_space *device_kernel_param_mem;
//create child kernel_info_t and index it with parameter_buffer address
- device_grid = new kernel_info_t(config.grid_dim, config.block_dim, device_kernel_entry);
+ gpgpu_t* gpu=thread->get_gpu();
+ device_grid = new kernel_info_t(config.grid_dim, config.block_dim, device_kernel_entry, gpu->getNameArrayMapping(), gpu->getNameInfoMapping());
device_grid->launch_cycle = gpu_sim_cycle + gpu_tot_sim_cycle;
kernel_info_t & parent_grid = thread->get_kernel();
DEV_RUNTIME_REPORT("child kernel launched by " << parent_grid.name() << ", cta (" <<
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index c85654c..11001d7 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -5258,11 +5258,18 @@ void tex_impl( const ptx_instruction *pI, ptx_thread_info *thread )
if (!ptx_tex_regs) ptx_tex_regs = new ptx_reg_t[4];
unsigned nelem = src2.get_vect_nelem();
thread->get_vector_operand_values(src2, ptx_tex_regs, nelem); //ptx_reg should be 4 entry vector type...coordinates into texture
-
+ /*
+ For programs with many streams, textures can be bound and unbound
+ asynchronously. This means we need to use the kernel's "snapshot" of
+ the state of the texture mappings when it was launched (so that we
+ don't try to access the incorrect texture mapping if it's been updated,
+ or that we don't access a mapping that has been unbound).
+ */
gpgpu_t *gpu = thread->get_gpu();
+ kernel_info_t &k = thread->get_kernel();
const struct textureReference* texref = gpu->get_texref(texname);
- const struct cudaArray* cuArray = gpu->get_texarray(texname);
- const struct textureInfo* texInfo = gpu->get_texinfo(texname);
+ const struct cudaArray* cuArray = k.get_texarray(texname);
+ const struct textureInfo* texInfo = k.get_texinfo(texname);
const struct textureReferenceAttr* texAttr = gpu->get_texattr(texname);
//assume always 2D f32 input
diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc
index ec570bf..3687448 100644
--- a/src/gpgpu-sim/gpu-sim.cc
+++ b/src/gpgpu-sim/gpu-sim.cc
@@ -499,6 +499,12 @@ void gpgpu_sim_config::reg_options(option_parser_t opp)
option_parser_register(opp, "-liveness_message_freq", OPT_INT64, &liveness_message_freq,
"Minimum number of seconds between simulation liveness messages (0 = always print)",
"1");
+ option_parser_register(opp, "-gpgpu_compute_capability_major", OPT_UINT32, &gpgpu_compute_capability_major,
+ "Major compute capability version number",
+ "7");
+ option_parser_register(opp, "-gpgpu_compute_capability_minor", OPT_UINT32, &gpgpu_compute_capability_minor,
+ "Minor compute capability version number",
+ "0");
option_parser_register(opp, "-gpgpu_flush_l1_cache", OPT_BOOL, &gpgpu_flush_l1_cache,
"Flush L1 cache at the end of each kernel call",
"0");
@@ -532,6 +538,14 @@ void gpgpu_sim_config::reg_options(option_parser_t opp)
option_parser_register(opp, "-visualizer_zlevel", OPT_INT32,
&g_visualizer_zlevel, "Compression level of the visualizer output log (0=no comp, 9=highest)",
"6");
+ option_parser_register(opp, "-gpgpu_stack_size_limit", OPT_INT32, &stack_size_limit,
+ "GPU thread stack size", "1024" );
+ option_parser_register(opp, "-gpgpu_heap_size_limit", OPT_INT32, &heap_size_limit,
+ "GPU malloc heap size ", "8388608" );
+ option_parser_register(opp, "-gpgpu_runtime_sync_depth_limit", OPT_INT32, &runtime_sync_depth_limit,
+ "GPU device runtime synchronize depth", "2" );
+ option_parser_register(opp, "-gpgpu_runtime_pending_launch_count_limit", OPT_INT32, &runtime_pending_launch_count_limit,
+ "GPU device runtime pending launch count", "2048" );
option_parser_register(opp, "-trace_enabled", OPT_BOOL,
&Trace::enabled, "Turn on traces",
"0");
@@ -794,6 +808,16 @@ void gpgpu_sim::set_prop( cudaDeviceProp *prop )
m_cuda_properties = prop;
}
+int gpgpu_sim::compute_capability_major() const
+{
+ return m_config.gpgpu_compute_capability_major;
+}
+
+int gpgpu_sim::compute_capability_minor() const
+{
+ return m_config.gpgpu_compute_capability_minor;
+}
+
const struct cudaDeviceProp *gpgpu_sim::get_prop() const
{
return m_cuda_properties;
diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h
index 6ce5524..d4cc18f 100644
--- a/src/gpgpu-sim/gpu-sim.h
+++ b/src/gpgpu-sim/gpu-sim.h
@@ -337,6 +337,11 @@ public:
unsigned get_max_concurrent_kernel() const { return max_concurrent_kernel; }
unsigned checkpoint_option;
+ size_t stack_limit() const {return stack_size_limit; }
+ size_t heap_limit() const {return heap_size_limit; }
+ size_t sync_depth_limit() const {return runtime_sync_depth_limit; }
+ size_t pending_launch_count_limit() const {return runtime_pending_launch_count_limit;}
+
private:
void init_clock_domains(void );
@@ -377,8 +382,15 @@ private:
int gpu_stat_sample_freq;
int gpu_runtime_stat_flag;
+ // Device Limits
+ size_t stack_size_limit;
+ size_t heap_size_limit;
+ size_t runtime_sync_depth_limit;
+ size_t runtime_pending_launch_count_limit;
-
+ //gpu compute capability options
+ unsigned int gpgpu_compute_capability_major;
+ unsigned int gpgpu_compute_capability_minor;
unsigned long long liveness_message_freq;
friend class gpgpu_sim;
@@ -438,6 +450,8 @@ public:
int shared_mem_size() const;
int shared_mem_per_block() const;
+ int compute_capability_major() const;
+ int compute_capability_minor() const;
int num_registers_per_core() const;
int num_registers_per_block() const;
int wrp_size() const;
generated by cgit v1.3 (git 2.53.0) at 2026-07-11 02:32:57 +0000