Assigned debug level 3 to stream manager output.

[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 10013]

4 files changed, 76 insertions, 39 deletions

diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 101fe8c..8bb0e01 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -415,7 +415,8 @@ __host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size)
 {
    CUctx_st* context = GPGPUSim_Context();
    *devPtr = context->get_device()->get_gpgpu()->gpu_malloc(size);
-   printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *devPtr);
+   if(g_debug_execution >= 3)
+      printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *devPtr);
    if ( *devPtr  ) {
        return g_last_cudaError = cudaSuccess;
    } else { 
@@ -495,7 +496,8 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t cou
 {
    //CUctx_st *context = GPGPUSim_Context();
    //gpgpu_t *gpu = context->get_device()->get_gpgpu();
-   printf("GPGPU-Sim PTX: cudaMemcpy(): devPtr = %p\n", dst);
+   if(g_debug_execution >= 3)
+      printf("GPGPU-Sim PTX: cudaMemcpy(): devPtr = %p\n", dst);
    if( kind == cudaMemcpyHostToDevice ) 
        g_stream_manager->push( stream_operation(src,(size_t)dst,count,0) );
    else if( kind == cudaMemcpyDeviceToHost ) 
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 7353cfc..e7667b2 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -319,8 +319,10 @@ addr_t generic_to_global( addr_t addr )
 void* gpgpu_t::gpu_malloc( size_t size )
 {
    unsigned long long result = m_dev_malloc;
-   printf("GPGPU-Sim PTX: allocating %zu bytes on GPU starting at address 0x%Lx\n", size, m_dev_malloc );
-   fflush(stdout);
+   if(g_debug_execution >= 3) {
+      printf("GPGPU-Sim PTX: allocating %zu bytes on GPU starting at address 0x%Lx\n", size, m_dev_malloc );
+      fflush(stdout);
+   }
    m_dev_malloc += size;
    if (size%64) m_dev_malloc += (64 - size%64); //align to 64 byte boundaries
    return(void*) result;
@@ -329,8 +331,10 @@ void* gpgpu_t::gpu_malloc( size_t size )
 void* gpgpu_t::gpu_mallocarray( size_t size )
 {
    unsigned long long result = m_dev_malloc;
-   printf("GPGPU-Sim PTX: allocating %zu bytes on GPU starting at address 0x%Lx\n", size, m_dev_malloc );
-   fflush(stdout);
+   if(g_debug_execution >= 3) {
+      printf("GPGPU-Sim PTX: allocating %zu bytes on GPU starting at address 0x%Lx\n", size, m_dev_malloc );
+      fflush(stdout);
+   }
    m_dev_malloc += size;
    if (size%64) m_dev_malloc += (64 - size%64); //align to 64 byte boundaries
    return(void*) result;
@@ -339,50 +343,66 @@ void* gpgpu_t::gpu_mallocarray( size_t size )
 
 void gpgpu_t::memcpy_to_gpu( size_t dst_start_addr, const void *src, size_t count )
 {
-   printf("GPGPU-Sim PTX: copying %zu bytes from CPU[0x%Lx] to GPU[0x%Lx] ... ", count, (unsigned long long) src, (unsigned long long) dst_start_addr );
-   fflush(stdout);
+   if(g_debug_execution >= 3) {
+      printf("GPGPU-Sim PTX: copying %zu bytes from CPU[0x%Lx] to GPU[0x%Lx] ... ", count, (unsigned long long) src, (unsigned long long) dst_start_addr );
+      fflush(stdout);
+   }
    char *src_data = (char*)src;
    for (unsigned n=0; n < count; n ++ ) 
       m_global_mem->write(dst_start_addr+n,1, src_data+n,NULL,NULL);
-   printf( " done.\n");
-   fflush(stdout);
+   if(g_debug_execution >= 3) {
+      printf( " done.\n");
+      fflush(stdout);
+   }
 }
 
 void gpgpu_t::memcpy_from_gpu( void *dst, size_t src_start_addr, size_t count )
 {
-   printf("GPGPU-Sim PTX: copying %zu bytes from GPU[0x%Lx] to CPU[0x%Lx] ...", count, (unsigned long long) src_start_addr, (unsigned long long) dst );
-   fflush(stdout);
+   if(g_debug_execution >= 3) {
+      printf("GPGPU-Sim PTX: copying %zu bytes from GPU[0x%Lx] to CPU[0x%Lx] ...", count, (unsigned long long) src_start_addr, (unsigned long long) dst );
+      fflush(stdout);
+   }
    unsigned char *dst_data = (unsigned char*)dst;
    for (unsigned n=0; n < count; n ++ ) 
       m_global_mem->read(src_start_addr+n,1,dst_data+n);
-   printf( " done.\n");
-   fflush(stdout);
+   if(g_debug_execution >= 3) {
+      printf( " done.\n");
+      fflush(stdout);
+   }
 }
 
 void gpgpu_t::memcpy_gpu_to_gpu( size_t dst, size_t src, size_t count )
 {
-   printf("GPGPU-Sim PTX: copying %zu bytes from GPU[0x%Lx] to GPU[0x%Lx] ...", count, 
+   if(g_debug_execution >= 3) {
+      printf("GPGPU-Sim PTX: copying %zu bytes from GPU[0x%Lx] to GPU[0x%Lx] ...", count,
           (unsigned long long) src, (unsigned long long) dst );
-   fflush(stdout);
+      fflush(stdout);
+   }
    for (unsigned n=0; n < count; n ++ ) {
       unsigned char tmp;
       m_global_mem->read(src+n,1,&tmp); 
       m_global_mem->write(dst+n,1, &tmp,NULL,NULL);
    }
-   printf( " done.\n");
-   fflush(stdout);
+   if(g_debug_execution >= 3) {
+      printf( " done.\n");
+      fflush(stdout);
+   }
 }
 
 void gpgpu_t::gpu_memset( size_t dst_start_addr, int c, size_t count )
 {
-   printf("GPGPU-Sim PTX: setting %zu bytes of memory to 0x%x starting at 0x%Lx... ", 
+   if(g_debug_execution >= 3) {
+      printf("GPGPU-Sim PTX: setting %zu bytes of memory to 0x%x starting at 0x%Lx... ",
           count, (unsigned char) c, (unsigned long long) dst_start_addr );
-   fflush(stdout);
+      fflush(stdout);
+   }
    unsigned char c_value = (unsigned char)c;
    for (unsigned n=0; n < count; n ++ ) 
       m_global_mem->write(dst_start_addr+n,1,&c_value,NULL,NULL);
-   printf( " done.\n");
-   fflush(stdout);
+   if(g_debug_execution >= 3) {
+      printf( " done.\n");
+      fflush(stdout);
+   }
 }
 
 void ptx_print_insn( address_type pc, FILE *fp )
diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc
index de81df8..ba3b9aa 100644
--- a/src/gpgpusim_entrypoint.cc
+++ b/src/gpgpusim_entrypoint.cc
@@ -89,13 +89,17 @@ void *gpgpu_sim_thread_concurrent(void*)
     // concurrent kernel execution simulation thread
     g_the_gpu->init();
     do {
-        printf("GPGPU-Sim: *** simulation thread starting and spinning waiting for work ***\n");
-        fflush(stdout);
+       if(g_debug_execution >= 3) {
+          printf("GPGPU-Sim: *** simulation thread starting and spinning waiting for work ***\n");
+          fflush(stdout);
+       }
         while( g_stream_manager->empty() && !g_sim_done )
             ;
-        printf("GPGPU-Sim: ** START simulation thread (detected work) **\n");
-        g_stream_manager->print(stdout);
-        fflush(stdout);
+        if(g_debug_execution >= 3) {
+           printf("GPGPU-Sim: ** START simulation thread (detected work) **\n");
+           g_stream_manager->print(stdout);
+           fflush(stdout);
+        }
         pthread_mutex_lock(&g_sim_lock);
         g_sim_active = true;
         pthread_mutex_unlock(&g_sim_lock);
@@ -119,16 +123,20 @@ void *gpgpu_sim_thread_concurrent(void*)
             g_the_gpu->deadlock_check();
             active = g_the_gpu->active() || !g_stream_manager->empty();
         } while( active );
-        printf("GPGPU-Sim: ** STOP simulation thread (no work) **\n");
-        fflush(stdout);
+        if(g_debug_execution >= 3) {
+           printf("GPGPU-Sim: ** STOP simulation thread (no work) **\n");
+           fflush(stdout);
+        }
         if( sim_cycles ) 
             g_the_gpu->print_stats();
         pthread_mutex_lock(&g_sim_lock);
         g_sim_active = false;
         pthread_mutex_unlock(&g_sim_lock);
     } while( !g_sim_done );
-    printf("GPGPU-Sim: *** simulation thread exiting ***\n");
-    fflush(stdout);
+    if(g_debug_execution >= 3) {
+       printf("GPGPU-Sim: *** simulation thread exiting ***\n");
+       fflush(stdout);
+    }
     sem_post(&g_sim_signal_exit);
     return NULL;
 }
diff --git a/src/stream_manager.cc b/src/stream_manager.cc
index 9e031b0..85778c9 100644
--- a/src/stream_manager.cc
+++ b/src/stream_manager.cc
@@ -116,30 +116,36 @@ void stream_operation::do_operation( gpgpu_sim *gpu )
         return;
 
     assert(!m_done && m_stream);
-    printf("GPGPU-Sim API: stream %u performing ", m_stream->get_uid() );
+    if(g_debug_execution >= 3)
+       printf("GPGPU-Sim API: stream %u performing ", m_stream->get_uid() );
     switch( m_type ) {
     case stream_memcpy_host_to_device:
-        printf("memcpy host-to-device\n");
+        if(g_debug_execution >= 3)
+            printf("memcpy host-to-device\n");
         gpu->memcpy_to_gpu(m_device_address_dst,m_host_address_src,m_cnt);
         m_stream->record_next_done();
         break;
-    case stream_memcpy_device_to_host: 
-        printf("memcpy device-to-host\n");
+    case stream_memcpy_device_to_host:
+        if(g_debug_execution >= 3)
+            printf("memcpy device-to-host\n");
         gpu->memcpy_from_gpu(m_host_address_dst,m_device_address_src,m_cnt);
         m_stream->record_next_done();
         break;
     case stream_memcpy_device_to_device:
-        printf("memcpy device-to-device\n");
+        if(g_debug_execution >= 3)
+            printf("memcpy device-to-device\n");
         gpu->memcpy_gpu_to_gpu(m_device_address_dst,m_device_address_src,m_cnt); 
         m_stream->record_next_done();
         break;
     case stream_memcpy_to_symbol:
-        printf("memcpy to symbol\n");
+        if(g_debug_execution >= 3)
+            printf("memcpy to symbol\n");
         gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_src,m_cnt,m_offset,1,gpu);
         m_stream->record_next_done();
         break;
     case stream_memcpy_from_symbol:
-        printf("memcpy from symbol\n");
+        if(g_debug_execution >= 3)
+            printf("memcpy from symbol\n");
         gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_dst,m_cnt,m_offset,0,gpu);
         m_stream->record_next_done();
         break;
@@ -330,7 +336,8 @@ void stream_manager::push( stream_operation op )
         op.set_stream(&m_stream_zero);
         m_stream_zero.push(op);
     }
-    print_impl(stdout);
+    if(g_debug_execution >= 3)
+       print_impl(stdout);
     pthread_mutex_unlock(&m_lock);
     if( m_cuda_launch_blocking || stream == NULL ) {
         unsigned int wait_amount = 100;