summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorTor Aamodt <[email protected]>2010-10-08 07:21:51 -0800
committerTor Aamodt <[email protected]>2010-10-08 07:21:51 -0800
commit8c2954c263b4cfd756188362919c1a16a7189788 (patch)
treee4827655db3edb26d4668fc42e78f1e5731afa02
parent474e6d99c8efa9e3b8518eecb059546364370fd9 (diff)
1. refactoring cuda api code for loading PTX (removing external PTX loading entirely)
2. some bug fixes for warp_inst_t 3. creating a new class, gpgpu_t, which contains the functional "memory" state visible to all threads running on a GPU (doing this as part of my continuing effort to hunt down and eradicate every global variable that is not the top level "the gpu") 4. other misc. changes Almost passing CUDA 3.1 regression? oclHistogram keeps failing under torque, but does not fail when run on the command line from the same directory. [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7827]
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 02:44:53 +0000