summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--libcuda/Makefile3
-rw-r--r--libcuda/cuda_runtime_api.cc503
-rw-r--r--libopencl/opencl_runtime_api.cc45
-rw-r--r--src/abstract_hardware_model.cc10
-rw-r--r--src/abstract_hardware_model.h56
-rw-r--r--src/cuda-sim/cuda-sim.cc109
-rw-r--r--src/cuda-sim/cuda-sim.h17
-rw-r--r--src/cuda-sim/cuda_device_printf.cc4
-rw-r--r--src/cuda-sim/cuda_device_printf.h2
-rw-r--r--src/cuda-sim/instructions.cc38
-rw-r--r--src/cuda-sim/ptx_ir.cc2
-rw-r--r--src/cuda-sim/ptx_ir.h15
-rw-r--r--src/cuda-sim/ptx_loader.cc19
-rw-r--r--src/cuda-sim/ptx_loader.h11
-rw-r--r--src/cuda-sim/ptx_sim.cc1
-rw-r--r--src/cuda-sim/ptx_sim.h21
-rw-r--r--src/debug.cc12
-rw-r--r--src/debug.h1
-rw-r--r--src/gpgpu-sim/gpu-sim.cc12
-rw-r--r--src/gpgpu-sim/gpu-sim.h9
-rw-r--r--src/gpgpu-sim/shader.cc31
-rw-r--r--src/gpgpu-sim/shader.h30
-rw-r--r--src/gpgpusim_entrypoint.cc5
-rw-r--r--src/gpgpusim_entrypoint.h4
24 files changed, 434 insertions, 526 deletions
diff --git a/libcuda/Makefile b/libcuda/Makefile
index bdb48f4..3a72a8a 100644
--- a/libcuda/Makefile
+++ b/libcuda/Makefile
@@ -76,9 +76,10 @@ CPP = g++ $(SNOW)
CC = gcc &(SNOW)
CREATELIBRARY = 1
DEBUG ?= 0
-CXXFLAGS += -O3 -g -Wall -fPIC $(GL)
ifeq ($(DEBUG),1)
CXXFLAGS += -Wall -g -fPIC $(GL)
+else
+ CXXFLAGS += -O3 -g -Wall -fPIC $(GL)
endif
PROG =cuda
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 3a857d7..c49cfe8 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -115,7 +115,6 @@
#include <GL/gl.h>
#endif
#endif
-#include <dirent.h>
#define __CUDA_RUNTIME_API_H__
@@ -130,18 +129,14 @@
#include "../src/cuda-sim/ptx_parser.h"
#include "../src/gpgpusim_entrypoint.h"
-static void gpgpu_ptx_sim_load_gpu_kernels();
-static const struct gpgpu_ptx_sim_kernel_info * get_kernel_info(const char *kernel_key);
-static void register_function_implementation( const char *name, class function_info *impl );
-static void ptxinfo_cuda_addinfo();
-static void gpgpu_ptx_sim_add_ptxstring( unsigned fat_cubin_handle,
- const char *ptx_string,
- const char *cubin_string,
- const char *sourcefname,
- unsigned capability );
+static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsigned max_gaddr, gpgpu_t *gpu );
+static int load_constants( symbol_table *symtab, addr_t min_gaddr, gpgpu_t *gpu );
+
static kernel_info_t gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key,
gpgpu_ptx_sim_arg_list_t args,
- struct dim3 gridDim, struct dim3 blockDim );
+ struct dim3 gridDim,
+ struct dim3 blockDim,
+ struct CUctx_st* context );
/*DEVICE_BUILTIN*/
struct cudaArray
@@ -168,12 +163,7 @@ cudaError_t g_last_cudaError = cudaSuccess;
void register_ptx_function( const char *name, function_info *impl )
{
- register_function_implementation( name, impl );
-}
-
-extern "C" void ptxinfo_addinfo()
-{
- ptxinfo_cuda_addinfo();
+ // no longer need this
}
#if defined __APPLE__
@@ -211,12 +201,57 @@ struct _cuda_device_id {
return m_gpgpu->get_prop();
}
unsigned get_id() const { return m_id; }
+
+ gpgpu_sim *get_gpgpu() { return m_gpgpu; }
private:
unsigned m_id;
class gpgpu_sim *m_gpgpu;
struct _cuda_device_id *m_next;
};
+struct CUctx_st {
+ CUctx_st( _cuda_device_id *gpu ) { m_gpu = gpu; }
+
+ _cuda_device_id *get_device() { return m_gpu; }
+
+ void add_binary( symbol_table *symtab, unsigned fat_cubin_handle )
+ {
+ m_code[fat_cubin_handle] = symtab;
+ m_last_fat_cubin_handle = fat_cubin_handle;
+ }
+
+ void add_ptxinfo( const char *deviceFun, const struct gpgpu_ptx_sim_kernel_info &info )
+ {
+ symbol *s = m_code[m_last_fat_cubin_handle]->lookup(deviceFun);
+ assert( s != NULL );
+ function_info *f = s->get_pc();
+ assert( f != NULL );
+ f->set_kernel_info(info);
+ }
+
+ void register_function( unsigned fat_cubin_handle, const char *hostFun, const char *deviceFun )
+ {
+ symbol *s = m_code[fat_cubin_handle]->lookup(deviceFun);
+ assert( s != NULL );
+ function_info *f = s->get_pc();
+ assert( f != NULL );
+ m_kernel_lookup[hostFun] = f;
+ }
+
+ function_info *get_kernel(const char *hostFun)
+ {
+ std::map<const void*,function_info*>::iterator i=m_kernel_lookup.find(hostFun);
+ assert( i != m_kernel_lookup.end() );
+ return i->second;
+ }
+
+private:
+ _cuda_device_id *m_gpu; // selected gpu
+ std::map<unsigned,symbol_table*> m_code; // fat binary handle => global symbol table
+ unsigned m_last_fat_cubin_handle;
+ std::map<const void*,function_info*> m_kernel_lookup; // unique id (CUDA app function address) => kernel entry point
+};
+
class _cuda_device_id *GPGPUSim_Init()
{
static _cuda_device_id *the_device = NULL;
@@ -247,12 +282,33 @@ class _cuda_device_id *GPGPUSim_Init()
#endif
the_gpu->set_prop(prop);
the_device = new _cuda_device_id(the_gpu);
-
- gpgpu_ptx_sim_load_gpu_kernels();
}
return the_device;
}
+static CUctx_st* GPGPUSim_Context()
+{
+ static CUctx_st *the_context = NULL;
+ if( the_context == NULL ) {
+ _cuda_device_id *the_gpu = GPGPUSim_Init();
+ the_context = new CUctx_st(the_gpu);
+ }
+ return the_context;
+}
+
+extern "C" void ptxinfo_addinfo()
+{
+ if( !strcmp("__cuda_dummy_entry__",get_ptxinfo_kname()) ) {
+ // this string produced by ptxas for empty ptx files (e.g., bandwidth test)
+ clear_ptxinfo();
+ return;
+ }
+ CUctx_st *context = GPGPUSim_Context();
+ print_ptxinfo();
+ context->add_ptxinfo( get_ptxinfo_kname(), get_ptxinfo_kinfo() );
+ clear_ptxinfo();
+}
+
void cuda_not_implemented( const char* func, unsigned line )
{
fflush(stdout);
@@ -378,8 +434,8 @@ extern "C" {
__host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size)
{
- GPGPUSim_Init();
- *devPtr = gpgpu_ptx_sim_malloc(size);
+ CUctx_st* context = GPGPUSim_Context();
+ *devPtr = context->get_device()->get_gpgpu()->gpgpu_ptx_sim_malloc(size);
printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *devPtr);
if ( *devPtr ) {
return g_last_cudaError = cudaSuccess;
@@ -388,8 +444,9 @@ __host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size)
}
}
-__host__ cudaError_t CUDARTAPI cudaMallocHost(void **ptr, size_t size){
- GPGPUSim_Init();
+__host__ cudaError_t CUDARTAPI cudaMallocHost(void **ptr, size_t size)
+{
+ GPGPUSim_Context();
*ptr = malloc(size);
if ( *ptr ) {
return g_last_cudaError = cudaSuccess;
@@ -399,10 +456,10 @@ __host__ cudaError_t CUDARTAPI cudaMallocHost(void **ptr, size_t size){
}
__host__ cudaError_t CUDARTAPI cudaMallocPitch(void **devPtr, size_t *pitch, size_t width, size_t height)
{
- GPGPUSim_Init();
unsigned malloc_width_inbytes = width;
printf("GPGPU-Sim PTX: cudaMallocPitch (width = %d)\n", malloc_width_inbytes);
- *devPtr = gpgpu_ptx_sim_malloc(malloc_width_inbytes*height);
+ CUctx_st* ctx = GPGPUSim_Context();
+ *devPtr = ctx->get_device()->get_gpgpu()->gpgpu_ptx_sim_malloc(malloc_width_inbytes*height);
pitch[0] = malloc_width_inbytes;
if ( *devPtr ) {
return g_last_cudaError = cudaSuccess;
@@ -414,14 +471,14 @@ __host__ cudaError_t CUDARTAPI cudaMallocHost(void **ptr, size_t size){
__host__ cudaError_t CUDARTAPI cudaMallocArray(struct cudaArray **array, const struct cudaChannelFormatDesc *desc, size_t width, size_t height __dv(1))
{
unsigned size = width * height * ((desc->x + desc->y + desc->z + desc->w)/8);
- GPGPUSim_Init();
+ CUctx_st* context = GPGPUSim_Context();
(*array) = (struct cudaArray*) malloc(sizeof(struct cudaArray));
(*array)->desc = *desc;
(*array)->width = width;
(*array)->height = height;
(*array)->size = size;
(*array)->dimensions = 2;
- ((*array)->devPtr32)= (int) (long long)gpgpu_ptx_sim_mallocarray(size);
+ ((*array)->devPtr32)= (int) (long long)context->get_device()->get_gpgpu()->gpgpu_ptx_sim_mallocarray(size);
printf("GPGPU-Sim PTX: cudaMallocArray: devPtr32 = %d\n", ((*array)->devPtr32));
((*array)->devPtr) = (void*) (long long) ((*array)->devPtr32);
if ( ((*array)->devPtr) ) {
@@ -457,14 +514,15 @@ __host__ cudaError_t CUDARTAPI cudaFree(void *devPtr)
__host__ cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind)
{
- gpgpu_ptx_sim_init_memory();
+ CUctx_st *context = GPGPUSim_Context();
+ gpgpu_t *gpu = context->get_device()->get_gpgpu();
printf("GPGPU-Sim PTX: cudaMemcpy(): devPtr = %p\n", dst);
if( kind == cudaMemcpyHostToDevice )
- gpgpu_ptx_sim_memcpy_to_gpu( (size_t)dst, src, count );
+ gpu->gpgpu_ptx_sim_memcpy_to_gpu( (size_t)dst, src, count );
else if( kind == cudaMemcpyDeviceToHost )
- gpgpu_ptx_sim_memcpy_from_gpu( dst, (size_t)src, count );
+ gpu->gpgpu_ptx_sim_memcpy_from_gpu( dst, (size_t)src, count );
else if( kind == cudaMemcpyDeviceToDevice )
- gpgpu_ptx_sim_memcpy_gpu_to_gpu( (size_t)dst, (size_t)src, count );
+ gpu->gpgpu_ptx_sim_memcpy_gpu_to_gpu( (size_t)dst, (size_t)src, count );
else {
printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported cudaMemcpyKind\n");
abort();
@@ -474,15 +532,16 @@ __host__ cudaError_t CUDARTAPI cudaFree(void *devPtr)
__host__ cudaError_t CUDARTAPI cudaMemcpyToArray(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind)
{
+ CUctx_st *context = GPGPUSim_Context();
+ gpgpu_t *gpu = context->get_device()->get_gpgpu();
size_t size = count;
printf("GPGPU-Sim PTX: cudaMemcpyToArray\n");
- gpgpu_ptx_sim_init_memory();
if( kind == cudaMemcpyHostToDevice )
- gpgpu_ptx_sim_memcpy_to_gpu( (size_t)(dst->devPtr), src, size);
+ gpu->gpgpu_ptx_sim_memcpy_to_gpu( (size_t)(dst->devPtr), src, size);
else if( kind == cudaMemcpyDeviceToHost )
- gpgpu_ptx_sim_memcpy_from_gpu( dst->devPtr, (size_t)src, size);
+ gpu->gpgpu_ptx_sim_memcpy_from_gpu( dst->devPtr, (size_t)src, size);
else if( kind == cudaMemcpyDeviceToDevice )
- gpgpu_ptx_sim_memcpy_gpu_to_gpu( (size_t)(dst->devPtr), (size_t)src, size);
+ gpu->gpgpu_ptx_sim_memcpy_gpu_to_gpu( (size_t)(dst->devPtr), (size_t)src, size);
else {
printf("GPGPU-Sim PTX: cudaMemcpyToArray - ERROR : unsupported cudaMemcpyKind\n");
abort();
@@ -508,17 +567,18 @@ __host__ cudaError_t CUDARTAPI cudaFree(void *devPtr)
__host__ cudaError_t CUDARTAPI cudaMemcpy2D(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind)
{
- gpgpu_ptx_sim_init_memory();
+ CUctx_st *context = GPGPUSim_Context();
+ gpgpu_t *gpu = context->get_device()->get_gpgpu();
struct cudaArray *cuArray_ptr;
size_t size = spitch*height;
cuArray_ptr = (cudaArray*)dst;
gpgpusim_ptx_assert( (dpitch==spitch), "different src and dst pitch not supported yet" );
if( kind == cudaMemcpyHostToDevice )
- gpgpu_ptx_sim_memcpy_to_gpu( (size_t)dst, src, size );
+ gpu->gpgpu_ptx_sim_memcpy_to_gpu( (size_t)dst, src, size );
else if( kind == cudaMemcpyDeviceToHost )
- gpgpu_ptx_sim_memcpy_from_gpu( dst, (size_t)src, size );
+ gpu->gpgpu_ptx_sim_memcpy_from_gpu( dst, (size_t)src, size );
else if( kind == cudaMemcpyDeviceToDevice )
- gpgpu_ptx_sim_memcpy_gpu_to_gpu( (size_t)dst, (size_t)src, size);
+ gpu->gpgpu_ptx_sim_memcpy_gpu_to_gpu( (size_t)dst, (size_t)src, size);
else {
printf("GPGPU-Sim PTX: cudaMemcpy2D - ERROR : unsupported cudaMemcpyKind\n");
abort();
@@ -529,8 +589,9 @@ __host__ cudaError_t CUDARTAPI cudaFree(void *devPtr)
__host__ cudaError_t CUDARTAPI cudaMemcpy2DToArray(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind)
{
+ CUctx_st *context = GPGPUSim_Context();
+ gpgpu_t *gpu = context->get_device()->get_gpgpu();
size_t size = spitch*height;
- gpgpu_ptx_sim_init_memory();
size_t channel_size = dst->desc.w+dst->desc.x+dst->desc.y+dst->desc.z;
gpgpusim_ptx_assert( ((channel_size%8) == 0), "none byte multiple destination channel size not supported (sz=%u)", channel_size );
unsigned elem_size = channel_size/8;
@@ -541,11 +602,11 @@ __host__ cudaError_t CUDARTAPI cudaFree(void *devPtr)
gpgpusim_ptx_assert( (elem_size*dst->width == width), "partial copy not supported" );
gpgpusim_ptx_assert( (spitch == width), "spitch != width not supported" );
if( kind == cudaMemcpyHostToDevice )
- gpgpu_ptx_sim_memcpy_to_gpu( (size_t)(dst->devPtr), src, size);
+ gpu->gpgpu_ptx_sim_memcpy_to_gpu( (size_t)(dst->devPtr), src, size);
else if( kind == cudaMemcpyDeviceToHost )
- gpgpu_ptx_sim_memcpy_from_gpu( dst->devPtr, (size_t)src, size);
+ gpu->gpgpu_ptx_sim_memcpy_from_gpu( dst->devPtr, (size_t)src, size);
else if( kind == cudaMemcpyDeviceToDevice )
- gpgpu_ptx_sim_memcpy_gpu_to_gpu( (size_t)dst->devPtr, (size_t)src, size);
+ gpu->gpgpu_ptx_sim_memcpy_gpu_to_gpu( (size_t)dst->devPtr, (size_t)src, size);
else {
printf("GPGPU-Sim PTX: cudaMemcpy2D - ERROR : unsupported cudaMemcpyKind\n");
abort();
@@ -571,18 +632,20 @@ __host__ cudaError_t CUDARTAPI cudaFree(void *devPtr)
__host__ cudaError_t CUDARTAPI cudaMemcpyToSymbol(const char *symbol, const void *src, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyHostToDevice))
{
+ CUctx_st *context = GPGPUSim_Context();
assert(kind == cudaMemcpyHostToDevice);
printf("GPGPU-Sim PTX: cudaMemcpyToSymbol: symbol = %p\n", symbol);
- gpgpu_ptx_sim_memcpy_symbol(symbol,src,count,offset,1);
+ gpgpu_ptx_sim_memcpy_symbol(symbol,src,count,offset,1,context->get_device()->get_gpgpu());
return g_last_cudaError = cudaSuccess;
}
__host__ cudaError_t CUDARTAPI cudaMemcpyFromSymbol(void *dst, const char *symbol, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToHost))
{
+ CUctx_st *context = GPGPUSim_Context();
assert(kind == cudaMemcpyDeviceToHost);
printf("GPGPU-Sim PTX: cudaMemcpyFromSymbol: symbol = %p\n", symbol);
- gpgpu_ptx_sim_memcpy_symbol(symbol,dst,count,offset,0);
+ gpgpu_ptx_sim_memcpy_symbol(symbol,dst,count,offset,0,context->get_device()->get_gpgpu());
return g_last_cudaError = cudaSuccess;
}
@@ -645,7 +708,9 @@ __host__ cudaError_t CUDARTAPI cudaFree(void *devPtr)
__host__ cudaError_t CUDARTAPI cudaMemset(void *mem, int c, size_t count)
{
- gpgpu_ptx_sim_memset((size_t)mem, c, count);
+ CUctx_st *context = GPGPUSim_Context();
+ gpgpu_t *gpu = context->get_device()->get_gpgpu();
+ gpu->gpgpu_ptx_sim_memset((size_t)mem, c, count);
return g_last_cudaError = cudaSuccess;
}
@@ -866,16 +931,17 @@ __host__ cudaError_t CUDARTAPI cudaSetupArgument(const void *arg, size_t size, s
}
-__host__ cudaError_t CUDARTAPI cudaLaunch(const char *symbol )
+__host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun )
{
+ CUctx_st* context = GPGPUSim_Context();
char *mode = getenv("PTX_SIM_MODE_FUNC");
if( mode )
sscanf(mode,"%u", &g_ptx_sim_mode);
- printf("\nGPGPU-Sim PTX: cudaLaunch for %p (mode=%s)\n", symbol,
+ printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s)\n", hostFun,
g_ptx_sim_mode?"functional simulation":"performance simulation");
gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" );
kernel_config &config = g_cuda_launch_stack.back();
- kernel_info_t grid = gpgpu_cuda_ptx_sim_init_grid(symbol,config.get_args(),config.grid_dim(),config.block_dim());
+ kernel_info_t grid = gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context);
if( g_ptx_sim_mode )
gpgpu_cuda_ptx_sim_main_func( grid, config.grid_dim(), config.block_dim(), config.get_args() );
else
@@ -1039,27 +1105,62 @@ int CUDARTAPI __cudaSynchronizeThreads(void**, void*)
void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin )
{
+ CUctx_st *context = GPGPUSim_Context();
+
static unsigned next_fat_bin_handle = 1;
+ static unsigned source_num=1;
unsigned fat_cubin_handle = next_fat_bin_handle++;
#if (CUDART_VERSION >= 2010)
__cudaFatCudaBinary *info = (__cudaFatCudaBinary *)fatCubin;
assert( info->version >= 3 );
unsigned num_ptx_versions=0;
- //unsigned max_capability=0;
- //unsigned selected_capability_offset=(unsigned)-1;
+ unsigned max_capability=0;
+ unsigned selected_capability=0;
+ bool found=false;
while( info->ptx[num_ptx_versions].gpuProfileName != NULL ) {
unsigned capability=0;
sscanf(info->ptx[num_ptx_versions].gpuProfileName,"compute_%u",&capability);
-
printf("GPGPU-Sim PTX: __cudaRegisterFatBinary found PTX versions for '%s', ", info->ident);
- printf("capability = %s\n", info->ptx[num_ptx_versions].gpuProfileName );
- gpgpu_ptx_sim_add_ptxstring( fat_cubin_handle, info->ptx[num_ptx_versions].ptx, info->cubin[num_ptx_versions].cubin, info->ident, capability );
-
+ printf("capability = %s\n", info->ptx[num_ptx_versions].gpuProfileName );
+ if( capability > max_capability ) {
+ found = true;
+ max_capability=capability;
+ selected_capability = num_ptx_versions;
+ }
+ if( capability == context->get_device()->get_gpgpu()->get_forced_capability() ) {
+ found = true;
+ selected_capability = num_ptx_versions;
+ }
num_ptx_versions++;
}
+ if( found ) {
+ printf("GPGPU-Sim PTX: Loading PTX for %s, capability = %s\n",
+ info->ident, info->ptx[selected_capability].gpuProfileName );
+ symbol_table *symtab;
+ const char *ptx = info->ptx[selected_capability].ptx;
+ if(context->get_device()->get_gpgpu()->convert_to_ptxplus() ) {
+ char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_to_ptxplus(ptx, info->cubin[selected_capability].cubin, source_num++,
+ context->get_device()->get_gpgpu()->saved_converted_ptxplus());
+ symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxplus_str,source_num);
+ context->add_binary(symtab,fat_cubin_handle);
+ gpgpu_ptxinfo_load_from_string(ptx,source_num);
+ delete[] ptxplus_str;
+ } else {
+ symtab=gpgpu_ptx_sim_load_ptx_from_string(ptx,source_num);
+ context->add_binary(symtab,fat_cubin_handle);
+ gpgpu_ptxinfo_load_from_string( ptx, source_num );
+ }
+ source_num++;
+ load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu());
+ load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu());
+ } else {
+ printf("GPGPU-Sim PTX: ERROR ** did not find an appropriate PTX in Cubin\n");
+ abort();
+ }
#endif
return (void**)fat_cubin_handle;
}
+
void __cudaUnregisterFatBinary(void **fatCubinHandle)
{
;
@@ -1078,8 +1179,11 @@ void CUDARTAPI __cudaRegisterFunction(
dim3 *gDim
)
{
- gpgpu_ptx_sim_register_kernel(fatCubinHandle,hostFun,deviceFun);
- return;
+ CUctx_st *context = GPGPUSim_Context();
+ unsigned fat_cubin_handle = (unsigned)(unsigned long long)fatCubinHandle;
+ printf("GPGPU-Sim PTX: __cudaRegisterFunction %s : hostFun 0x%p, fat_cubin_handle = %u\n",
+ deviceFun, hostFun, fat_cubin_handle);
+ context->register_function( fat_cubin_handle, hostFun, deviceFun );
}
extern void __cudaRegisterVar(
@@ -1166,7 +1270,7 @@ cudaError_t cudaGLMapBufferObject(void** devPtr, GLuint bufferObj)
{
#ifdef OPENGL_SUPPORT
GLint buffer_size=0;
- GPGPUSim_Init();
+ CUctx_st* ctx = GPGPUSim_Context();
glbmap_entry_t *p = g_glbmap;
while ( p && p->m_bufferObj != bufferObj )
@@ -1175,7 +1279,7 @@ cudaError_t cudaGLMapBufferObject(void** devPtr, GLuint bufferObj)
glBindBuffer(GL_ARRAY_BUFFER,bufferObj);
glGetBufferParameteriv(GL_ARRAY_BUFFER,GL_BUFFER_SIZE,&buffer_size);
assert( buffer_size != 0 );
- *devPtr = gpgpu_ptx_sim_malloc(buffer_size);
+ *devPtr = ctx->get_device()->get_gpgpu()->gpgpu_ptx_sim_malloc(buffer_size);
// create entry and insert to front of list
glbmap_entry_t *n = (glbmap_entry_t *) calloc(1,sizeof(glbmap_entry_t));
@@ -1274,9 +1378,11 @@ cudaError_t CUDARTAPI cudaSetDeviceFlags( int flags )
return g_last_cudaError = cudaErrorUnknown;
}
-cudaError_t CUDARTAPI cudaFuncGetAttributes(struct cudaFuncAttributes *attr, const char *func)
+cudaError_t CUDARTAPI cudaFuncGetAttributes(struct cudaFuncAttributes *attr, const char *hostFun )
{
- const struct gpgpu_ptx_sim_kernel_info *kinfo = get_kernel_info(func);
+ CUctx_st *context = GPGPUSim_Context();
+ function_info *entry = context->get_kernel(hostFun);
+ const struct gpgpu_ptx_sim_kernel_info *kinfo = entry->get_kernel_info();
attr->sharedSizeBytes = kinfo->smem;
attr->constSizeBytes = kinfo->cmem;
@@ -1366,18 +1472,6 @@ int CUDARTAPI __cudaSynchronizeThreads(void**, void*)
-struct ptx_info_t {
- unsigned fat_cubin_handle;
- char *str;
- char *cubin_str;
- char *fname;
- ptx_info_t *next;
- unsigned capability;
-};
-
-static ptx_info_t *g_ptx_source_array = NULL;
-static unsigned g_ptx_max_capability = 0;
-
extern "C" int ptx_parse();
extern "C" int ptx__scan_string(const char*);
extern "C" FILE *ptx_in;
@@ -1417,69 +1511,12 @@ function_info *get_kernel(const char *kernel_key, std::string &kernel_func_name_
return (*g_kernel_name_to_function_lookup)[kernel_func_name_mangled];
}
-void gpgpu_ptx_sim_register_kernel(void **fatCubinHandle,const char *hostFun, const char *deviceFun)
-{
- const void* key=hostFun;
- print_splash();
- if ( g_host_to_kernel_entrypoint_name_lookup == NULL )
- g_host_to_kernel_entrypoint_name_lookup = new std::map<const void*,std::string>;
- if( g_kernel_name_to_function_lookup == NULL )
- g_kernel_name_to_function_lookup = new std::map<std::string,function_info*>;
- if ( g_host_to_kernel_entrypoint_name_lookup->find(key) !=
- g_host_to_kernel_entrypoint_name_lookup->end() ) {
- printf("GPGPU-Sim Loader error: Don't know how to identify PTX kernels during cudaLaunch\n"
- " for this application.\n");
- abort();
- }
- (*g_host_to_kernel_entrypoint_name_lookup)[key] = deviceFun;
- if( g_kernel_name_to_function_lookup->find(deviceFun) ==
- g_kernel_name_to_function_lookup->end() ) {
- (*g_kernel_name_to_function_lookup)[deviceFun] = NULL; // we set this later, set keys now for error checking
- }
-
- printf("GPGPU-Sim PTX: __cudaRegisterFunction %s : 0x%Lx\n", deviceFun, (unsigned long long)hostFun);
-}
-
-void ptxinfo_cuda_addinfo()
-{
- if( !strcmp("__cuda_dummy_entry__",get_ptxinfo_kname()) ) {
- // this string produced by ptxas for empty ptx files (e.g., bandwidth test)
- clear_ptxinfo();
- return;
- }
- if ( g_kernel_name_to_function_lookup ) {
- std::map<std::string,function_info*>::iterator i=g_kernel_name_to_function_lookup->find(get_ptxinfo_kname());
- if ( (g_kernel_name_to_function_lookup == NULL) || (i == g_kernel_name_to_function_lookup->end()) ) {
- printf ("GPGPU-Sim PTX: ERROR ** implementation for '%s' not found.\n", get_ptxinfo_kname() );
- abort();
- } else {
- print_ptxinfo();
- function_info *fi = i->second;
- assert(fi!=NULL);
- fi->set_kernel_info(get_ptxinfo_kinfo());
- }
- } else {
- printf ("GPGPU-Sim PTX: ERROR ** Kernel '%s' not found (no kernels registered).\n", get_ptxinfo_kname() );
- abort();
- }
- clear_ptxinfo();
-}
-static void register_function_implementation( const char *name, function_info *impl )
+void ptxinfo_cuda_addinfo( CUctx_st *context )
{
- printf("GPGPU-Sim PTX: parsing function %s\n", name );
- if( g_kernel_name_to_function_lookup == NULL )
- g_kernel_name_to_function_lookup = new std::map<std::string,function_info*>;
-
- std::map<std::string,function_info*>::iterator i_kernel = g_kernel_name_to_function_lookup->find(name);
- if (i_kernel != g_kernel_name_to_function_lookup->end() && i_kernel->second != NULL) {
- printf("GPGPU-Sim PTX: WARNING: Function already parsed once. Overwriting.\n");
- }
- (*g_kernel_name_to_function_lookup)[name] = impl;
}
-
-static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsigned max_gaddr)
+static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsigned max_gaddr, gpgpu_t *gpu )
{
printf( "GPGPU-Sim PTX: loading globals with explicit initializers... \n" );
fflush(stdout);
@@ -1503,7 +1540,7 @@ static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsign
operand_info op = *i;
ptx_reg_t value = op.get_literal_value();
assert( (addr+offset+nbytes) < min_gaddr ); // min_gaddr is start of "heap" for cudaMalloc
- g_global_mem->write(addr+offset,nbytes,&value,NULL,NULL); // assuming little endian here
+ gpu->get_global_memory()->write(addr+offset,nbytes,&value,NULL,NULL); // assuming little endian here
offset+=nbytes;
ng_bytes+=nbytes;
}
@@ -1515,7 +1552,7 @@ static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsign
return ng_bytes;
}
-static int load_constants( symbol_table *symtab, addr_t min_gaddr )
+static int load_constants( symbol_table *symtab, addr_t min_gaddr, gpgpu_t *gpu )
{
printf( "GPGPU-Sim PTX: loading constants with explicit initializers... " );
fflush(stdout);
@@ -1547,7 +1584,7 @@ static int load_constants( symbol_table *symtab, addr_t min_gaddr )
unsigned addr=constant->get_address() + nbytes_written;
assert( addr+nbytes < min_gaddr );
- g_global_mem->write(addr,nbytes,&value,NULL,NULL); // assume little endian (so u8 is the first byte in u32)
+ gpu->get_global_memory()->write(addr,nbytes,&value,NULL,NULL); // assume little endian (so u8 is the first byte in u32)
nc_bytes+=nbytes;
nbytes_written += nbytes;
}
@@ -1558,199 +1595,19 @@ static int load_constants( symbol_table *symtab, addr_t min_gaddr )
return nc_bytes;
}
-static FILE *open_ptxinfo (const char* ptx_filename)
+kernel_info_t gpgpu_cuda_ptx_sim_init_grid( const char *hostFun,
+ gpgpu_ptx_sim_arg_list_t args,
+ struct dim3 gridDim,
+ struct dim3 blockDim,
+ CUctx_st* context )
{
- const int ptx_fnamelen = strlen(ptx_filename);
- char *ptxi_fname = new char[ptx_fnamelen+5];
- strcpy (ptxi_fname, ptx_filename);
- strcpy (ptxi_fname+ptx_fnamelen, "info");
-
- //ptxinfo_debug=1;
- g_ptxinfo_filename = ptxi_fname;
- FILE *f = fopen (ptxi_fname, "rt");
- return f;
-}
-
-
-static int ptx_file_filter(
-#if !defined(__APPLE__)
- const
-#endif
- struct dirent *de )
-{
- const char *tmp = strstr(de->d_name,".ptx");
- if ( tmp != NULL && tmp[4] == 0 ) {
- return 1;
- }
- return 0;
-}
-
-extern unsigned g_ptx_force_max_capability;
-extern int g_ptx_convert_to_ptxplus;
-
-void gpgpu_ptx_sim_add_ptxstring( unsigned fat_cubin_handle, const char *ptx_string, const char *cubin_string, const char *sourcefname, unsigned capability )
-{
- ptx_info_t *t = new ptx_info_t;
- t->next = NULL;
- t->str = strdup(ptx_string);
- t->fat_cubin_handle = fat_cubin_handle;
- if (cubin_string != NULL) {
- t->cubin_str = strdup(cubin_string);
- } else {
- if(g_ptx_convert_to_ptxplus != 0) {
- printf("GPGPU-Sim PTX: ERROR cubin information, required for ptxplus, missing from fat bin object\n");
- printf("GPGPU-Sim PTX: ptxplus currently requires CUDA 2.3 or earlier... disable ptxplus by removing\n");
- printf("GPGPU-Sim PTX: \'-gpgpu_ptx_convert_to_ptxplus 1\' in \'gpgpusim.config\'\n");
- abort();
- }
- t->cubin_str = NULL;
- }
- t->fname = strdup(sourcefname);
- t->capability = capability;
-
- // Set overall max capability
- if(g_ptx_max_capability < capability)
- g_ptx_max_capability = capability;
-
- // put ptx source into a fifo
- if (g_ptx_source_array == NULL) {
- // first ptx source
- g_ptx_source_array = t;
- } else {
- // insert subsequent ptx source at the end of queue
- ptx_info_t *l_ptx_source = g_ptx_source_array;
- while (l_ptx_source->next != NULL) {
- l_ptx_source = l_ptx_source->next;
- }
- l_ptx_source->next = t;
- }
-}
-
-void gpgpu_ptx_sim_load_gpu_kernels()
-{
- static unsigned source_num = 0;
- ptx_in = NULL;
- if ( g_filename )
- ptx_in = fopen( g_filename, "r" );
- gpgpu_ptx_sim_init_memory();
- if (ptx_in) {
- symbol_table *symtab=init_parser(g_filename);
- ptx_parse();
- ptxinfo_in = open_ptxinfo(g_filename);
- ptxinfo_parse();
- load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF);
- load_constants(symtab,STATIC_ALLOC_LIMIT);
- } else {
- if (!g_override_embedded_ptx) {
- printf("GPGPU-Sim PTX: USING EMBEDDED .ptx files...\n");
- unsigned selected_capability = 0;
-
- if(g_ptx_force_max_capability == 0) {
- // No forced max capability, selected the highest capability
- //assert(g_ptx_max_capability > 0);
- selected_capability = g_ptx_max_capability;
- printf("GPGPU-Sim PTX: Loading PTX, selected capability: compute_%u\n", selected_capability);
- } else {
- // Forced max capability, select the highest capability less than or equal to forced capability
- ptx_info_t *pti;
- for( pti=g_ptx_source_array; pti!=NULL; pti=pti->next ){
- if(selected_capability < pti->capability && pti->capability <= g_ptx_force_max_capability)
- selected_capability = pti->capability;
- }
- //assert(selected_capability > 0);
- printf("GPGPU-Sim PTX: Loading PTX, max forced capability: compute_%u, selected capability: compute_%u\n",
- g_ptx_force_max_capability, selected_capability);
- }
-
- ptx_info_t *s;
- for ( s=g_ptx_source_array; s!=NULL; s=s->next ) {
- if(s->capability != selected_capability)
- continue;
-
- printf("GPGPU-Sim PTX: Loading PTX for %s, capability = compute_%u\n", s->fname, s->capability);
-
- symbol_table *symtab;
- source_num++;
- if(g_ptx_convert_to_ptxplus) {
- char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_to_ptxplus(s->str, s->cubin_str, source_num);
- symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxplus_str, s->str, source_num);
- delete[] ptxplus_str;
- } else {
- symtab=gpgpu_ptx_sim_load_ptx_from_string(s->str, s->str, source_num);
- }
- load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF);
- load_constants(symtab,STATIC_ALLOC_LIMIT);
- }
- } else {
- if(g_ptx_convert_to_ptxplus) {
- perror("GPGPU-Sim PTX: convert_to_ptxplus option enabled. Cannot use this option with external ptx files.\n");
- assert(0);
- }
- const char *filename = NULL;
- struct dirent **namelist;
- int n = scandir(".", &namelist, ptx_file_filter, alphasort);
- if (n < 0)
- perror("GPGPU-Sim PTX: no PTX files returned by scandir");
- else {
- while (n--) {
- filename = namelist[n]->d_name;
- printf("Parsing %s..\n", filename);
- ptx_in = fopen( filename, "r" );
- symbol_table *symtab=init_parser(filename);
- ptx_parse ();
- ptxinfo_in = open_ptxinfo(filename);
- ptxinfo_parse();
- load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF);
- load_constants(symtab,STATIC_ALLOC_LIMIT);
-
- free(namelist[n]);
- }
- free(namelist);
- }
- }
- }
-
- if ( ptx_in == NULL && g_override_embedded_ptx ) {
- printf("GPGPU-Sim PTX Simulator error: Could find/open .ptx file for reading\n");
- printf(" This means there are no .ptx files in the current directory.\n");
- printf(" Either place a .ptx file in the current directory, or ensure\n" );
- printf(" the PTX_SIM_KERNELFILE environment variable points to .ptx file.\n");
- printf(" PTX_SIM_KERNELFILE=\"%s\"\n", g_filename );
- exit(1);
- }
-
- if ( g_error_detected ) {
- printf( "GPGPU-Sim PTX: PTX parsing errors detected -- exiting.\n" );
- exit(1);
- }
- printf( "GPGPU-Sim PTX: Program parsing completed\n" );
-
- if ( g_kernel_name_to_function_lookup ) {
- for ( std::map<std::string,function_info*>::iterator f=g_kernel_name_to_function_lookup->begin();
- f != g_kernel_name_to_function_lookup->end(); f++ ) {
- gpgpu_ptx_assemble(f->first,f->second);
- }
- }
-}
-
-const struct gpgpu_ptx_sim_kernel_info * get_kernel_info(const char *kernel_key)
-{
- std::string kname;
- function_info *finfo = get_kernel(kernel_key,kname);
- return finfo->get_kernel_info();
-}
-
-kernel_info_t gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, gpgpu_ptx_sim_arg_list_t args,
- struct dim3 gridDim, struct dim3 blockDim )
-{
- std::string kname;
- function_info *entry = get_kernel(kernel_key,kname);
+ function_info *entry = context->get_kernel(hostFun);
+ std::string kname = entry->get_name();
printf("GPGPU-Sim PTX: Launching kernel \'%s\' gridDim= (%u,%u,%u) blockDim = (%u,%u,%u); ntuid=%u\n",
kname.c_str(), gridDim.x,gridDim.y,gridDim.z,blockDim.x,blockDim.y,blockDim.z,
g_ptx_thread_info_uid_next );
-
unsigned argcount=args.size();
unsigned argn=1;
for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) {
@@ -1758,7 +1615,7 @@ kernel_info_t gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, gpgpu_ptx_si
argn++;
}
- entry->finalize(g_param_mem);
+ entry->finalize(context->get_device()->get_gpgpu()->get_param_memory());
g_ptx_kernel_count++;
fflush(stdout);
diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc
index c4b9c2c..b331f92 100644
--- a/libopencl/opencl_runtime_api.cc
+++ b/libopencl/opencl_runtime_api.cc
@@ -148,7 +148,7 @@ private:
};
struct _cl_mem {
- _cl_mem( cl_mem_flags flags, size_t size , void *host_ptr, cl_int *errcode_ret );
+ _cl_mem( cl_mem_flags flags, size_t size , void *host_ptr, cl_int *errcode_ret, cl_device_id gpu );
cl_mem device_ptr();
void* host_ptr();
bool is_on_host() { return m_is_on_host; }
@@ -295,7 +295,8 @@ _cl_mem::_cl_mem(
cl_mem_flags flags,
size_t size ,
void * host_ptr,
- cl_int * errcode_ret )
+ cl_int * errcode_ret,
+ cl_device_id gpu )
{
if( errcode_ret )
*errcode_ret = CL_SUCCESS;
@@ -305,7 +306,6 @@ _cl_mem::_cl_mem(
m_size = size;
m_host_ptr = host_ptr;
m_device_ptr = 0;
- gpgpu_ptx_sim_init_memory();
if( (flags & (CL_MEM_USE_HOST_PTR|CL_MEM_COPY_HOST_PTR)) && host_ptr == NULL ) {
if( errcode_ret != NULL )
@@ -330,9 +330,9 @@ _cl_mem::_cl_mem(
}
if( !(flags & (CL_MEM_USE_HOST_PTR|CL_MEM_ALLOC_HOST_PTR)) ) {
// if not allocating on host, then allocate GPU memory and make a copy
- m_device_ptr = (size_t) gpgpu_ptx_sim_malloc(size);
+ m_device_ptr = (size_t) gpu->the_device()->gpgpu_ptx_sim_malloc(size);
if( host_ptr )
- gpgpu_ptx_sim_memcpy_to_gpu( m_device_ptr, host_ptr, size );
+ gpu->the_device()->gpgpu_ptx_sim_memcpy_to_gpu( m_device_ptr, host_ptr, size );
}
}
@@ -356,7 +356,7 @@ cl_mem _cl_context::CreateBuffer(
if( host_ptr && (m_hostptr_to_cl_mem.find(host_ptr) != m_hostptr_to_cl_mem.end()) ) {
printf("GPGPU-Sim OpenCL API: WARNING ** clCreateBuffer - buffer already created for this host variable\n");
}
- cl_mem result = new _cl_mem(flags,size,host_ptr,errcode_ret);
+ cl_mem result = new _cl_mem(flags,size,host_ptr,errcode_ret,m_gpu);
m_devptr_to_cl_mem[result->device_ptr()] = result;
if( host_ptr )
m_hostptr_to_cl_mem[host_ptr] = result;
@@ -516,7 +516,8 @@ void _cl_program::Build(const char *options)
}
}
info.m_asm = tmp;
- info.m_symtab = gpgpu_ptx_sim_load_ptx_from_string( tmp, tmp, source_num );
+ info.m_symtab = gpgpu_ptx_sim_load_ptx_from_string( tmp, source_num );
+ gpgpu_ptxinfo_load_from_string( tmp, source_num );
free(tmp);
}
printf("GPGPU-Sim OpenCL API: finished compiling OpenCL kernels.\n");
@@ -828,16 +829,19 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
if ( err_val != CL_SUCCESS )
return err_val;
- gpgpu_ptx_sim_memcpy_symbol( "%_global_size", _global_size, 3 * sizeof(int), 0, 1 );
- gpgpu_ptx_sim_memcpy_symbol( "%_work_dim", &work_dim, 1 * sizeof(int), 0, 1 );
- gpgpu_ptx_sim_memcpy_symbol( "%_global_num_groups", &GridDim, 3 * sizeof(int), 0, 1 );
- gpgpu_ptx_sim_memcpy_symbol( "%_global_launch_offset", zeros, 3 * sizeof(int), 0, 1 );
- gpgpu_ptx_sim_memcpy_symbol( "%_global_block_offset", zeros, 3 * sizeof(int), 0, 1 );
+ gpgpu_t *gpu = command_queue->get_device()->the_device();
+ gpgpu_ptx_sim_memcpy_symbol( "%_global_size", _global_size, 3 * sizeof(int), 0, 1, gpu );
+ gpgpu_ptx_sim_memcpy_symbol( "%_work_dim", &work_dim, 1 * sizeof(int), 0, 1, gpu );
+ gpgpu_ptx_sim_memcpy_symbol( "%_global_num_groups", &GridDim, 3 * sizeof(int), 0, 1, gpu );
+ gpgpu_ptx_sim_memcpy_symbol( "%_global_launch_offset", zeros, 3 * sizeof(int), 0, 1, gpu );
+ gpgpu_ptx_sim_memcpy_symbol( "%_global_block_offset", zeros, 3 * sizeof(int), 0, 1, gpu );
+
+ kernel_info_t grid = gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu);
if ( g_ptx_sim_mode )
- gpgpu_opencl_ptx_sim_main_func( kernel->get_implementation(), GridDim, BlockDim, params );
+ gpgpu_opencl_ptx_sim_main_func( grid, GridDim, BlockDim, params );
else
- gpgpu_opencl_ptx_sim_main_perf( kernel->get_implementation(), GridDim, BlockDim, params );
+ gpgpu_opencl_ptx_sim_main_perf( grid, GridDim, BlockDim, params );
return CL_SUCCESS;
}
@@ -854,7 +858,8 @@ clEnqueueReadBuffer(cl_command_queue command_queue,
{
if( !blocking_read )
gpgpusim_opencl_warning(__my_func__,__LINE__, "non-blocking read treated as blocking read");
- gpgpu_ptx_sim_memcpy_from_gpu( ptr, (size_t)buffer, cb );
+ gpgpu_t *gpu = command_queue->get_device()->the_device();
+ gpu->gpgpu_ptx_sim_memcpy_from_gpu( ptr, (size_t)buffer, cb );
return CL_SUCCESS;
}
@@ -871,7 +876,8 @@ clEnqueueWriteBuffer(cl_command_queue command_queue,
{
if( !blocking_write )
gpgpusim_opencl_warning(__my_func__,__LINE__, "non-blocking write treated as blocking write");
- gpgpu_ptx_sim_memcpy_to_gpu( (size_t)buffer, ptr, cb );
+ gpgpu_t *gpu = command_queue->get_device()->the_device();
+ gpu->gpgpu_ptx_sim_memcpy_to_gpu( (size_t)buffer, ptr, cb );
return CL_SUCCESS;
}
@@ -1143,12 +1149,13 @@ clEnqueueCopyBuffer(cl_command_queue command_queue,
if( src == NULL || dst == NULL )
return CL_INVALID_MEM_OBJECT;
+ gpgpu_t *gpu = command_queue->get_device()->the_device();
if( src->is_on_host() && !dst->is_on_host() )
- gpgpu_ptx_sim_memcpy_to_gpu( ((size_t)dst->device_ptr())+dst_offset, ((char*)src->host_ptr())+src_offset, cb );
+ gpu->gpgpu_ptx_sim_memcpy_to_gpu( ((size_t)dst->device_ptr())+dst_offset, ((char*)src->host_ptr())+src_offset, cb );
else if( !src->is_on_host() && dst->is_on_host() )
- gpgpu_ptx_sim_memcpy_from_gpu( ((char*)dst->host_ptr())+dst_offset, ((size_t)src->device_ptr())+src_offset, cb );
+ gpu->gpgpu_ptx_sim_memcpy_from_gpu( ((char*)dst->host_ptr())+dst_offset, ((size_t)src->device_ptr())+src_offset, cb );
else if( !src->is_on_host() && !dst->is_on_host() )
- gpgpu_ptx_sim_memcpy_gpu_to_gpu( ((size_t)dst->device_ptr())+dst_offset, ((size_t)src->device_ptr())+src_offset, cb );
+ gpu->gpgpu_ptx_sim_memcpy_gpu_to_gpu( ((size_t)dst->device_ptr())+dst_offset, ((size_t)src->device_ptr())+src_offset, cb );
else
opencl_not_implemented(__my_func__,__LINE__);
return CL_SUCCESS;
diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc
index f5149eb..a0e21f7 100644
--- a/src/abstract_hardware_model.cc
+++ b/src/abstract_hardware_model.cc
@@ -1,4 +1,5 @@
#include "abstract_hardware_model.h"
+#include "cuda-sim/memory.h"
void move_warp( warp_inst_t *&dst, warp_inst_t *&src )
{
@@ -9,3 +10,12 @@ void move_warp( warp_inst_t *&dst, warp_inst_t *&src )
src->clear();
}
+gpgpu_t::gpgpu_t()
+{
+ g_global_mem = new memory_space_impl<8192>("global",64*1024);
+ g_param_mem = new memory_space_impl<8192>("param",64*1024);
+ g_tex_mem = new memory_space_impl<8192>("tex",64*1024);
+ g_surf_mem = new memory_space_impl<8192>("surf",64*1024);
+
+ g_dev_malloc=GLOBAL_HEAP_START;
+}
diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h
index 15a7296..c90d56e 100644
--- a/src/abstract_hardware_model.h
+++ b/src/abstract_hardware_model.h
@@ -139,6 +139,44 @@ public:
virtual class gpgpu_sim *get_gpu() = 0;
};
+#define GLOBAL_HEAP_START 0x80000000
+ // start allocating from this address (lower values used for allocating globals in .ptx file)
+#define SHARED_MEM_SIZE_MAX (64*1024)
+#define LOCAL_MEM_SIZE_MAX (16*1024)
+#define MAX_STREAMING_MULTIPROCESSORS 64
+#define MAX_THREAD_PER_SM 1024
+#define TOTAL_LOCAL_MEM_PER_SM (MAX_THREAD_PER_SM*LOCAL_MEM_SIZE_MAX)
+#define TOTAL_SHARED_MEM (MAX_STREAMING_MULTIPROCESSORS*SHARED_MEM_SIZE_MAX)
+#define TOTAL_LOCAL_MEM (MAX_STREAMING_MULTIPROCESSORS*MAX_THREAD_PER_SM*LOCAL_MEM_SIZE_MAX)
+#define SHARED_GENERIC_START (GLOBAL_HEAP_START-TOTAL_SHARED_MEM)
+#define LOCAL_GENERIC_START (SHARED_GENERIC_START-TOTAL_LOCAL_MEM)
+#define STATIC_ALLOC_LIMIT (GLOBAL_HEAP_START - (TOTAL_LOCAL_MEM+TOTAL_SHARED_MEM))
+
+class gpgpu_t {
+public:
+ gpgpu_t();
+ void* gpgpu_ptx_sim_malloc( size_t size );
+ void* gpgpu_ptx_sim_mallocarray( size_t count );
+ void gpgpu_ptx_sim_memcpy_to_gpu( size_t dst_start_addr, const void *src, size_t count );
+ void gpgpu_ptx_sim_memcpy_from_gpu( void *dst, size_t src_start_addr, size_t count );
+ void gpgpu_ptx_sim_memcpy_gpu_to_gpu( size_t dst, size_t src, size_t count );
+ void gpgpu_ptx_sim_memset( size_t dst_start_addr, int c, size_t count );
+
+ class memory_space *get_global_memory() { return g_global_mem; }
+ class memory_space *get_tex_memory() { return g_tex_mem; }
+ class memory_space *get_surf_memory() { return g_surf_mem; }
+ class memory_space *get_param_memory() { return g_param_mem; }
+
+protected:
+ // functional simulation state
+ class memory_space *g_global_mem;
+ class memory_space *g_tex_mem;
+ class memory_space *g_surf_mem;
+ class memory_space *g_param_mem;
+
+ unsigned long long g_dev_malloc;
+};
+
struct gpgpu_ptx_sim_kernel_info
{
// Holds properties of the kernel (Kernel's resource use).
@@ -306,6 +344,24 @@ public:
m_per_scalar_thread[lane_id].callback.instruction = inst;
m_per_scalar_thread[lane_id].callback.thread = thread;
}
+ void set_active( std::vector<unsigned> &active )
+ {
+ warp_active_mask.reset();
+ for( std::vector<unsigned>::iterator i=active.begin(); i!=active.end(); ++i ) {
+ unsigned t = *i;
+ assert( t < m_warp_size );
+ warp_active_mask.set(t);
+ }
+ if( m_isatomic ) {
+ for( unsigned i=0; i < m_warp_size; i++ ) {
+ if( !warp_active_mask.test(i) ) {
+ m_per_scalar_thread[i].callback.function = NULL;
+ m_per_scalar_thread[i].callback.instruction = NULL;
+ m_per_scalar_thread[i].callback.thread = NULL;
+ }
+ }
+ }
+ }
// accessors
virtual void print_insn(FILE *fp) const
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index dcf1fd4..7f4eebd 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -309,20 +309,6 @@ void function_info::ptx_assemble()
m_assembled = true;
}
-
-
-void gpgpu_ptx_sim_init_memory()
-{
- static bool initialized = false;
- if ( !initialized ) {
- g_global_mem = new memory_space_impl<8192>("global",64*1024);
- g_param_mem = new memory_space_impl<8192>("param",64*1024);
- g_tex_mem = new memory_space_impl<8192>("tex",64*1024);
- g_surf_mem = new memory_space_impl<8192>("surf",64*1024);
- initialized = true;
- }
-}
-
addr_t shared_to_generic( unsigned smid, addr_t addr )
{
assert( addr < SHARED_MEM_SIZE_MAX );
@@ -392,9 +378,7 @@ addr_t generic_to_global( addr_t addr )
}
-unsigned long long g_dev_malloc=GLOBAL_HEAP_START;
-
-void* gpgpu_ptx_sim_malloc( size_t size )
+void* gpgpu_t::gpgpu_ptx_sim_malloc( size_t size )
{
unsigned long long result = g_dev_malloc;
printf("GPGPU-Sim PTX: allocating %zu bytes on GPU starting at address 0x%Lx\n", size, g_dev_malloc );
@@ -404,7 +388,7 @@ void* gpgpu_ptx_sim_malloc( size_t size )
return(void*) result;
}
-void* gpgpu_ptx_sim_mallocarray( size_t size )
+void* gpgpu_t::gpgpu_ptx_sim_mallocarray( size_t size )
{
unsigned long long result = g_dev_malloc;
printf("GPGPU-Sim PTX: allocating %zu bytes on GPU starting at address 0x%Lx\n", size, g_dev_malloc );
@@ -415,7 +399,7 @@ void* gpgpu_ptx_sim_mallocarray( size_t size )
}
-void gpgpu_ptx_sim_memcpy_to_gpu( size_t dst_start_addr, const void *src, size_t count )
+void gpgpu_t::gpgpu_ptx_sim_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);
@@ -426,7 +410,7 @@ void gpgpu_ptx_sim_memcpy_to_gpu( size_t dst_start_addr, const void *src, size_t
fflush(stdout);
}
-void gpgpu_ptx_sim_memcpy_from_gpu( void *dst, size_t src_start_addr, size_t count )
+void gpgpu_t::gpgpu_ptx_sim_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);
@@ -437,7 +421,7 @@ void gpgpu_ptx_sim_memcpy_from_gpu( void *dst, size_t src_start_addr, size_t cou
fflush(stdout);
}
-void gpgpu_ptx_sim_memcpy_gpu_to_gpu( size_t dst, size_t src, size_t count )
+void gpgpu_t::gpgpu_ptx_sim_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,
(unsigned long long) src, (unsigned long long) dst );
@@ -451,7 +435,7 @@ void gpgpu_ptx_sim_memcpy_gpu_to_gpu( size_t dst, size_t src, size_t count )
fflush(stdout);
}
-void gpgpu_ptx_sim_memset( size_t dst_start_addr, int c, size_t count )
+void gpgpu_t::gpgpu_ptx_sim_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... ",
count, (unsigned char) c, (unsigned long long) dst_start_addr );
@@ -890,6 +874,12 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id )
}
}
if( !skip ) {
+ ptx_instruction *pJ = NULL;
+ if( pI->get_opcode() == VOTE_OP ) {
+ pJ = new ptx_instruction(*pI);
+ *((warp_inst_t*)pJ) = inst;
+ pI = pJ;
+ }
switch ( pI->get_opcode() ) {
#define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,this); op_classification = CLASSIFICATION; break;
#include "opcodes.def"
@@ -899,6 +889,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id )
printf( "Execution error: Invalid opcode (0x%x)\n", pI->get_opcode() );
break;
}
+ delete pJ;
// Run exit instruction if exit option included
if(pI->is_exit())
@@ -1045,11 +1036,6 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id )
}
}
-std::list<ptx_thread_info *> g_active_threads;
-std::map<unsigned,memory_space*> g_shared_memory_lookup;
-std::map<unsigned,ptx_cta_info*> g_ptx_cta_lookup;
-std::map<unsigned,std::map<unsigned,memory_space*> > g_local_memory_lookup;
-
void set_param_gpgpu_num_shaders(int num_shaders)
{
gpgpu_param_num_shaders = num_shaders;
@@ -1073,8 +1059,14 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel,
unsigned num_threads,
core_t *core,
unsigned hw_cta_id,
- unsigned hw_warp_id )
+ unsigned hw_warp_id,
+ gpgpu_t *gpu )
{
+ static std::list<ptx_thread_info *> active_threads;
+ static std::map<unsigned,memory_space*> shared_memory_lookup;
+ static std::map<unsigned,ptx_cta_info*> ptx_cta_lookup;
+ static std::map<unsigned,std::map<unsigned,memory_space*> > local_memory_lookup;
+
if ( *thread_info != NULL ) {
ptx_thread_info *thd = *thread_info;
assert( thd->is_done() );
@@ -1091,16 +1083,12 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel,
*thread_info = NULL;
}
- if ( !g_active_threads.empty() ) { //if g_active_threads not empty...
- assert( g_active_threads.size() <= threads_left );
- ptx_thread_info *thd = g_active_threads.front();
- g_active_threads.pop_front();
+ if ( !active_threads.empty() ) { //if g_active_threads not empty...
+ assert( active_threads.size() <= threads_left );
+ ptx_thread_info *thd = active_threads.front();
+ active_threads.pop_front();
*thread_info = thd;
- thd->set_hw_tid(tid);
- thd->set_hw_wid(hw_warp_id);
- thd->set_hw_ctaid(hw_cta_id);
- thd->set_core(core);
- thd->set_hw_sid(sid);
+ thd->init(gpu, core, sid, hw_cta_id, hw_warp_id, tid );
return 1;
}
@@ -1127,7 +1115,7 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel,
unsigned sm_idx = (tid/cta_size)*gpgpu_param_num_shaders + sid;
- if ( g_shared_memory_lookup.find(sm_idx) == g_shared_memory_lookup.end() ) {
+ if ( shared_memory_lookup.find(sm_idx) == shared_memory_lookup.end() ) {
if ( g_debug_execution >= 1 ) {
printf(" <CTA alloc> : sm_idx=%u sid=%u max_cta_per_sm=%u\n",
sm_idx, sid, max_cta_per_sm );
@@ -1135,20 +1123,20 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel,
char buf[512];
snprintf(buf,512,"shared_%u", sid);
shared_mem = new memory_space_impl<16*1024>(buf,4);
- g_shared_memory_lookup[sm_idx] = shared_mem;
+ shared_memory_lookup[sm_idx] = shared_mem;
cta_info = new ptx_cta_info(sm_idx);
- g_ptx_cta_lookup[sm_idx] = cta_info;
+ ptx_cta_lookup[sm_idx] = cta_info;
} else {
if ( g_debug_execution >= 1 ) {
printf(" <CTA realloc> : sm_idx=%u sid=%u max_cta_per_sm=%u\n",
sm_idx, sid, max_cta_per_sm );
}
- shared_mem = g_shared_memory_lookup[sm_idx];
- cta_info = g_ptx_cta_lookup[sm_idx];
+ shared_mem = shared_memory_lookup[sm_idx];
+ cta_info = ptx_cta_lookup[sm_idx];
cta_info->check_cta_thread_status_and_reset();
}
- std::map<unsigned,memory_space*> &local_mem_lookup = g_local_memory_lookup[sid];
+ std::map<unsigned,memory_space*> &local_mem_lookup = local_memory_lookup[sid];
while( kernel.more_threads_in_cta() ) {
dim3 ctaid3d = kernel.get_next_cta_id();
unsigned new_tid = kernel.get_next_thread_id();
@@ -1174,11 +1162,6 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel,
thd->set_tid(tid3d);
if( kernel.entry()->get_ptx_version().extensions() )
thd->cpy_tid_to_reg(tid3d);
- thd->set_hw_tid((unsigned)-1);
- thd->set_hw_wid((unsigned)-1);
- thd->set_hw_ctaid((unsigned)-1);
- thd->set_core(NULL);
- thd->set_hw_sid((unsigned)-1);
thd->set_valid();
thd->m_shared_mem = shared_mem;
function_info *finfo = thd->func_info();
@@ -1192,7 +1175,7 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel,
ctaid3d.x,ctaid3d.y,ctaid3d.z,tid3d.x,tid3d.y,tid3d.z, (unsigned long long)thd );
fflush(stdout);
}
- g_active_threads.push_back(thd);
+ active_threads.push_back(thd);
}
if ( g_debug_execution==-1 ) {
printf("GPGPU-Sim PTX simulator: <-- FINISHING THREAD ALLOCATION\n");
@@ -1201,16 +1184,10 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel,
kernel.increment_cta_id();
- assert( g_active_threads.size() <= threads_left );
-
- *thread_info = g_active_threads.front();
- (*thread_info)->set_hw_tid(tid);
- (*thread_info)->set_hw_wid(hw_warp_id);
- (*thread_info)->set_hw_ctaid(hw_cta_id);
- (*thread_info)->set_core(core);
- (*thread_info)->set_hw_sid(sid);
- g_active_threads.pop_front();
-
+ assert( active_threads.size() <= threads_left );
+ *thread_info = active_threads.front();
+ (*thread_info)->init(gpu, core, sid, hw_cta_id, hw_warp_id, tid );
+ active_threads.pop_front();
return 1;
}
@@ -1220,7 +1197,11 @@ size_t get_kernel_code_size( class function_info *entry )
}
-kernel_info_t gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,gpgpu_ptx_sim_arg_list_t args, struct dim3 gridDim, struct dim3 blockDim )
+kernel_info_t gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,
+ gpgpu_ptx_sim_arg_list_t args,
+ struct dim3 gridDim,
+ struct dim3 blockDim,
+ gpgpu_t *gpu )
{
unsigned argcount=args.size();
unsigned argn=1;
@@ -1228,7 +1209,7 @@ kernel_info_t gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,gpgpu_pt
entry->add_param_data(argcount-argn,&(*a));
argn++;
}
- entry->finalize(g_param_mem);
+ entry->finalize(gpu->get_param_memory());
g_ptx_kernel_count++;
fflush(stdout);
@@ -1263,7 +1244,7 @@ void gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceNam
g_global_name_lookup[hostVar] = deviceName;
}
-void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to )
+void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu )
{
printf("GPGPU-Sim PTX: starting gpgpu_ptx_sim_memcpy_symbol with hostVar 0x%p\n", hostVar);
bool found_sym = false;
@@ -1314,11 +1295,11 @@ void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t co
unsigned dst = sym->get_address() + offset;
switch (mem_region.get_type()) {
case const_space:
- mem = g_global_mem;
+ mem = gpu->get_global_memory();
mem_name = "global";
break;
case global_space:
- mem = g_global_mem;
+ mem = gpu->get_global_memory();
mem_name = "global";
break;
default:
diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h
index 322e8a2..b72a157 100644
--- a/src/cuda-sim/cuda-sim.h
+++ b/src/cuda-sim/cuda-sim.h
@@ -13,7 +13,6 @@ class symbol_table;
extern const char *g_gpgpusim_version_string;
extern int g_ptx_sim_mode;
-extern memory_space *g_global_mem;
extern int g_debug_execution;
extern int g_debug_thread_uid;
extern void ** g_inst_classification_stat;
@@ -25,20 +24,13 @@ extern FILE* ptx_inst_debug_file;
extern class kernel_info_t gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,
gpgpu_ptx_sim_arg_list_t args,
struct dim3 gridDim,
- struct dim3 blockDim );
+ struct dim3 blockDim,
+ class gpgpu_t *gpu );
extern void gpgpu_cuda_ptx_sim_main_func( kernel_info_t kernel, dim3 gridDim, dim3 blockDim, gpgpu_ptx_sim_arg_list_t args);
extern void print_splash();
-extern void* gpgpu_ptx_sim_malloc( size_t count );
-extern void* gpgpu_ptx_sim_mallocarray( size_t count );
-extern void gpgpu_ptx_sim_memcpy_to_gpu( size_t dst_start_addr, const void *src, size_t count );
-extern void gpgpu_ptx_sim_memcpy_from_gpu( void *dst, size_t src_start_addr, size_t count );
-extern void gpgpu_ptx_sim_memcpy_gpu_to_gpu( size_t dst, size_t src, size_t count );
-extern void gpgpu_ptx_sim_memset( size_t dst_start_addr, int c, size_t count );
-extern void gpgpu_ptx_sim_init_memory();
-extern void gpgpu_ptx_sim_register_kernel(void **fatCubinHandle,const char *hostFun, const char *deviceFun);
extern void gpgpu_ptx_sim_register_const_variable(void*, const char *deviceName, size_t size );
extern void gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size );
-extern void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to );
+extern void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu );
extern void gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* texref, const struct cudaArray* array);
extern void gpgpu_ptx_sim_bindNameToTexture(const char* name, const struct textureReference* texref);
@@ -55,7 +47,8 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel,
unsigned num_threads,
class core_t *core,
unsigned hw_cta_id,
- unsigned hw_warp_id );
+ unsigned hw_warp_id,
+ gpgpu_t *gpu );
const warp_inst_t *ptx_fetch_inst( address_type pc );
const struct gpgpu_ptx_sim_kernel_info* ptx_sim_kernel_info(class function_info *kernel);
void ptx_print_insn( address_type pc, FILE *fp );
diff --git a/src/cuda-sim/cuda_device_printf.cc b/src/cuda-sim/cuda_device_printf.cc
index 9edeb5e..8483dd9 100644
--- a/src/cuda-sim/cuda_device_printf.cc
+++ b/src/cuda-sim/cuda_device_printf.cc
@@ -65,7 +65,7 @@
#include "cuda_device_printf.h"
#include "ptx_ir.h"
-void decode_space( memory_space_t &space, const ptx_thread_info *thread, const operand_info &op, memory_space *&mem, addr_t &addr);
+void decode_space( memory_space_t &space, ptx_thread_info *thread, const operand_info &op, memory_space *&mem, addr_t &addr);
void my_cuda_printf(const char *fmtstr,const char *arg_list)
{
@@ -104,7 +104,7 @@ void my_cuda_printf(const char *fmtstr,const char *arg_list)
}
}
-void gpgpusim_cuda_vprintf(const ptx_instruction * pI, const ptx_thread_info * thread, const function_info * target_func )
+void gpgpusim_cuda_vprintf(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func )
{
char *fmtstr = NULL;
char *arg_list = NULL;
diff --git a/src/cuda-sim/cuda_device_printf.h b/src/cuda-sim/cuda_device_printf.h
index 022b555..7b2a5ef 100644
--- a/src/cuda-sim/cuda_device_printf.h
+++ b/src/cuda-sim/cuda_device_printf.h
@@ -65,6 +65,6 @@
#ifndef CUDA_DEVICE_PRINTF_INCLUDED
#define CUDA_DEVICE_PRINTF_INCLUDED
-void gpgpusim_cuda_vprintf(const class ptx_instruction * pI, const class ptx_thread_info * thread, const class function_info * target_func );
+void gpgpusim_cuda_vprintf(const class ptx_instruction * pI, class ptx_thread_info * thread, const class function_info * target_func );
#endif
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index 23d6cc8..335ea59 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -249,7 +249,7 @@ ptx_reg_t ptx_thread_info::get_operand_value( const operand_info &op, operand_in
//complete other cases for reading from memory, such as reading from other const memory
if((op.get_addr_space() == 1)&&(derefFlag)) {
// global memory - g[4], g[$r0]
- mem = g_global_mem;
+ mem = thread->get_global_memory();
type_info_key::type_decode(opType,size,t);
mem->read(result.u32,size/8,&finalResult.u128);
thread->m_last_effective_address = result.u32;
@@ -269,7 +269,7 @@ ptx_reg_t ptx_thread_info::get_operand_value( const operand_info &op, operand_in
sign_extend(finalResult,size,dstInfo);
} else if((op.get_addr_space() == 3)&&(derefFlag)) {
// const memory - ce0c1[4], ce0c1[$r0]
- mem = g_global_mem;
+ mem = thread->get_global_memory();
type_info_key::type_decode(opType,size,t);
mem->read((result.u32 + op.get_const_mem_offset()),size/8,&finalResult.u128);
thread->m_last_effective_address = result.u32;
@@ -606,7 +606,7 @@ void ptx_thread_info::set_operand_value( const operand_info &dst, const ptx_reg_
else if(dst.get_addr_space() == 1)
{
dstData = thread->get_operand_value(dst, dst, type, thread, 0);
- mem = g_global_mem;
+ mem = thread->get_global_memory();
type_info_key::type_decode(type,size,t);
mem->write(dstData.u32,size/8,&data.u128,thread,pI);
@@ -863,7 +863,7 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread )
// Copy value pointed to in operand 'a' into register 'd'
// (i.e. copy src1_data to dst)
- g_global_mem->read(src1_data.u32,size/8,&data.s64);
+ thread->get_global_memory()->read(src1_data.u32,size/8,&data.s64);
thread->set_operand_value(dst, data, to_type, thread, pI); // Write value into register 'd'
// Get the atomic operation to be performed
@@ -1086,7 +1086,7 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread )
// Write operation result into global memory
// (i.e. copy src1_data to dst)
- g_global_mem->write(src1_data.u32,size/8,&op_result.s64,thread,pI);
+ thread->get_global_memory()->write(src1_data.u32,size/8,&op_result.s64,thread,pI);
gpgpu_sim *gpu = thread->get_gpu();
gpu->decrement_atomic_count(thread->get_hw_sid(),thread->get_hw_wid());
}
@@ -1998,7 +1998,7 @@ void isspacep_impl( const ptx_instruction *pI, ptx_thread_info *thread )
thread->set_reg(dst.get_symbol(),p);
}
-void decode_space( memory_space_t &space, const ptx_thread_info *thread, const operand_info &op, memory_space *&mem, addr_t &addr)
+void decode_space( memory_space_t &space, ptx_thread_info *thread, const operand_info &op, memory_space *&mem, addr_t &addr)
{
unsigned smid = thread->get_hw_sid();
unsigned hwtid = thread->get_hw_tid();
@@ -2018,23 +2018,23 @@ void decode_space( memory_space_t &space, const ptx_thread_info *thread, const o
}
}
switch ( space.get_type() ) {
- case global_space: mem = g_global_mem; break;
+ case global_space: mem = thread->get_global_memory(); break;
case param_space_local:
case local_space:
mem = thread->m_local_mem;
addr += thread->get_local_mem_stack_pointer();
break;
- case tex_space: mem = g_tex_mem; break;
- case surf_space: mem = g_surf_mem; break;
- case param_space_kernel: mem = g_param_mem; break;
+ case tex_space: mem = thread->get_tex_memory(); break;
+ case surf_space: mem = thread->get_surf_memory(); break;
+ case param_space_kernel: mem = thread->get_param_memory(); break;
case shared_space: mem = thread->m_shared_mem; break;
- case const_space: mem = g_global_mem; break;
+ case const_space: mem = thread->get_global_memory(); break;
case generic_space:
if( thread->get_ptx_version().ver() >= 2.0 ) {
// convert generic address to memory space address
space = whichspace(addr);
switch ( space.get_type() ) {
- case global_space: mem = g_global_mem; addr = generic_to_global(addr); break;
+ case global_space: mem = thread->get_global_memory(); addr = generic_to_global(addr); break;
case local_space: mem = thread->m_local_mem; addr = generic_to_local(smid,hwtid,addr); break;
case shared_space: mem = thread->m_shared_mem; addr = generic_to_shared(smid,addr); break;
default: abort();
@@ -3555,7 +3555,7 @@ void tex_impl( const ptx_instruction *pI, ptx_thread_info *thread )
//assume always 2D f32 input
//access array with src2 coordinates
- memory_space *mem = g_global_mem;
+ memory_space *mem = thread->get_global_memory();
float x_f32, y_f32;
size_t size;
int t;
@@ -3799,12 +3799,10 @@ void vote_impl( const ptx_instruction *pI, ptx_thread_info *thread )
threads_in_warp.clear();
and_all = true;
or_all = false;
- unsigned mask=0x80000000;
- unsigned offset=31;
- while( mask && !pI->active(mask) ) {
- mask = mask>>1;
+ int offset=31;
+ while( (offset>=0) && !pI->active(offset) )
offset--;
- }
+ assert( offset >= 0 );
last_tid = (thread->get_hw_tid() - (thread->get_hw_tid()%pI->warp_size())) + offset;
}
@@ -3883,7 +3881,7 @@ ptx_reg_t srcOperandModifiers(ptx_reg_t opData, operand_info opInfo, operand_inf
//complete other cases for reading from memory, such as reading from other const memory
if(opInfo.get_addr_space() == 1)
{
- mem = g_global_mem;
+ mem = thread->get_global_memory();
type_info_key::type_decode(type,size,t);
mem->read(opData.u32,size/8,&result.u64);
if( type == S16_TYPE || type == S32_TYPE )
@@ -3901,7 +3899,7 @@ ptx_reg_t srcOperandModifiers(ptx_reg_t opData, operand_info opInfo, operand_inf
}
else if(opInfo.get_addr_space() == 3)
{
- mem = g_global_mem;
+ mem = thread->get_global_memory();
type_info_key::type_decode(type,size,t);
mem->read((opData.u32 + opInfo.get_const_mem_offset()),size/8,&result.u64);
diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc
index b7ec3ac..8172f80 100644
--- a/src/cuda-sim/ptx_ir.cc
+++ b/src/cuda-sim/ptx_ir.cc
@@ -1162,7 +1162,7 @@ void ptx_instruction::print_insn( FILE *fp ) const
snprintf(buf,1024,"%s", m_source.c_str());
p = strtok(buf,";");
if( !is_label() )
- fprintf(fp," PC=%3u [%3u] ", m_PC, m_instr_mem_index );
+ fprintf(fp," PC=0x%03x ", m_PC );
else
fprintf(fp," " );
fprintf(fp,"(%s:%u) %s", m_source_file.c_str(), m_source_line, p );
diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h
index 6010caf..612326c 100644
--- a/src/cuda-sim/ptx_ir.h
+++ b/src/cuda-sim/ptx_ir.h
@@ -803,7 +803,6 @@ public:
const char *source,
unsigned warp_size );
-
void print_insn() const;
virtual void print_insn( FILE *fp ) const;
unsigned inst_size() const { return m_inst_size; }
@@ -1404,20 +1403,6 @@ struct textureInfo {
extern std::map<std::string,symbol_table*> g_sym_name_to_symbol_table;
-#define GLOBAL_HEAP_START 0x80000000
- // start allocating from this address (lower values used for allocating globals in .ptx file)
-#define SHARED_MEM_SIZE_MAX (64*1024)
-#define LOCAL_MEM_SIZE_MAX (16*1024)
-#define MAX_STREAMING_MULTIPROCESSORS 64
-#define MAX_THREAD_PER_SM 1024
-#define TOTAL_LOCAL_MEM_PER_SM (MAX_THREAD_PER_SM*LOCAL_MEM_SIZE_MAX)
-#define TOTAL_SHARED_MEM (MAX_STREAMING_MULTIPROCESSORS*SHARED_MEM_SIZE_MAX)
-#define TOTAL_LOCAL_MEM (MAX_STREAMING_MULTIPROCESSORS*MAX_THREAD_PER_SM*LOCAL_MEM_SIZE_MAX)
-#define SHARED_GENERIC_START (GLOBAL_HEAP_START-TOTAL_SHARED_MEM)
-#define LOCAL_GENERIC_START (SHARED_GENERIC_START-TOTAL_LOCAL_MEM)
-#define STATIC_ALLOC_LIMIT (GLOBAL_HEAP_START - (TOTAL_LOCAL_MEM+TOTAL_SHARED_MEM))
-
-
extern bool g_keep_intermediate_files;
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index d4f3822..2387528 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -87,9 +87,6 @@ extern "C" int ptxinfo_parse();
extern "C" int ptxinfo_debug;
extern "C" FILE *ptxinfo_in;
-extern int g_ptx_save_converted_ptxplus;
-
-
static bool g_save_embedded_ptx;
bool g_keep_intermediate_files;
@@ -129,7 +126,7 @@ void print_ptx_file( const char *p, unsigned source_num, const char *filename )
fflush(stdout);
}
-char* gpgpu_ptx_sim_convert_ptx_to_ptxplus(const char *ptx_str, const char *cubin_str, unsigned source_num)
+char* gpgpu_ptx_sim_convert_ptx_to_ptxplus(const char *ptx_str, const char *cubin_str, unsigned source_num, bool save_converted )
{
printf("GPGPU-Sim PTX: converting EMBEDDED .ptx file to ptxplus \n");
@@ -201,7 +198,7 @@ char* gpgpu_ptx_sim_convert_ptx_to_ptxplus(const char *ptx_str, const char *cubi
strcpy(ptxplus_str, text.c_str());
// Save ptxplus to file if specified
- if(g_ptx_save_converted_ptxplus) {
+ if(save_converted) {
char fname_ptxplus_save[1024];
snprintf(fname_ptxplus_save,1024,"_%u.ptxplus", source_num );
printf("GPGPU-Sim PTX: saving converted ptxplus to file \"%s\"\n", fname_ptxplus_save);
@@ -224,11 +221,10 @@ char* gpgpu_ptx_sim_convert_ptx_to_ptxplus(const char *ptx_str, const char *cubi
printf("GPGPU-Sim PTX: DONE converting EMBEDDED .ptx file to ptxplus \n");
return ptxplus_str;
-
}
-symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, const char *p_for_info, unsigned source_num )
+symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num )
{
char buf[1024];
snprintf(buf,1024,"_%u.ptx", source_num );
@@ -242,7 +238,7 @@ symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, const char *p_f
int errors = ptx_parse ();
if ( errors ) {
char fname[1024];
- snprintf(fname,1024,"_ptx_XXXXXX");
+ snprintf(fname,1024,"_ptx_errors_XXXXXX");
int fd=mkstemp(fname);
close(fd);
printf("GPGPU-Sim PTX: parser error detected, exiting... but first extracting .ptx to \"%s\"\n", fname);
@@ -257,7 +253,11 @@ symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, const char *p_f
print_ptx_file(p,source_num,buf);
printf("GPGPU-Sim PTX: finished parsing EMBEDDED .ptx file %s\n",buf);
+ return symtab;
+}
+void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num )
+{
char fname[1024];
snprintf(fname,1024,"_ptx_XXXXXX");
int fd=mkstemp(fname);
@@ -287,9 +287,11 @@ symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, const char *p_f
char commandline[1024];
char extra_flags[1024];
extra_flags[0]=0;
+
#if CUDART_VERSION >= 3000
snprintf(extra_flags,1024,"--gpu-name=sm_20");
#endif
+
snprintf(commandline,1024,"ptxas %s -v %s --output-file /dev/null 2> %s",
extra_flags, fname2, tempfile_ptxinfo);
printf("GPGPU-Sim PTX: generating ptxinfo using \"%s\"\n", commandline);
@@ -310,6 +312,5 @@ symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, const char *p_f
printf("GPGPU-Sim PTX: ERROR ** while loading PTX (c) %d\n", result);
exit(1);
}
- return symtab;
}
diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h
index e9efd9b..da0d10e 100644
--- a/src/cuda-sim/ptx_loader.h
+++ b/src/cuda-sim/ptx_loader.h
@@ -66,15 +66,10 @@
#ifndef PTX_LOADER_H_INCLUDED
#define PTX_LOADER_H_INCLUDED
-class memory_space;
-
-extern memory_space *g_global_mem;
-extern memory_space *g_tex_mem;
-extern memory_space *g_surf_mem;
-extern memory_space *g_param_mem;
extern bool g_override_embedded_ptx;
-class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, const char *p_for_info, unsigned source_num );
-char* gpgpu_ptx_sim_convert_ptx_to_ptxplus(const char *ptx_str, const char *cubin_str, unsigned source_num);
+class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num );
+void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num );
+char* gpgpu_ptx_sim_convert_ptx_to_ptxplus(const char *ptx_str, const char *cubin_str, unsigned source_num, bool save_converted );
#endif
diff --git a/src/cuda-sim/ptx_sim.cc b/src/cuda-sim/ptx_sim.cc
index 2cd0ba6..59f4522 100644
--- a/src/cuda-sim/ptx_sim.cc
+++ b/src/cuda-sim/ptx_sim.cc
@@ -245,6 +245,7 @@ ptx_thread_info::ptx_thread_info()
m_last_was_call = false;
m_enable_debug_trace = false;
m_local_mem_stack_pointer = 0;
+ m_gpu = NULL;
}
const ptx_version &ptx_thread_info::get_ptx_version() const
diff --git a/src/cuda-sim/ptx_sim.h b/src/cuda-sim/ptx_sim.h
index 80ccde3..298122e 100644
--- a/src/cuda-sim/ptx_sim.h
+++ b/src/cuda-sim/ptx_sim.h
@@ -276,6 +276,16 @@ public:
~ptx_thread_info();
ptx_thread_info();
+ void init(gpgpu_t *gpu, core_t *core, unsigned sid, unsigned cta_id, unsigned wid, unsigned tid )
+ {
+ m_gpu = gpu;
+ m_core = core;
+ m_hw_sid=sid;
+ m_hw_ctaid=cta_id;
+ m_hw_wid=wid;
+ m_hw_tid=tid;
+ }
+
void ptx_fetch_inst( inst_t &inst ) const;
void ptx_exec_inst( warp_inst_t &inst, unsigned lane_id );
@@ -310,11 +320,6 @@ public:
unsigned get_hw_ctaid() const { return m_hw_ctaid;}
unsigned get_hw_wid() const { return m_hw_wid;}
unsigned get_hw_sid() const { return m_hw_sid;}
- void set_hw_tid(unsigned tid) { m_hw_tid=tid;}
- void set_hw_wid(unsigned wid) { m_hw_wid=wid;}
- void set_hw_sid(unsigned sid) { m_hw_sid=sid;}
- void set_hw_ctaid(unsigned cta_id) { m_hw_ctaid=cta_id;}
- void set_core(core_t *core) { m_core = core; }
core_t *get_core() { return m_core; }
unsigned get_icount() const { return m_icount;}
@@ -421,6 +426,11 @@ public:
void enable_debug_trace() { m_enable_debug_trace = true; }
unsigned get_local_mem_stack_pointer() const { return m_local_mem_stack_pointer; }
+ memory_space *get_global_memory() { return m_gpu->get_global_memory(); }
+ memory_space *get_tex_memory() { return m_gpu->get_tex_memory(); }
+ memory_space *get_surf_memory() { return m_gpu->get_surf_memory(); }
+ memory_space *get_param_memory() { return m_gpu->get_param_memory(); }
+
public:
addr_t m_last_effective_address;
bool m_branch_taken;
@@ -436,6 +446,7 @@ private:
unsigned m_uid;
core_t *m_core;
+ gpgpu_t *m_gpu;
bool m_valid;
dim3 m_ntid;
dim3 m_tid;
diff --git a/src/debug.cc b/src/debug.cc
index 7bc72ac..c1aa3c6 100644
--- a/src/debug.cc
+++ b/src/debug.cc
@@ -61,7 +61,8 @@ void gpgpu_sim::gpgpu_debug()
brk_pt &b=i->second;
if( b.is_watchpoint() ) {
unsigned addr = b.get_addr();
- unsigned new_value = read_location(addr);
+ unsigned new_value;
+ g_global_mem->read(addr,4,&new_value);
if( new_value != b.get_value() || g_watchpoint_hits.find(num) != g_watchpoint_hits.end() ) {
printf( "GPGPU-Sim PTX DBG: watch point %u triggered (old value=%x, new value=%x)\n",
num,b.get_value(),new_value );
@@ -155,7 +156,8 @@ void gpgpu_sim::gpgpu_debug()
tok = strtok(NULL," \t\n");
unsigned addr;
sscanf(tok,"%x",&addr);
- unsigned value = read_location(addr);
+ unsigned value;
+ g_global_mem->read(addr,4,&value);
g_global_mem->set_watch(addr,next_brkpt);
breakpoints[next_brkpt++] = brk_pt(addr,value);
} else if( !strcmp(tok,"l") ) {
@@ -199,9 +201,3 @@ bool thread_at_brkpt( ptx_thread_info *thread, const struct brk_pt &b )
return b.is_equal(thread->get_location(),thread->get_uid());
}
-unsigned read_location( addr_t addr )
-{
- unsigned result=0;
- g_global_mem->read(addr,4,&result);
- return result;
-}
diff --git a/src/debug.h b/src/debug.h
index 7c123a3..0884c99 100644
--- a/src/debug.h
+++ b/src/debug.h
@@ -61,7 +61,6 @@ extern int gpgpu_ptx_instruction_classification ;
class ptx_thread_info;
class ptx_instruction;
bool thread_at_brkpt( ptx_thread_info *thd_info, const struct brk_pt &b );
-unsigned read_location( addr_t addr );
void hit_watchpoint( unsigned watchpoint_num, ptx_thread_info *thd, const ptx_instruction *pI );
#endif
diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc
index 96d5c62..77c6aae 100644
--- a/src/gpgpu-sim/gpu-sim.cc
+++ b/src/gpgpu-sim/gpu-sim.cc
@@ -174,10 +174,6 @@ int g_ptx_inst_debug_to_file;
char* g_ptx_inst_debug_file;
int g_ptx_inst_debug_thread_uid;
-int g_ptx_convert_to_ptxplus;
-int g_ptx_save_converted_ptxplus;
-unsigned g_ptx_force_max_capability;
-
void visualizer_options(option_parser_t opp);
void gpgpu_sim::reg_options(option_parser_t opp)
@@ -371,15 +367,15 @@ void gpgpu_sim::reg_options(option_parser_t opp)
"Thread UID for executed instructions' debug output",
"1");
option_parser_register(opp, "-gpgpu_ptx_convert_to_ptxplus", OPT_BOOL,
- &g_ptx_convert_to_ptxplus,
+ &m_ptx_convert_to_ptxplus,
"Convert embedded ptx to ptxplus",
"0");
option_parser_register(opp, "-gpgpu_ptx_save_converted_ptxplus", OPT_BOOL,
- &g_ptx_save_converted_ptxplus,
+ &m_ptx_save_converted_ptxplus,
"Saved converted ptxplus to a file",
"0");
option_parser_register(opp, "-gpgpu_ptx_force_max_capability", OPT_UINT32,
- &g_ptx_force_max_capability,
+ &m_ptx_force_max_capability,
"Force maximum compute capability",
"0");
option_parser_register(opp, "-gpgpu_operand_collector", OPT_BOOL, &m_shader_config->gpgpu_operand_collector,
@@ -1080,7 +1076,7 @@ void shader_core_ctx::issue_block2core( kernel_info_t &kernel )
for (unsigned i = start_thread; i<end_thread; i++) {
m_thread[i].m_cta_id = free_cta_hw_id;
unsigned warp_id = i/m_config->warp_size;
- nthreads_in_block += ptx_sim_init_thread(kernel,&m_thread[i].m_functional_model_thread_state,m_sid,i,cta_size-(i-start_thread),m_config->n_thread_per_shader,this,free_cta_hw_id,warp_id);
+ nthreads_in_block += ptx_sim_init_thread(kernel,&m_thread[i].m_functional_model_thread_state,m_sid,i,cta_size-(i-start_thread),m_config->n_thread_per_shader,this,free_cta_hw_id,warp_id,m_gpu);
warps.set( warp_id );
}
assert( nthreads_in_block > 0 && nthreads_in_block <= m_config->n_thread_per_shader); // should be at least one, but less than max
diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h
index ac03cbc..d690fab 100644
--- a/src/gpgpu-sim/gpu-sim.h
+++ b/src/gpgpu-sim/gpu-sim.h
@@ -178,7 +178,7 @@ extern int g_ptx_inst_debug_thread_uid;
-class gpgpu_sim {
+class gpgpu_sim : public gpgpu_t {
public:
gpgpu_sim();
@@ -211,6 +211,10 @@ public:
void gpu_print_stat() const;
void dump_pipeline( int mask, int s, int m ) const;
+ unsigned get_forced_capability() const { return m_ptx_force_max_capability; }
+ bool convert_to_ptxplus() const { return m_ptx_convert_to_ptxplus; }
+ bool saved_converted_ptxplus() const { return m_ptx_save_converted_ptxplus; }
+
private:
// clocks
void init_clock_domains(void);
@@ -271,6 +275,9 @@ private:
// options
bool gpu_deadlock_detect;
+ int m_ptx_convert_to_ptxplus;
+ int m_ptx_save_converted_ptxplus;
+ unsigned m_ptx_force_max_capability;
// stats
struct shader_core_stats *m_shader_stats;
diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc
index 061952a..54e4ffa 100644
--- a/src/gpgpu-sim/shader.cc
+++ b/src/gpgpu-sim/shader.cc
@@ -593,7 +593,7 @@ void pdom_warp_ctx_t::print (FILE *fout) const
}
for (unsigned m=1,j=0; j<m_warp_size; j++, m<<=1)
fprintf(fout, "%c", ((warp->m_active_mask[k] & m)?'1':'0') );
- fprintf(fout, " pc: %4u", warp->m_pc[k] );
+ fprintf(fout, " pc: 0x%03x", warp->m_pc[k] );
if ( warp->m_recvg_pc[k] == (unsigned)-1 ) {
fprintf(fout," rp: ---- cd: %2u ", warp->m_calldepth[k] );
} else {
@@ -639,13 +639,14 @@ void shader_core_ctx::fetch_new()
// decode 1 or 2 instructions and place them into ibuffer
address_type pc = m_inst_fetch_buffer.m_pc;
const warp_inst_t* pI1 = ptx_fetch_inst(pc);
- assert(pI1);
m_warp[m_inst_fetch_buffer.m_warp_id].ibuffer_fill(0,pI1);
m_warp[m_inst_fetch_buffer.m_warp_id].inc_inst_in_pipeline();
- const warp_inst_t* pI2 = ptx_fetch_inst(pc+pI1->isize);
- if( pI2 ) {
- m_warp[m_inst_fetch_buffer.m_warp_id].ibuffer_fill(1,pI2);
- m_warp[m_inst_fetch_buffer.m_warp_id].inc_inst_in_pipeline();
+ if( pI1 ) {
+ const warp_inst_t* pI2 = ptx_fetch_inst(pc+pI1->isize);
+ if( pI2 ) {
+ m_warp[m_inst_fetch_buffer.m_warp_id].ibuffer_fill(1,pI2);
+ m_warp[m_inst_fetch_buffer.m_warp_id].inc_inst_in_pipeline();
+ }
}
m_inst_fetch_buffer.m_valid = false;
}
@@ -767,10 +768,12 @@ void shader_core_ctx::decode_new()
unsigned issued=0;
while( !m_warp[warp_id].waiting() && !m_warp[warp_id].ibuffer_empty() && (checked < 2) && (issued < 2) ) {
unsigned active_mask = m_pdom_warp[warp_id]->get_active_mask();
- const warp_inst_t *pI = m_warp[warp_id].ibuffer_next();
+ const warp_inst_t *pI = m_warp[warp_id].ibuffer_next_inst();
+ bool valid = m_warp[warp_id].ibuffer_next_valid();
unsigned pc,rpc;
m_pdom_warp[warp_id]->get_pdom_stack_top_info(&pc,&rpc);
if( pI ) {
+ assert(valid);
if( pc != pI->pc ) {
// control hazard
m_warp[warp_id].set_next_pc(pc);
@@ -785,6 +788,10 @@ void shader_core_ctx::decode_new()
issued++;
}
}
+ } else if( valid ) {
+ // this case can happen after a return instruction in diverged warp
+ m_warp[warp_id].set_next_pc(pc);
+ m_warp[warp_id].ibuffer_flush();
}
m_warp[warp_id].ibuffer_step();
checked++;
@@ -792,7 +799,7 @@ void shader_core_ctx::decode_new()
if ( issued ) {
m_last_warp_issued=warp_id;
break;
- }
+ }
}
}
@@ -852,7 +859,9 @@ mshr_entry* mshr_shader_unit::add_mshr(mem_access_t &access, warp_inst_t* warp)
// creates an mshr based on the access struct information
mshr_entry* mshr = alloc_free_mshr(access.space == tex_space);
mshr->init(access.addr,access.iswrite,access.space,warp->warp_id());
- mshr->add_inst(*warp);
+ warp_inst_t inst = *warp;
+ inst.set_active(access.warp_indices);
+ mshr->add_inst(inst);
if( m_shader_config->gpgpu_interwarp_mshr_merge ) {
mshr_entry* mergehit = m_mshr_lookup.shader_get_mergeable_mshr(mshr);
if (mergehit) {
@@ -2025,8 +2034,10 @@ void shd_warp_t::print_ibuffer( FILE *fout ) const
{
fprintf(fout," ibuffer[%2u] : ", m_warp_id );
for( unsigned i=0; i < IBUFFER_SIZE; i++) {
- const inst_t *inst = m_ibuffer[i];
+ const inst_t *inst = m_ibuffer[i].m_inst;
if( inst ) inst->print_insn(fout);
+ else if( m_ibuffer[i].m_valid )
+ fprintf(fout," <invalid instruction> ");
else fprintf(fout," <empty> ");
}
fprintf(fout,"\n");
diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h
index d9a8825..cdbaf37 100644
--- a/src/gpgpu-sim/shader.h
+++ b/src/gpgpu-sim/shader.h
@@ -146,8 +146,6 @@ public:
m_done_exit=false;
m_last_fetch=0;
m_next=0;
- for(unsigned i=0;i<IBUFFER_SIZE;i++)
- m_ibuffer[i]=NULL;
}
void init( address_type start_pc, unsigned cta_id, unsigned wid, unsigned active )
{
@@ -186,31 +184,32 @@ public:
void ibuffer_fill( unsigned slot, const warp_inst_t *pI )
{
assert(slot < IBUFFER_SIZE );
- m_ibuffer[slot]=pI;
+ m_ibuffer[slot].m_inst=pI;
+ m_ibuffer[slot].m_valid=true;
m_next=0;
}
bool ibuffer_empty() const
{
for( unsigned i=0; i < IBUFFER_SIZE; i++)
- if(m_ibuffer[i])
+ if(m_ibuffer[i].m_valid)
return false;
return true;
}
void ibuffer_flush()
{
for(unsigned i=0;i<IBUFFER_SIZE;i++) {
- if( m_ibuffer[i] )
+ if( m_ibuffer[i].m_valid )
dec_inst_in_pipeline();
- m_ibuffer[i]=NULL;
+ m_ibuffer[i].m_inst=NULL;
+ m_ibuffer[i].m_valid=false;
}
}
- const warp_inst_t *ibuffer_next()
- {
- return m_ibuffer[m_next];
- }
+ const warp_inst_t *ibuffer_next_inst() { return m_ibuffer[m_next].m_inst; }
+ bool ibuffer_next_valid() { return m_ibuffer[m_next].m_valid; }
void ibuffer_free()
{
- m_ibuffer[m_next] = NULL;
+ m_ibuffer[m_next].m_inst = NULL;
+ m_ibuffer[m_next].m_valid = false;
}
void ibuffer_step()
{
@@ -252,8 +251,13 @@ private:
unsigned n_completed; // number of threads in warp completed
class mshr_entry *m_imiss_pending;
-
- const warp_inst_t *m_ibuffer[IBUFFER_SIZE];
+
+ struct ibuffer_entry {
+ ibuffer_entry() { m_valid = false; m_inst = NULL; }
+ const warp_inst_t *m_inst;
+ bool m_valid;
+ };
+ ibuffer_entry m_ibuffer[IBUFFER_SIZE];
unsigned m_next;
unsigned m_n_atomic; // number of outstanding atomic operations
diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc
index 7947a8f..6f13cd2 100644
--- a/src/gpgpusim_entrypoint.cc
+++ b/src/gpgpusim_entrypoint.cc
@@ -167,19 +167,18 @@ int gpgpu_cuda_ptx_sim_main_perf( kernel_info_t grid,
return 0;
}
-int gpgpu_opencl_ptx_sim_main_perf( class function_info *entry,
+int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t grid,
struct dim3 gridDim,
struct dim3 blockDim,
gpgpu_ptx_sim_arg_list_t grid_params )
{
- kernel_info_t grid = gpgpu_opencl_ptx_sim_init_grid(entry,grid_params,gridDim,blockDim);
g_the_gpu.launch(grid);
sem_post(&g_sim_signal_start);
sem_wait(&g_sim_signal_finish);
return 0;
}
-int gpgpu_opencl_ptx_sim_main_func( class function_info *entry,
+int gpgpu_opencl_ptx_sim_main_func( kernel_info_t grid,
struct dim3 gridDim,
struct dim3 blockDim,
gpgpu_ptx_sim_arg_list_t grid_params )
diff --git a/src/gpgpusim_entrypoint.h b/src/gpgpusim_entrypoint.h
index f77a906..689dcf9 100644
--- a/src/gpgpusim_entrypoint.h
+++ b/src/gpgpusim_entrypoint.h
@@ -77,12 +77,12 @@ int gpgpu_cuda_ptx_sim_main_perf( kernel_info_t grid,
struct dim3 blockDim,
gpgpu_ptx_sim_arg_list_t grid_params );
-int gpgpu_opencl_ptx_sim_main_perf( class function_info *entry,
+int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t grid,
struct dim3 gridDim,
struct dim3 blockDim,
gpgpu_ptx_sim_arg_list_t grid_params );
-int gpgpu_opencl_ptx_sim_main_func( class function_info *entry,
+int gpgpu_opencl_ptx_sim_main_func( kernel_info_t grid,
struct dim3 gridDim,
struct dim3 blockDim,
gpgpu_ptx_sim_arg_list_t grid_params );
generated by cgit v1.3 (git 2.53.0) at 2026-07-11 00:00:08 +0000