summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--libcuda/cuda_runtime_api.cc25
-rw-r--r--src/cuda-sim/instructions.cc28
-rw-r--r--src/cuda-sim/ptx_sim.cc4
-rw-r--r--src/gpgpu-sim/visualizer.cc30
4 files changed, 62 insertions, 25 deletions
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 6e562fe..3e0d0da 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -290,6 +290,8 @@ struct cudaDeviceProp **gpgpu_cuda_devices;
static dim3 g_cudaGridDim;
static dim3 g_cudaBlockDim;
+static cudaStream_t g_stream_id = 0;
+
/*******************************************************************************
* *
* *
@@ -858,26 +860,33 @@ __host__ cudaError_t CUDARTAPI cudaLaunch(const char *symbol )
__host__ cudaError_t CUDARTAPI cudaStreamCreate(cudaStream_t *stream)
{
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaErrorUnknown;
+ printf("GPGPU-Sim PTX: WARNING: This stub implementation of %s can only support a single stream! \n", __my_func__);
+ assert(stream != NULL);
+ *stream = g_stream_id;
+ assert(g_stream_id == 0);
+ g_stream_id += 1;
+ return g_last_cudaError = cudaSuccess;
}
__host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream)
{
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaErrorUnknown;
+ printf("GPGPU-Sim PTX: WARNING: This stub implementation of %s can only support a single stream! \n", __my_func__);
+ assert(stream == 0);
+ g_stream_id -= 1;
+ assert(g_stream_id == 0);
+ return g_last_cudaError = cudaSuccess;
}
__host__ cudaError_t CUDARTAPI cudaStreamSynchronize(cudaStream_t stream)
{
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaErrorUnknown;
+ printf("GPGPU-Sim PTX: WARNING: This implementation of %s can only a stub! \n", __my_func__);
+ return g_last_cudaError = cudaSuccess; // it is stub because all cuda calls are synchronous
}
__host__ cudaError_t CUDARTAPI cudaStreamQuery(cudaStream_t stream)
{
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaErrorUnknown;
+ printf("GPGPU-Sim PTX: WARNING: This implementation of %s can only a stub! \n", __my_func__);
+ return g_last_cudaError = cudaSuccess; // it is always success because all cuda calls are synchronous
}
/*******************************************************************************
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index 5decd2b..7cdf0de 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -2721,7 +2721,33 @@ void tex_impl( const ptx_instruction *pI, ptx_thread_info *thread )
switch (dimension) {
case GEOM_MODIFIER_1D:
- x = ptx_tex_regs[0].u64;
+ width = cuArray->width;
+ height = cuArray->height;
+ if (texref->normalized) {
+ x_f32 = ptx_tex_regs[0].f32;
+ if (texref->addressMode[0] == cudaAddressModeClamp) {
+ x_f32 = (x_f32 > 1.0)? 1.0 : x_f32;
+ x_f32 = (x_f32 < 0.0)? 0.0 : x_f32;
+ } else if (texref->addressMode[0] == cudaAddressModeWrap) {
+ x_f32 = x_f32 - floor(x_f32);
+ }
+
+ if( texref->filterMode == cudaFilterModeLinear ) {
+ float xb = x_f32 * width - 0.5;
+ alpha = xb - floor(xb);
+ alpha = reduce_precision(alpha,9);
+ beta = 0.0;
+
+ x = (int)floor(xb);
+ y = 0;
+ } else {
+ x = (int) floor(x_f32 * width);
+ y = 0;
+ }
+ } else {
+ x = ptx_tex_regs[0].u64;
+ }
+ width *= (cuArray->desc.w+cuArray->desc.x+cuArray->desc.y+cuArray->desc.z)/8;
x *= (cuArray->desc.w+cuArray->desc.x+cuArray->desc.y+cuArray->desc.z)/8;
tex_array_index = tex_array_base + x;
diff --git a/src/cuda-sim/ptx_sim.cc b/src/cuda-sim/ptx_sim.cc
index a1a77fd..6a62783 100644
--- a/src/cuda-sim/ptx_sim.cc
+++ b/src/cuda-sim/ptx_sim.cc
@@ -316,6 +316,10 @@ static void print_reg( std::string name, ptx_reg_t value )
return;
}
const type_info *t = sym->type();
+ if( t == NULL ) {
+ printf("<unknown type> 0x%llx\n", (unsigned long long ) value.u64 );
+ return;
+ }
type_info_key ti = t->get_key();
switch ( ti.scalar_type() ) {
diff --git a/src/gpgpu-sim/visualizer.cc b/src/gpgpu-sim/visualizer.cc
index 3772fc2..d0f406c 100644
--- a/src/gpgpu-sim/visualizer.cc
+++ b/src/gpgpu-sim/visualizer.cc
@@ -129,14 +129,12 @@ void visualizer_options(option_parser_t opp)
void visualizer_printstat()
{
- static unsigned int *last_shader_cycle_distro = NULL;
gzFile visualizer_file = NULL; // gzFile is basically a pointer to a struct, so it is fine to initialize it as NULL
unsigned i;
if ( !g_visualizer_enabled )
return;
- if (!last_shader_cycle_distro)
- last_shader_cycle_distro = (unsigned int*) calloc(warp_size + 3, sizeof(unsigned int));
+ // initialize file name if it is not set
if ( g_visualizer_filename == NULL ) {
time_t curr_time;
time(&curr_time);
@@ -149,22 +147,19 @@ void visualizer_printstat()
}
char buf[1024];
snprintf(buf,1024,"gpgpusim_visualizer__%s.log.gz",date);
- visualizer_file = gzopen(buf, "w");
- if (visualizer_file == NULL) {
- printf("error - could not open visualizer trace file.\n");
- exit(1);
- }
- gzsetparams(visualizer_file, g_visualizer_zlevel, Z_DEFAULT_STRATEGY);
g_visualizer_filename = strdup(buf);
- } else {
- visualizer_file = gzopen(g_visualizer_filename,"a");
- if (visualizer_file == NULL) {
- printf("error - could not open visualizer trace file.\n");
- exit(1);
- }
- gzsetparams(visualizer_file, g_visualizer_zlevel, Z_DEFAULT_STRATEGY);
}
+ // clean the content of the visualizer log if it is the first time, otherwise attach at the end
+ static bool visualizer_first_printstat = true;
+ visualizer_file = gzopen(g_visualizer_filename, (visualizer_first_printstat)? "w" : "a");
+ if (visualizer_file == NULL) {
+ printf("error - could not open visualizer trace file.\n");
+ exit(1);
+ }
+ gzsetparams(visualizer_file, g_visualizer_zlevel, Z_DEFAULT_STRATEGY);
+ visualizer_first_printstat = false;
+
// instruction count per shader core
gzprintf(visualizer_file, "shaderinsncount: ");
for (i=0;i<gpu_n_shader;i++) {
@@ -299,6 +294,9 @@ void visualizer_printstat()
gzprintf(visualizer_file, "gpu_stall_by_MSHRwb: %d\n", gpu_stall_by_MSHRwb);
// warp divergence breakdown
+ static unsigned int *last_shader_cycle_distro = NULL;
+ if (!last_shader_cycle_distro)
+ last_shader_cycle_distro = (unsigned int*) calloc(warp_size + 3, sizeof(unsigned int));
time_vector_print_interval2gzfile(visualizer_file);
gzprintf(visualizer_file, "WarpDivergenceBreakdown:");
unsigned int total=0;
generated by cgit v1.3 (git 2.53.0) at 2026-07-10 23:32:17 +0000