summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--Makefile8
-rw-r--r--configs/Fermi/gpgpusim.config14
-rw-r--r--configs/QuadroFX5800/gpgpusim.config15
-rw-r--r--decuda_to_ptxplus/decudaInstList.cc65
-rw-r--r--decuda_to_ptxplus/decudaInstList.h12
-rw-r--r--decuda_to_ptxplus/ptx_parser.h22
-rw-r--r--libcuda/cuda_runtime_api.cc2102
-rw-r--r--libopencl/Makefile141
-rw-r--r--src/abstract_hardware_model.cc11
-rw-r--r--src/abstract_hardware_model.h2
-rw-r--r--src/cuda-sim/Makefile73
-rw-r--r--src/cuda-sim/cuda-sim.cc1
-rw-r--r--src/cuda-sim/instructions.cc34
-rw-r--r--src/cuda-sim/ptx.l3
-rw-r--r--src/cuda-sim/ptx.y2
-rw-r--r--src/cuda-sim/ptx_loader.cc187
-rw-r--r--src/cuda-sim/ptx_loader.h3
-rw-r--r--src/cuda-sim/ptx_parser.cc11
-rw-r--r--src/gpgpu-sim/dram.cc29
-rw-r--r--src/gpgpu-sim/dram.h10
-rw-r--r--src/gpgpu-sim/gpu-sim.cc4
-rw-r--r--src/gpgpu-sim/gpu-sim.h22
-rw-r--r--src/gpgpu-sim/shader.cc17
-rw-r--r--src/gpgpusim_entrypoint.cc11
24 files changed, 1693 insertions, 1106 deletions
diff --git a/Makefile b/Makefile
index 52fb93f..a146436 100644
--- a/Makefile
+++ b/Makefile
@@ -75,7 +75,9 @@ $(SIM_LIB_DIR)/libcudart.so: $(LIBS) cudalib
$(SIM_OBJ_FILES_DIR)/cuda-sim/decuda_pred_table/*.o \
$(SIM_OBJ_FILES_DIR)/gpgpu-sim/*.o \
$(SIM_OBJ_FILES_DIR)/intersim/*.o \
- $(SIM_OBJ_FILES_DIR)/*.o -lm -lz -lGL -pthread \
+ $(SIM_OBJ_FILES_DIR)/*.o -lm -lz -lGL -pthread \
+ $(GPGPUSIM_ROOT)/cuobjdump_to_ptxplus/cuobjdump_parser.o \
+ $(GPGPUSIM_ROOT)/cuobjdump_to_ptxplus/cuobjdump_lexer.o \
-o $(SIM_LIB_DIR)/libcudart.so
if [ ! -f $(SIM_LIB_DIR)/libcudart.so.2 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.2; fi
if [ ! -f $(SIM_LIB_DIR)/libcudart.so.3 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.3; fi
@@ -89,6 +91,8 @@ $(SIM_LIB_DIR)/libcudart.dylib: $(LIBS) cudalib
$(SIM_OBJ_FILES_DIR)/gpgpu-sim/*.o \
$(SIM_OBJ_FILES_DIR)/intersim/*.o \
$(SIM_OBJ_FILES_DIR)/*.o -lm -lz -pthread \
+ $(GPGPUSIM_ROOT)/cuobjdump_to_ptxplus/cuobjdump_parser.o \
+ $(GPGPUSIM_ROOT)/cuobjdump_to_ptxplus/cuobjdump_lexer.o \
-o $(SIM_LIB_DIR)/libcudart.dylib
$(SIM_LIB_DIR)/libOpenCL.so: $(LIBS) opencllib
@@ -106,6 +110,7 @@ $(SIM_LIB_DIR)/libOpenCL.so: $(LIBS) opencllib
cudalib: cuda-sim
$(MAKE) -C ./libcuda/ depend
$(MAKE) -C ./libcuda/
+ $(MAKE) -C cuobjdump_to_ptxplus/
cuda-sim:
$(MAKE) -C ./src/cuda-sim/ depend
@@ -165,5 +170,6 @@ endif
$(MAKE) clean -C ./src/gpgpu-sim/
$(MAKE) clean -C ./src/
$(MAKE) clean -C ./decuda_to_ptxplus/
+ $(MAKE) clean -C ./cuobjdump_to_ptxplus/
rm -rf $(SIM_LIB_DIR)
rm -rf $(SIM_OBJ_FILES_DIR)
diff --git a/configs/Fermi/gpgpusim.config b/configs/Fermi/gpgpusim.config
index b20fe86..e1bcf1c 100644
--- a/configs/Fermi/gpgpusim.config
+++ b/configs/Fermi/gpgpusim.config
@@ -66,8 +66,9 @@
-gpgpu_mem_addr_mapping dramid@8;00000000.00000000.00000000.00000000.0000RRRR.RRRRRRRR.RRBBBCCC.CCCSSSSS
# GDDR5 timing from hynix H5GQ1H24AFR
-# {nbk:tCCD:tRRD:tRCD:tRAS:tRP:tRC:CL:WL:tCDLR:tWR}
--gpgpu_dram_timing_opt 16:2:5:12:28:12:35:10:7:6:12
+# to disable bank groups, set nbkgrp to 1 and tCCDL and tRTPL to 0
+# {nbk:nbkgrp:tCCD:tCCDL:tRRD:tRCD:tRAS:tRP:tRTPL:tRC:CL:WL:tCDLR:tWR}
+-gpgpu_dram_timing_opt 16:4:2:3:5:12:28:12:2:35:10:7:6:12
# GDDR3
#-gpgpu_dram_timing_opt 8:2:8:12:25:10:35:10:7:6:11
@@ -80,9 +81,12 @@
-gpgpu_runtime_stat 500
-enable_ptx_file_line_stats 1
-# enable SASS execution
-#-gpgpu_ptx_convert_to_ptxplus 1
-#-gpgpu_ptx_save_converted_ptxplus 1
+# Using cuobjdump to extract ptx/SASS
+-gpgpu_ptx_use_cuobjdump 0
+
+# SASS execution (only supported with CUDA >= 4.0)
+-gpgpu_ptx_convert_to_ptxplus 0
+-gpgpu_ptx_save_converted_ptxplus 0
# enable operand collector
-gpgpu_operand_collector_num_units_sp 6
diff --git a/configs/QuadroFX5800/gpgpusim.config b/configs/QuadroFX5800/gpgpusim.config
index b5e1f01..6c33a85 100644
--- a/configs/QuadroFX5800/gpgpusim.config
+++ b/configs/QuadroFX5800/gpgpusim.config
@@ -37,19 +37,20 @@
-gpgpu_mem_address_mask 1
-gpgpu_mem_addr_mapping dramid@8;00000000.00000000.00000000.00000000.0000RRRR.RRRRRRRR.RRBBBCCC.CCCSSSSS
# GDDR3 timing from Samsung K4J52324QH-HC12 @ 800MHz
-# {nbk:tCCD:tRRD:tRCD:tRAS:tRP:tRC:CL:WL:tCDLR:tWR}
--gpgpu_dram_timing_opt 8:2:8:12:25:10:35:10:7:6:11
+# {nbk:nbkgrp:tCCD:tCCDL:tRRD:tRCD:tRAS:tRP:tRTPL:tRC:CL:WL:tCDLR:tWR}
+-gpgpu_dram_timing_opt 8:1:2:0:8:12:25:10:0:35:10:7:6:11
# stat collection
-gpgpu_memlatency_stat 14
-gpgpu_runtime_stat 500
-enable_ptx_file_line_stats 1
-# enable SASS execution
-# For this to work correctly, you need to build the binary using older cuda (<3.0)
-# Make sure this condition is met before enabling ptxplus
-#-gpgpu_ptx_convert_to_ptxplus 1
-#-gpgpu_ptx_save_converted_ptxplus 1
+# Using cuobjdump to extract ptx/SASS
+-gpgpu_ptx_use_cuobjdump 0
+
+# SASS execution (only supported with CUDA >= 4.0)
+-gpgpu_ptx_convert_to_ptxplus 0
+-gpgpu_ptx_save_converted_ptxplus 0
# enable operand collector
-gpgpu_operand_collector_num_units_sp 6
diff --git a/decuda_to_ptxplus/decudaInstList.cc b/decuda_to_ptxplus/decudaInstList.cc
index 3a77de3..4035d76 100644
--- a/decuda_to_ptxplus/decudaInstList.cc
+++ b/decuda_to_ptxplus/decudaInstList.cc
@@ -27,7 +27,14 @@
#include "decudaInstList.h"
-
+#define P_DEBUG 1
+#define DPRINTF(...) \
+ if(P_DEBUG) { \
+ printf("(%s:%u) ", __FILE__, __LINE__); \
+ printf(__VA_ARGS__); \
+ printf("\n"); \
+ fflush(stdout); \
+ }
extern void output(const char * text);
//Constructor
@@ -129,6 +136,13 @@ void decudaInstList::printHeaderInstList()
break;
}
}
+ for ( std::list<std::string>::iterator iter = m_realTexList.begin();
+ iter != m_realTexList.end();
+ iter ++) {
+ output(".tex .u64 ");
+ output((*iter).c_str());
+ output(";\n");
+ }
}
void decudaInstList::printNewPtxList(decudaInstList* headerInfo)
@@ -426,7 +440,7 @@ void decudaInstList::addDoublePredReg(std::string pred, std::string reg, bool lo
void decudaInstList::addTex(std::string tex)
{
std::string origTex = tex;
-
+ DPRINTF("decudaInstList::addTex tex=%s", tex.c_str());
// If $tex# tex from decuda, then use index to get real tex name
if(tex.substr(0, 4) == "$tex") {
tex = tex.substr(4, tex.size()-4);
@@ -446,9 +460,9 @@ void decudaInstList::addTex(std::string tex)
}
// Add the tex to instruction operand list
- char* texName = new char [strlen(origTex.c_str())+1];
- strcpy(texName, origTex.c_str());
- getListEnd().addOperand(texName);
+ //char* texName = new char [strlen(origTex.c_str())+1];
+ //strcpy(texName, origTex.c_str());
+ //getListEnd().addOperand(texName);
}
@@ -478,24 +492,34 @@ void decudaInstList::addConstMemoryValue(std::string constMemoryValue)
m_constMemoryList.back().m_constMemory.push_back(constMemoryValue);
}
+void decudaInstList::addConstMemoryValue2(std::string constMemoryValue)
+{
+ m_constMemoryList2.back().m_constMemory.push_back(constMemoryValue);
+}
+
// set type of constant memory
void decudaInstList::setConstMemoryType(const char* type)
{
m_constMemoryList.back().type = type;
}
+void decudaInstList::setConstMemoryType2(const char* type)
+{
+ m_constMemoryList2.back().type = type;
+}
+
// print const memory directive
void decudaInstList::printMemory()
{
// Constant memory
- std::list<constMemory>::iterator i;
- for(i=m_constMemoryList.begin(); i!=m_constMemoryList.end(); ++i) {
+
+ for(std::list<constMemory>::iterator i=m_constMemoryList.begin(); i!=m_constMemoryList.end(); ++i) {
char line[40];
// Global or entry specific
if(i->entryIndex == 0)
- sprintf(line, ".const %s c%d[%d] = {", i->type, i->index, i->m_constMemory.size());
+ sprintf(line, ".const %s constant0[%d] = {", i->type, i->m_constMemory.size());
else
sprintf(line, ".const %s ce%dc%d[%d] = {", i->type, i->entryIndex, i->index, i->m_constMemory.size());
@@ -512,6 +536,26 @@ void decudaInstList::printMemory()
output("\n};\n\n");
}
+
+ for(std::list<constMemory2>::iterator i=m_constMemoryList2.begin(); i!=m_constMemoryList2.end(); ++i) {
+ char line[1024];
+
+ // Global or entry specific
+ sprintf(line, ".const %s constant1%s[%d] = {", i->type, i->kernel, i->m_constMemory.size());
+
+ output(line);
+
+ std::list<std::string>::iterator j;
+ int l=0;
+ for(j=i->m_constMemory.begin(); j!=i->m_constMemory.end(); ++j) {
+ if(j!=i->m_constMemory.begin())
+ output(", ");
+ if( (l++ % 4) == 0) output("\n ");
+ output(j->c_str());
+ }
+ output("\n};\n\n");
+ }
+
// Next, print out the local memory declaration
std::list<decudaEntry>::iterator e;
int eIndex=1; // entry index starts from 1 from the first blank entry is missing here (only in header entry list)
@@ -565,13 +609,16 @@ void decudaInstList::addVector(char* vector, int vectorSize) {
// memType: 0=constant, 1=shared, 2=global, 3=local
void decudaInstList::addMemoryOperand(std::string mem, int memType) {
std::string origMem = mem;
+ assert(0 && "This function is obsolete and shouldn't be used");
// If constant memory type, add prefix for entry specific constant memory
if(memType == 0) {
// Entry-specific constant memory c1
if(mem.substr(0, 3) == "c1[") {
std::stringstream out;
- out << "ce" << m_entryList.size() << mem;
+ //out << "ce" << m_entryList.size() << mem;
+ mem = mem.substr(2);
+ out << "constant1" << this->m_entryList.back().m_entryName << mem;
mem = out.str();
}
// Global memory c14
diff --git a/decuda_to_ptxplus/decudaInstList.h b/decuda_to_ptxplus/decudaInstList.h
index da0d48a..d7387ee 100644
--- a/decuda_to_ptxplus/decudaInstList.h
+++ b/decuda_to_ptxplus/decudaInstList.h
@@ -34,7 +34,7 @@
#include <iostream>
#include <sstream>
#include <stdlib.h>
-#include <regex.hpp>
+#include <boost/regex.hpp>
#include <map>
// Used for entry specific constant memory segments (c1)
@@ -46,6 +46,13 @@ struct constMemory
std::list<std::string> m_constMemory;
};
+struct constMemory2
+{
+ const char* kernel;
+ const char* type;
+ std::list<std::string> m_constMemory;
+};
+
// Used for uninitialized constant memory (globally defined)
struct constMemoryPtr
{
@@ -102,6 +109,7 @@ protected:
// Const memory list
std::list<constMemory> m_constMemoryList;
+ std::list<constMemory2> m_constMemoryList2;
// Const memory pointers list
std::list<constMemoryPtr> m_constMemoryPtrList;
@@ -157,7 +165,9 @@ public:
void addEntryConstMemory(int index); // add entry specific const memory
void addConstMemory(int index); // add global const memory
void setConstMemoryType(const char* type); // set type of constant memory
+ void setConstMemoryType2(const char* type); // set type of constant memory
void addConstMemoryValue(std::string constMemoryValue); // add const memory
+ void addConstMemoryValue2(std::string constMemoryValue); // add const memory
std::list<std::string> getRealTexList(); // get the list of real tex names
void setRealTexList(std::list<std::string> realTexList); // set the list of real tex names
diff --git a/decuda_to_ptxplus/ptx_parser.h b/decuda_to_ptxplus/ptx_parser.h
index c429ead..bed82bb 100644
--- a/decuda_to_ptxplus/ptx_parser.h
+++ b/decuda_to_ptxplus/ptx_parser.h
@@ -37,6 +37,14 @@
#define NON_ARRAY_IDENTIFIER 1
#define ARRAY_IDENTIFIER_NO_DIM 2
#define ARRAY_IDENTIFIER 3
+#define P_DEBUG 1
+#define DPRINTF(...) \
+ if(P_DEBUG) { \
+ printf("(%s:%u) ", __FILE__, __LINE__); \
+ printf(__VA_ARGS__); \
+ printf("\n"); \
+ fflush(stdout); \
+ }
enum _memory_space_t {
@@ -128,21 +136,30 @@ void add_function_name( const char *headerInput )
//void add_space_spec(int headerInput)
void add_space_spec( enum _memory_space_t spec, int value )
{
+ DPRINTF("add_space_spec: spec=%u", spec);
+ decudaInst *instEntry;
switch(spec)
{
case param_space_unclassified:
if(inEntryDirective && inParamDirective) {
- decudaInst *instEntry = new decudaInst();
+ instEntry = new decudaInst();
instEntry->setBase(".param");
g_headerList->add(instEntry);
}
break;
case tex_space:
inTexDirective = true;
- decudaInst *instEntry = new decudaInst();
+ instEntry = new decudaInst();
instEntry->setBase(".tex");
g_headerList->add(instEntry);
break;
+ case const_space:
+ if(inEntryDirective) {
+ instEntry = new decudaInst();
+ instEntry->setBase(".const");
+ g_headerList->add(instEntry);
+ }
+ break;
}
}
@@ -277,6 +294,7 @@ void func_header(const char* headerBase)
// If start of an entry
if((strcmp(headerBase, ".entry")==0)||(strcmp(headerBase, ".func")==0)) {
+ inEntryDirective = true;
g_headerList->addEntry("");
decudaInst *instEntry = new decudaInst();
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 7a651bb..c0d7be3 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -109,11 +109,11 @@
#include <stdarg.h>
#ifdef OPENGL_SUPPORT
#define GL_GLEXT_PROTOTYPES
- #ifdef __APPLE__
- #include <GLUT/glut.h> // Apple's version of GLUT is here
- #else
- #include <GL/gl.h>
- #endif
+#ifdef __APPLE__
+#include <GLUT/glut.h> // Apple's version of GLUT is here
+#else
+#include <GL/gl.h>
+#endif
#endif
#define __CUDA_RUNTIME_API_H__
@@ -140,27 +140,27 @@ static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsign
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 CUctx_st* context );
+ gpgpu_ptx_sim_arg_list_t args,
+ struct dim3 gridDim,
+ struct dim3 blockDim,
+ struct CUctx_st* context );
/*DEVICE_BUILTIN*/
struct cudaArray
{
- void *devPtr;
- int devPtr32;
- struct cudaChannelFormatDesc desc;
- int width;
- int height;
- int size; //in bytes
- unsigned dimensions;
+ void *devPtr;
+ int devPtr32;
+ struct cudaChannelFormatDesc desc;
+ int width;
+ int height;
+ int size; //in bytes
+ unsigned dimensions;
};
#if !defined(__dv)
#if defined(__cplusplus)
#define __dv(v) \
- = v
+ = v
#else /* __cplusplus */
#define __dv(v)
#endif /* __cplusplus */
@@ -172,7 +172,7 @@ extern stream_manager *g_stream_manager;
void register_ptx_function( const char *name, function_info *impl )
{
- // no longer need this
+ // no longer need this
}
#if defined __APPLE__
@@ -190,174 +190,174 @@ void register_ptx_function( const char *name, function_info *impl )
#endif
struct _cuda_device_id {
- _cuda_device_id(gpgpu_sim* gpu) {m_id = 0; m_next = NULL; m_gpgpu=gpu;}
- struct _cuda_device_id *next() { return m_next; }
- unsigned num_shader() const { return m_gpgpu->get_config().num_shader(); }
- int num_devices() const {
- if( m_next == NULL ) return 1;
- else return 1 + m_next->num_devices();
- }
- struct _cuda_device_id *get_device( unsigned n )
- {
- assert( n < (unsigned)num_devices() );
- struct _cuda_device_id *p=this;
- for(unsigned i=0; i<n; i++)
- p = p->m_next;
- return p;
- }
- const struct cudaDeviceProp *get_prop() const
- {
- return m_gpgpu->get_prop();
- }
- unsigned get_id() const { return m_id; }
+ _cuda_device_id(gpgpu_sim* gpu) {m_id = 0; m_next = NULL; m_gpgpu=gpu;}
+ struct _cuda_device_id *next() { return m_next; }
+ unsigned num_shader() const { return m_gpgpu->get_config().num_shader(); }
+ int num_devices() const {
+ if( m_next == NULL ) return 1;
+ else return 1 + m_next->num_devices();
+ }
+ struct _cuda_device_id *get_device( unsigned n )
+ {
+ assert( n < (unsigned)num_devices() );
+ struct _cuda_device_id *p=this;
+ for(unsigned i=0; i<n; i++)
+ p = p->m_next;
+ return p;
+ }
+ const struct cudaDeviceProp *get_prop() const
+ {
+ return m_gpgpu->get_prop();
+ }
+ unsigned get_id() const { return m_id; }
- gpgpu_sim *get_gpgpu() { return m_gpgpu; }
+ gpgpu_sim *get_gpgpu() { return m_gpgpu; }
private:
- unsigned m_id;
- class gpgpu_sim *m_gpgpu;
- struct _cuda_device_id *m_next;
+ 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; }
+ CUctx_st( _cuda_device_id *gpu ) { m_gpu = gpu; }
- _cuda_device_id *get_device() { return m_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_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 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 )
- {
- if( m_code.find(fat_cubin_handle) != m_code.end() ) {
- 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;
- } else {
- m_kernel_lookup[hostFun] = NULL;
- }
- }
+ void register_function( unsigned fat_cubin_handle, const char *hostFun, const char *deviceFun )
+ {
+ if( m_code.find(fat_cubin_handle) != m_code.end() ) {
+ 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;
+ } else {
+ m_kernel_lookup[hostFun] = NULL;
+ }
+ }
- 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;
- }
+ 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
+ _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 kernel_config {
public:
- kernel_config( dim3 GridDim, dim3 BlockDim, size_t sharedMem, struct CUstream_st *stream )
- {
- m_GridDim=GridDim;
- m_BlockDim=BlockDim;
- m_sharedMem=sharedMem;
- m_stream = stream;
- }
- void set_arg( const void *arg, size_t size, size_t offset )
- {
- m_args.push_front( gpgpu_ptx_sim_arg(arg,size,offset) );
- }
- dim3 grid_dim() const { return m_GridDim; }
- dim3 block_dim() const { return m_BlockDim; }
- gpgpu_ptx_sim_arg_list_t get_args() { return m_args; }
- struct CUstream_st *get_stream() { return m_stream; }
-
+ kernel_config( dim3 GridDim, dim3 BlockDim, size_t sharedMem, struct CUstream_st *stream )
+ {
+ m_GridDim=GridDim;
+ m_BlockDim=BlockDim;
+ m_sharedMem=sharedMem;
+ m_stream = stream;
+ }
+ void set_arg( const void *arg, size_t size, size_t offset )
+ {
+ m_args.push_front( gpgpu_ptx_sim_arg(arg,size,offset) );
+ }
+ dim3 grid_dim() const { return m_GridDim; }
+ dim3 block_dim() const { return m_BlockDim; }
+ gpgpu_ptx_sim_arg_list_t get_args() { return m_args; }
+ struct CUstream_st *get_stream() { return m_stream; }
+
private:
- dim3 m_GridDim;
- dim3 m_BlockDim;
- size_t m_sharedMem;
- struct CUstream_st *m_stream;
- gpgpu_ptx_sim_arg_list_t m_args;
+ dim3 m_GridDim;
+ dim3 m_BlockDim;
+ size_t m_sharedMem;
+ struct CUstream_st *m_stream;
+ gpgpu_ptx_sim_arg_list_t m_args;
};
class _cuda_device_id *GPGPUSim_Init()
{
- static _cuda_device_id *the_device = NULL;
- if( !the_device ) {
- gpgpu_sim *the_gpu = gpgpu_ptx_sim_init_perf();
+ static _cuda_device_id *the_device = NULL;
+ if( !the_device ) {
+ gpgpu_sim *the_gpu = gpgpu_ptx_sim_init_perf();
- cudaDeviceProp *prop = (cudaDeviceProp *) calloc(sizeof(cudaDeviceProp),1);
- snprintf(prop->name,256,"GPGPU-Sim_v%s", g_gpgpusim_version_string );
- prop->major = 2;
- prop->minor = 0;
- prop->totalGlobalMem = 0x40000000 /* 1 GB */;
- prop->memPitch = 0;
- prop->maxThreadsPerBlock = 512;
- prop->maxThreadsDim[0] = 512;
- prop->maxThreadsDim[1] = 512;
- prop->maxThreadsDim[2] = 512;
- prop->maxGridSize[0] = 0x40000000;
- prop->maxGridSize[1] = 0x40000000;
- prop->maxGridSize[2] = 0x40000000;
- prop->totalConstMem = 0x40000000;
- prop->textureAlignment = 0;
- prop->sharedMemPerBlock = the_gpu->shared_mem_size();
- prop->regsPerBlock = the_gpu->num_registers_per_core();
- prop->warpSize = the_gpu->wrp_size();
- prop->clockRate = the_gpu->shader_clock();
+ cudaDeviceProp *prop = (cudaDeviceProp *) calloc(sizeof(cudaDeviceProp),1);
+ snprintf(prop->name,256,"GPGPU-Sim_v%s", g_gpgpusim_version_string );
+ prop->major = 2;
+ prop->minor = 0;
+ prop->totalGlobalMem = 0x40000000 /* 1 GB */;
+ prop->memPitch = 0;
+ prop->maxThreadsPerBlock = 512;
+ prop->maxThreadsDim[0] = 512;
+ prop->maxThreadsDim[1] = 512;
+ prop->maxThreadsDim[2] = 512;
+ prop->maxGridSize[0] = 0x40000000;
+ prop->maxGridSize[1] = 0x40000000;
+ prop->maxGridSize[2] = 0x40000000;
+ prop->totalConstMem = 0x40000000;
+ prop->textureAlignment = 0;
+ prop->sharedMemPerBlock = the_gpu->shared_mem_size();
+ prop->regsPerBlock = the_gpu->num_registers_per_core();
+ prop->warpSize = the_gpu->wrp_size();
+ prop->clockRate = the_gpu->shader_clock();
#if (CUDART_VERSION >= 2010)
- prop->multiProcessorCount = the_gpu->get_config().num_shader();
+ prop->multiProcessorCount = the_gpu->get_config().num_shader();
#endif
- the_gpu->set_prop(prop);
- the_device = new _cuda_device_id(the_gpu);
- }
- start_sim_thread(1);
- return the_device;
+ the_gpu->set_prop(prop);
+ the_device = new _cuda_device_id(the_gpu);
+ }
+ start_sim_thread(1);
+ 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;
+ 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();
+ 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);
- fflush(stderr);
- printf("\n\nGPGPU-Sim PTX: Execution error: CUDA API function \"%s()\" has not been implemented yet.\n"
- " [$GPGPUSIM_ROOT/libcuda/%s around line %u]\n\n\n",
- func,__FILE__, line );
- fflush(stdout);
- abort();
+ fflush(stdout);
+ fflush(stderr);
+ printf("\n\nGPGPU-Sim PTX: Execution error: CUDA API function \"%s()\" has not been implemented yet.\n"
+ " [$GPGPUSIM_ROOT/libcuda/%s around line %u]\n\n\n",
+ func,__FILE__, line );
+ fflush(stdout);
+ abort();
}
@@ -366,27 +366,27 @@ void cuda_not_implemented( const char* func, unsigned line )
void gpgpusim_ptx_error_impl( const char *func, const char *file, unsigned line, const char *msg, ... )
{
- va_list ap;
- char buf[1024];
- va_start(ap,msg);
- vsnprintf(buf,1024,msg,ap);
- va_end(ap);
+ va_list ap;
+ char buf[1024];
+ va_start(ap,msg);
+ vsnprintf(buf,1024,msg,ap);
+ va_end(ap);
- printf("GPGPU-Sim CUDA API: %s\n", buf);
- printf(" [%s:%u : %s]\n", file, line, func );
- abort();
+ printf("GPGPU-Sim CUDA API: %s\n", buf);
+ printf(" [%s:%u : %s]\n", file, line, func );
+ abort();
}
void gpgpusim_ptx_assert_impl( int test_value, const char *func, const char *file, unsigned line, const char *msg, ... )
{
- va_list ap;
- char buf[1024];
- va_start(ap,msg);
- vsnprintf(buf,1024,msg,ap);
- va_end(ap);
+ va_list ap;
+ char buf[1024];
+ va_start(ap,msg);
+ vsnprintf(buf,1024,msg,ap);
+ va_end(ap);
- if ( test_value == 0 )
- gpgpusim_ptx_error_impl(func, file, line, msg);
+ if ( test_value == 0 )
+ gpgpusim_ptx_error_impl(func, file, line, msg);
}
@@ -398,860 +398,1142 @@ int g_active_device = 0; //active gpu that runs the code
std::list<kernel_config> g_cuda_launch_stack;
/*******************************************************************************
-* *
-* *
-* *
-*******************************************************************************/
+ * *
+ * *
+ * *
+ *******************************************************************************/
extern "C" {
/*******************************************************************************
-* *
-* *
-* *
-*******************************************************************************/
+ * *
+ * *
+ * *
+ *******************************************************************************/
__host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size)
{
- CUctx_st* context = GPGPUSim_Context();
- *devPtr = context->get_device()->get_gpgpu()->gpu_malloc(size);
- if(g_debug_execution >= 3)
- printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *devPtr);
- if ( *devPtr ) {
- return g_last_cudaError = cudaSuccess;
- } else {
- return g_last_cudaError = cudaErrorMemoryAllocation;
- }
+ CUctx_st* context = GPGPUSim_Context();
+ *devPtr = context->get_device()->get_gpgpu()->gpu_malloc(size);
+ if(g_debug_execution >= 3)
+ printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *devPtr);
+ if ( *devPtr ) {
+ return g_last_cudaError = cudaSuccess;
+ } else {
+ return g_last_cudaError = cudaErrorMemoryAllocation;
+ }
}
-
+
__host__ cudaError_t CUDARTAPI cudaMallocHost(void **ptr, size_t size)
{
- GPGPUSim_Context();
- *ptr = malloc(size);
- if ( *ptr ) {
- return g_last_cudaError = cudaSuccess;
- } else {
- return g_last_cudaError = cudaErrorMemoryAllocation;
- }
- }
- __host__ cudaError_t CUDARTAPI cudaMallocPitch(void **devPtr, size_t *pitch, size_t width, size_t height)
+ GPGPUSim_Context();
+ *ptr = malloc(size);
+ if ( *ptr ) {
+ return g_last_cudaError = cudaSuccess;
+ } else {
+ return g_last_cudaError = cudaErrorMemoryAllocation;
+ }
+}
+__host__ cudaError_t CUDARTAPI cudaMallocPitch(void **devPtr, size_t *pitch, size_t width, size_t height)
{
- unsigned malloc_width_inbytes = width;
- printf("GPGPU-Sim PTX: cudaMallocPitch (width = %d)\n", malloc_width_inbytes);
- CUctx_st* ctx = GPGPUSim_Context();
- *devPtr = ctx->get_device()->get_gpgpu()->gpu_malloc(malloc_width_inbytes*height);
- pitch[0] = malloc_width_inbytes;
- if ( *devPtr ) {
- return g_last_cudaError = cudaSuccess;
- } else {
- return g_last_cudaError = cudaErrorMemoryAllocation;
- }
+ unsigned malloc_width_inbytes = width;
+ printf("GPGPU-Sim PTX: cudaMallocPitch (width = %d)\n", malloc_width_inbytes);
+ CUctx_st* ctx = GPGPUSim_Context();
+ *devPtr = ctx->get_device()->get_gpgpu()->gpu_malloc(malloc_width_inbytes*height);
+ pitch[0] = malloc_width_inbytes;
+ if ( *devPtr ) {
+ return g_last_cudaError = cudaSuccess;
+ } else {
+ return g_last_cudaError = cudaErrorMemoryAllocation;
+ }
}
__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);
- 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)context->get_device()->get_gpgpu()->gpu_mallocarray(size);
- printf("GPGPU-Sim PTX: cudaMallocArray: devPtr32 = %d\n", ((*array)->devPtr32));
- ((*array)->devPtr) = (void*) (long long) ((*array)->devPtr32);
- if ( ((*array)->devPtr) ) {
- return g_last_cudaError = cudaSuccess;
- } else {
- return g_last_cudaError = cudaErrorMemoryAllocation;
- }
+ unsigned size = width * height * ((desc->x + desc->y + desc->z + desc->w)/8);
+ 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)context->get_device()->get_gpgpu()->gpu_mallocarray(size);
+ printf("GPGPU-Sim PTX: cudaMallocArray: devPtr32 = %d\n", ((*array)->devPtr32));
+ ((*array)->devPtr) = (void*) (long long) ((*array)->devPtr32);
+ if ( ((*array)->devPtr) ) {
+ return g_last_cudaError = cudaSuccess;
+ } else {
+ return g_last_cudaError = cudaErrorMemoryAllocation;
+ }
}
__host__ cudaError_t CUDARTAPI cudaFree(void *devPtr)
{
- // TODO... manage g_global_mem space?
- return g_last_cudaError = cudaSuccess;
+ // TODO... manage g_global_mem space?
+ return g_last_cudaError = cudaSuccess;
}
- __host__ cudaError_t CUDARTAPI cudaFreeHost(void *ptr)
+__host__ cudaError_t CUDARTAPI cudaFreeHost(void *ptr)
{
- free (ptr); // this will crash the system if called twice
- return g_last_cudaError = cudaSuccess;
+ free (ptr); // this will crash the system if called twice
+ return g_last_cudaError = cudaSuccess;
}
- __host__ cudaError_t CUDARTAPI cudaFreeArray(struct cudaArray *array)
+__host__ cudaError_t CUDARTAPI cudaFreeArray(struct cudaArray *array)
{
- // TODO... manage g_global_mem space?
- return g_last_cudaError = cudaSuccess;
+ // TODO... manage g_global_mem space?
+ return g_last_cudaError = cudaSuccess;
};
/*******************************************************************************
-* *
-* *
-* *
-*******************************************************************************/
+ * *
+ * *
+ * *
+ *******************************************************************************/
__host__ cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind)
{
- //CUctx_st *context = GPGPUSim_Context();
- //gpgpu_t *gpu = context->get_device()->get_gpgpu();
- if(g_debug_execution >= 3)
- printf("GPGPU-Sim PTX: cudaMemcpy(): devPtr = %p\n", dst);
- if( kind == cudaMemcpyHostToDevice )
- g_stream_manager->push( stream_operation(src,(size_t)dst,count,0) );
- else if( kind == cudaMemcpyDeviceToHost )
- g_stream_manager->push( stream_operation((size_t)src,dst,count,0) );
- else if( kind == cudaMemcpyDeviceToDevice )
- g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,0) );
- else {
- printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported cudaMemcpyKind\n");
- abort();
- }
- return g_last_cudaError = cudaSuccess;
+ //CUctx_st *context = GPGPUSim_Context();
+ //gpgpu_t *gpu = context->get_device()->get_gpgpu();
+ if(g_debug_execution >= 3)
+ printf("GPGPU-Sim PTX: cudaMemcpy(): devPtr = %p\n", dst);
+ if( kind == cudaMemcpyHostToDevice )
+ g_stream_manager->push( stream_operation(src,(size_t)dst,count,0) );
+ else if( kind == cudaMemcpyDeviceToHost )
+ g_stream_manager->push( stream_operation((size_t)src,dst,count,0) );
+ else if( kind == cudaMemcpyDeviceToDevice )
+ g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,0) );
+ else {
+ printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported cudaMemcpyKind\n");
+ abort();
+ }
+ return g_last_cudaError = cudaSuccess;
}
- __host__ cudaError_t CUDARTAPI cudaMemcpyToArray(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind)
+__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");
- if( kind == cudaMemcpyHostToDevice )
- gpu->memcpy_to_gpu( (size_t)(dst->devPtr), src, size);
- else if( kind == cudaMemcpyDeviceToHost )
- gpu->memcpy_from_gpu( dst->devPtr, (size_t)src, size);
- else if( kind == cudaMemcpyDeviceToDevice )
- gpu->memcpy_gpu_to_gpu( (size_t)(dst->devPtr), (size_t)src, size);
- else {
- printf("GPGPU-Sim PTX: cudaMemcpyToArray - ERROR : unsupported cudaMemcpyKind\n");
- abort();
- }
- dst->devPtr32 = (unsigned) (size_t)(dst->devPtr);
- return g_last_cudaError = cudaSuccess;
+ CUctx_st *context = GPGPUSim_Context();
+ gpgpu_t *gpu = context->get_device()->get_gpgpu();
+ size_t size = count;
+ printf("GPGPU-Sim PTX: cudaMemcpyToArray\n");
+ if( kind == cudaMemcpyHostToDevice )
+ gpu->memcpy_to_gpu( (size_t)(dst->devPtr), src, size);
+ else if( kind == cudaMemcpyDeviceToHost )
+ gpu->memcpy_from_gpu( dst->devPtr, (size_t)src, size);
+ else if( kind == cudaMemcpyDeviceToDevice )
+ gpu->memcpy_gpu_to_gpu( (size_t)(dst->devPtr), (size_t)src, size);
+ else {
+ printf("GPGPU-Sim PTX: cudaMemcpyToArray - ERROR : unsupported cudaMemcpyKind\n");
+ abort();
+ }
+ dst->devPtr32 = (unsigned) (size_t)(dst->devPtr);
+ return g_last_cudaError = cudaSuccess;
}
- __host__ cudaError_t CUDARTAPI cudaMemcpyFromArray(void *dst, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t count, enum cudaMemcpyKind kind)
+__host__ cudaError_t CUDARTAPI cudaMemcpyFromArray(void *dst, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t count, enum cudaMemcpyKind kind)
{
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaErrorUnknown;
+ cuda_not_implemented(__my_func__,__LINE__);
+ return g_last_cudaError = cudaErrorUnknown;
}
- __host__ cudaError_t CUDARTAPI cudaMemcpyArrayToArray(struct cudaArray *dst, size_t wOffsetDst, size_t hOffsetDst, const struct cudaArray *src, size_t wOffsetSrc, size_t hOffsetSrc, size_t count, enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToDevice))
+__host__ cudaError_t CUDARTAPI cudaMemcpyArrayToArray(struct cudaArray *dst, size_t wOffsetDst, size_t hOffsetDst, const struct cudaArray *src, size_t wOffsetSrc, size_t hOffsetSrc, size_t count, enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToDevice))
{
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaErrorUnknown;
+ cuda_not_implemented(__my_func__,__LINE__);
+ return g_last_cudaError = cudaErrorUnknown;
}
- __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)
+__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)
{
- 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 )
- gpu->memcpy_to_gpu( (size_t)dst, src, size );
- else if( kind == cudaMemcpyDeviceToHost )
- gpu->memcpy_from_gpu( dst, (size_t)src, size );
- else if( kind == cudaMemcpyDeviceToDevice )
- gpu->memcpy_gpu_to_gpu( (size_t)dst, (size_t)src, size);
- else {
- printf("GPGPU-Sim PTX: cudaMemcpy2D - ERROR : unsupported cudaMemcpyKind\n");
- abort();
- }
- return g_last_cudaError = cudaSuccess;
+ 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 )
+ gpu->memcpy_to_gpu( (size_t)dst, src, size );
+ else if( kind == cudaMemcpyDeviceToHost )
+ gpu->memcpy_from_gpu( dst, (size_t)src, size );
+ else if( kind == cudaMemcpyDeviceToDevice )
+ gpu->memcpy_gpu_to_gpu( (size_t)dst, (size_t)src, size);
+ else {
+ printf("GPGPU-Sim PTX: cudaMemcpy2D - ERROR : unsupported cudaMemcpyKind\n");
+ abort();
+ }
+ return g_last_cudaError = cudaSuccess;
}
- __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)
+__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;
- 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;
- gpgpusim_ptx_assert( (dst->dimensions==2), "copy to none 2D array not supported" );
- gpgpusim_ptx_assert( (wOffset==0), "non-zero wOffset not yet supported" );
- gpgpusim_ptx_assert( (hOffset==0), "non-zero hOffset not yet supported" );
- gpgpusim_ptx_assert( (dst->height == (int)height), "partial copy not supported" );
- gpgpusim_ptx_assert( (elem_size*dst->width == width), "partial copy not supported" );
- gpgpusim_ptx_assert( (spitch == width), "spitch != width not supported" );
- if( kind == cudaMemcpyHostToDevice )
- gpu->memcpy_to_gpu( (size_t)(dst->devPtr), src, size);
- else if( kind == cudaMemcpyDeviceToHost )
- gpu->memcpy_from_gpu( dst->devPtr, (size_t)src, size);
- else if( kind == cudaMemcpyDeviceToDevice )
- gpu->memcpy_gpu_to_gpu( (size_t)dst->devPtr, (size_t)src, size);
- else {
- printf("GPGPU-Sim PTX: cudaMemcpy2D - ERROR : unsupported cudaMemcpyKind\n");
- abort();
- }
- dst->devPtr32 = (unsigned) (size_t)(dst->devPtr);
- return g_last_cudaError = cudaSuccess;
+ CUctx_st *context = GPGPUSim_Context();
+ gpgpu_t *gpu = context->get_device()->get_gpgpu();
+ size_t size = spitch*height;
+ 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;
+ gpgpusim_ptx_assert( (dst->dimensions==2), "copy to none 2D array not supported" );
+ gpgpusim_ptx_assert( (wOffset==0), "non-zero wOffset not yet supported" );
+ gpgpusim_ptx_assert( (hOffset==0), "non-zero hOffset not yet supported" );
+ gpgpusim_ptx_assert( (dst->height == (int)height), "partial copy not supported" );
+ gpgpusim_ptx_assert( (elem_size*dst->width == width), "partial copy not supported" );
+ gpgpusim_ptx_assert( (spitch == width), "spitch != width not supported" );
+ if( kind == cudaMemcpyHostToDevice )
+ gpu->memcpy_to_gpu( (size_t)(dst->devPtr), src, size);
+ else if( kind == cudaMemcpyDeviceToHost )
+ gpu->memcpy_from_gpu( dst->devPtr, (size_t)src, size);
+ else if( kind == cudaMemcpyDeviceToDevice )
+ gpu->memcpy_gpu_to_gpu( (size_t)dst->devPtr, (size_t)src, size);
+ else {
+ printf("GPGPU-Sim PTX: cudaMemcpy2D - ERROR : unsupported cudaMemcpyKind\n");
+ abort();
+ }
+ dst->devPtr32 = (unsigned) (size_t)(dst->devPtr);
+ return g_last_cudaError = cudaSuccess;
}
- __host__ cudaError_t CUDARTAPI cudaMemcpy2DFromArray(void *dst, size_t dpitch, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t width, size_t height, enum cudaMemcpyKind kind)
+__host__ cudaError_t CUDARTAPI cudaMemcpy2DFromArray(void *dst, size_t dpitch, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t width, size_t height, enum cudaMemcpyKind kind)
{
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaErrorUnknown;
+ cuda_not_implemented(__my_func__,__LINE__);
+ return g_last_cudaError = cudaErrorUnknown;
}
- __host__ cudaError_t CUDARTAPI cudaMemcpy2DArrayToArray(struct cudaArray *dst, size_t wOffsetDst, size_t hOffsetDst, const struct cudaArray *src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToDevice))
+__host__ cudaError_t CUDARTAPI cudaMemcpy2DArrayToArray(struct cudaArray *dst, size_t wOffsetDst, size_t hOffsetDst, const struct cudaArray *src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToDevice))
{
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaErrorUnknown;
+ cuda_not_implemented(__my_func__,__LINE__);
+ return g_last_cudaError = cudaErrorUnknown;
}
- __host__ cudaError_t CUDARTAPI cudaMemcpyToSymbol(const char *symbol, const void *src, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyHostToDevice))
+__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);
- //stream_operation( const char *symbol, const void *src, size_t count, size_t offset )
- g_stream_manager->push( stream_operation(src,symbol,count,offset,0) );
- //gpgpu_ptx_sim_memcpy_symbol(symbol,src,count,offset,1,context->get_device()->get_gpgpu());
- return g_last_cudaError = cudaSuccess;
+ //CUctx_st *context = GPGPUSim_Context();
+ assert(kind == cudaMemcpyHostToDevice);
+ printf("GPGPU-Sim PTX: cudaMemcpyToSymbol: symbol = %p\n", symbol);
+ //stream_operation( const char *symbol, const void *src, size_t count, size_t offset )
+ g_stream_manager->push( stream_operation(src,symbol,count,offset,0) );
+ //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))
+__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);
- g_stream_manager->push( stream_operation(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;
+ //CUctx_st *context = GPGPUSim_Context();
+ assert(kind == cudaMemcpyDeviceToHost);
+ printf("GPGPU-Sim PTX: cudaMemcpyFromSymbol: symbol = %p\n", symbol);
+ g_stream_manager->push( stream_operation(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;
}
/*******************************************************************************
-* *
-* *
-* *
-*******************************************************************************/
+ * *
+ * *
+ * *
+ *******************************************************************************/
- __host__ cudaError_t CUDARTAPI cudaMemcpyAsync(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream)
+__host__ cudaError_t CUDARTAPI cudaMemcpyAsync(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream)
{
- struct CUstream_st *s = (struct CUstream_st *)stream;
- switch( kind ) {
- case cudaMemcpyHostToDevice: g_stream_manager->push( stream_operation(src,(size_t)dst,count,s) ); break;
- case cudaMemcpyDeviceToHost: g_stream_manager->push( stream_operation((size_t)src,dst,count,s) ); break;
- case cudaMemcpyDeviceToDevice: g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,s) ); break;
- default:
- abort();
- }
- return g_last_cudaError = cudaSuccess;
+ struct CUstream_st *s = (struct CUstream_st *)stream;
+ switch( kind ) {
+ case cudaMemcpyHostToDevice: g_stream_manager->push( stream_operation(src,(size_t)dst,count,s) ); break;
+ case cudaMemcpyDeviceToHost: g_stream_manager->push( stream_operation((size_t)src,dst,count,s) ); break;
+ case cudaMemcpyDeviceToDevice: g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,s) ); break;
+ default:
+ abort();
+ }
+ return g_last_cudaError = cudaSuccess;
}
- __host__ cudaError_t CUDARTAPI cudaMemcpyToArrayAsync(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream)
+__host__ cudaError_t CUDARTAPI cudaMemcpyToArrayAsync(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream)
{
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaErrorUnknown;
+ cuda_not_implemented(__my_func__,__LINE__);
+ return g_last_cudaError = cudaErrorUnknown;
}
- __host__ cudaError_t CUDARTAPI cudaMemcpyFromArrayAsync(void *dst, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream)
+__host__ cudaError_t CUDARTAPI cudaMemcpyFromArrayAsync(void *dst, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream)
{
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaErrorUnknown;
+ cuda_not_implemented(__my_func__,__LINE__);
+ return g_last_cudaError = cudaErrorUnknown;
}
- __host__ cudaError_t CUDARTAPI cudaMemcpy2DAsync(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream)
+__host__ cudaError_t CUDARTAPI cudaMemcpy2DAsync(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream)
{
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaErrorUnknown;
+ cuda_not_implemented(__my_func__,__LINE__);
+ return g_last_cudaError = cudaErrorUnknown;
}
- __host__ cudaError_t CUDARTAPI cudaMemcpy2DToArrayAsync(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream)
+__host__ cudaError_t CUDARTAPI cudaMemcpy2DToArrayAsync(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream)
{
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaErrorUnknown;
+ cuda_not_implemented(__my_func__,__LINE__);
+ return g_last_cudaError = cudaErrorUnknown;
}
- __host__ cudaError_t CUDARTAPI cudaMemcpy2DFromArrayAsync(void *dst, size_t dpitch, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream)
+__host__ cudaError_t CUDARTAPI cudaMemcpy2DFromArrayAsync(void *dst, size_t dpitch, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream)
{
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaErrorUnknown;
+ cuda_not_implemented(__my_func__,__LINE__);
+ return g_last_cudaError = cudaErrorUnknown;
}
/*******************************************************************************
-* *
-* *
-* *
-*******************************************************************************/
+ * *
+ * *
+ * *
+ *******************************************************************************/
__host__ cudaError_t CUDARTAPI cudaMemset(void *mem, int c, size_t count)
{
- CUctx_st *context = GPGPUSim_Context();
- gpgpu_t *gpu = context->get_device()->get_gpgpu();
- gpu->gpu_memset((size_t)mem, c, count);
- return g_last_cudaError = cudaSuccess;
+ CUctx_st *context = GPGPUSim_Context();
+ gpgpu_t *gpu = context->get_device()->get_gpgpu();
+ gpu->gpu_memset((size_t)mem, c, count);
+ return g_last_cudaError = cudaSuccess;
}
- __host__ cudaError_t CUDARTAPI cudaMemset2D(void *mem, size_t pitch, int c, size_t width, size_t height)
+__host__ cudaError_t CUDARTAPI cudaMemset2D(void *mem, size_t pitch, int c, size_t width, size_t height)
{
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaErrorUnknown;
+ cuda_not_implemented(__my_func__,__LINE__);
+ return g_last_cudaError = cudaErrorUnknown;
}
/*******************************************************************************
-* *
-* *
-* *
-*******************************************************************************/
+ * *
+ * *
+ * *
+ *******************************************************************************/
- __host__ cudaError_t CUDARTAPI cudaGetSymbolAddress(void **devPtr, const char *symbol)
+__host__ cudaError_t CUDARTAPI cudaGetSymbolAddress(void **devPtr, const char *symbol)
{
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaErrorUnknown;
+ cuda_not_implemented(__my_func__,__LINE__);
+ return g_last_cudaError = cudaErrorUnknown;
}
- __host__ cudaError_t CUDARTAPI cudaGetSymbolSize(size_t *size, const char *symbol)
+__host__ cudaError_t CUDARTAPI cudaGetSymbolSize(size_t *size, const char *symbol)
{
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaErrorUnknown;
+ cuda_not_implemented(__my_func__,__LINE__);
+ return g_last_cudaError = cudaErrorUnknown;
}
/*******************************************************************************
-* *
-* *
-* *
-*******************************************************************************/
- __host__ cudaError_t CUDARTAPI cudaGetDeviceCount(int *count)
+ * *
+ * *
+ * *
+ *******************************************************************************/
+__host__ cudaError_t CUDARTAPI cudaGetDeviceCount(int *count)
{
- _cuda_device_id *dev = GPGPUSim_Init();
- *count = dev->num_devices();
- return g_last_cudaError = cudaSuccess;
+ _cuda_device_id *dev = GPGPUSim_Init();
+ *count = dev->num_devices();
+ return g_last_cudaError = cudaSuccess;
}
- __host__ cudaError_t CUDARTAPI cudaGetDeviceProperties(struct cudaDeviceProp *prop, int device)
+__host__ cudaError_t CUDARTAPI cudaGetDeviceProperties(struct cudaDeviceProp *prop, int device)
{
- _cuda_device_id *dev = GPGPUSim_Init();
- if (device <= dev->num_devices() ) {
- *prop= *dev->get_prop();
- return g_last_cudaError = cudaSuccess;
- } else {
- return g_last_cudaError = cudaErrorInvalidDevice;
- }
+ _cuda_device_id *dev = GPGPUSim_Init();
+ if (device <= dev->num_devices() ) {
+ *prop= *dev->get_prop();
+ return g_last_cudaError = cudaSuccess;
+ } else {
+ return g_last_cudaError = cudaErrorInvalidDevice;
+ }
}
- __host__ cudaError_t CUDARTAPI cudaChooseDevice(int *device, const struct cudaDeviceProp *prop)
+__host__ cudaError_t CUDARTAPI cudaChooseDevice(int *device, const struct cudaDeviceProp *prop)
{
- _cuda_device_id *dev = GPGPUSim_Init();
- *device = dev->get_id();
- return g_last_cudaError = cudaSuccess;
+ _cuda_device_id *dev = GPGPUSim_Init();
+ *device = dev->get_id();
+ return g_last_cudaError = cudaSuccess;
}
-
- __host__ cudaError_t CUDARTAPI cudaSetDevice(int device)
+
+__host__ cudaError_t CUDARTAPI cudaSetDevice(int device)
{
- //set the active device to run cuda
- if ( device <= GPGPUSim_Init()->num_devices() ) {
- g_active_device = device;
- return g_last_cudaError = cudaSuccess;
- } else {
- return g_last_cudaError = cudaErrorInvalidDevice;
- }
+ //set the active device to run cuda
+ if ( device <= GPGPUSim_Init()->num_devices() ) {
+ g_active_device = device;
+ return g_last_cudaError = cudaSuccess;
+ } else {
+ return g_last_cudaError = cudaErrorInvalidDevice;
+ }
}
-
- __host__ cudaError_t CUDARTAPI cudaGetDevice(int *device)
+
+__host__ cudaError_t CUDARTAPI cudaGetDevice(int *device)
{
- *device = g_active_device;
- return g_last_cudaError = cudaSuccess;
+ *device = g_active_device;
+ return g_last_cudaError = cudaSuccess;
}
/*******************************************************************************
-* *
-* *
-* *
-*******************************************************************************/
+ * *
+ * *
+ * *
+ *******************************************************************************/
- __host__ cudaError_t CUDARTAPI cudaBindTexture(size_t *offset,
- const struct textureReference *texref,
- const void *devPtr,
- const struct cudaChannelFormatDesc *desc,
- size_t size __dv(UINT_MAX))
+__host__ cudaError_t CUDARTAPI cudaBindTexture(size_t *offset,
+ const struct textureReference *texref,
+ const void *devPtr,
+ const struct cudaChannelFormatDesc *desc,
+ size_t size __dv(UINT_MAX))
{
- CUctx_st *context = GPGPUSim_Context();
- gpgpu_t *gpu = context->get_device()->get_gpgpu();
- printf("GPGPU-Sim PTX: in cudaBindTexture: sizeof(struct textureReference) = %zu\n", sizeof(struct textureReference));
- struct cudaArray *array;
- array = (struct cudaArray*) malloc(sizeof(struct cudaArray));
- array->desc = *desc;
- array->size = size;
- array->width = size;
- array->height = 1;
- array->dimensions = 1;
- array->devPtr = (void*)devPtr;
- array->devPtr32 = (int)(long long)devPtr;
- offset = 0;
- printf("GPGPU-Sim PTX: size = %zu\n", size);
- printf("GPGPU-Sim PTX: texref = %p, array = %p\n", texref, array);
- printf("GPGPU-Sim PTX: devPtr32 = %x\n", array->devPtr32);
- printf("GPGPU-Sim PTX: Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref));
- printf("GPGPU-Sim PTX: ChannelFormatDesc: x=%d, y=%d, z=%d, w=%d\n", desc->x, desc->y, desc->z, desc->w);
- printf("GPGPU-Sim PTX: Texture Normalized? = %d\n", texref->normalized);
- gpu->gpgpu_ptx_sim_bindTextureToArray(texref, array);
- devPtr = (void*)(long long)array->devPtr32;
- printf("GPGPU-Sim PTX: devPtr = %p\n", devPtr);
- return g_last_cudaError = cudaSuccess;
+ CUctx_st *context = GPGPUSim_Context();
+ gpgpu_t *gpu = context->get_device()->get_gpgpu();
+ printf("GPGPU-Sim PTX: in cudaBindTexture: sizeof(struct textureReference) = %zu\n", sizeof(struct textureReference));
+ struct cudaArray *array;
+ array = (struct cudaArray*) malloc(sizeof(struct cudaArray));
+ array->desc = *desc;
+ array->size = size;
+ array->width = size;
+ array->height = 1;
+ array->dimensions = 1;
+ array->devPtr = (void*)devPtr;
+ array->devPtr32 = (int)(long long)devPtr;
+ offset = 0;
+ printf("GPGPU-Sim PTX: size = %zu\n", size);
+ printf("GPGPU-Sim PTX: texref = %p, array = %p\n", texref, array);
+ printf("GPGPU-Sim PTX: devPtr32 = %x\n", array->devPtr32);
+ printf("GPGPU-Sim PTX: Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref));
+ printf("GPGPU-Sim PTX: ChannelFormatDesc: x=%d, y=%d, z=%d, w=%d\n", desc->x, desc->y, desc->z, desc->w);
+ printf("GPGPU-Sim PTX: Texture Normalized? = %d\n", texref->normalized);
+ gpu->gpgpu_ptx_sim_bindTextureToArray(texref, array);
+ devPtr = (void*)(long long)array->devPtr32;
+ printf("GPGPU-Sim PTX: devPtr = %p\n", devPtr);
+ return g_last_cudaError = cudaSuccess;
}
__host__ cudaError_t CUDARTAPI cudaBindTextureToArray(const struct textureReference *texref, const struct cudaArray *array, const struct cudaChannelFormatDesc *desc)
{
- CUctx_st *context = GPGPUSim_Context();
- gpgpu_t *gpu = context->get_device()->get_gpgpu();
- printf("GPGPU-Sim PTX: in cudaBindTextureToArray: %p %p\n", texref, array);
- printf("GPGPU-Sim PTX: devPtr32 = %x\n", array->devPtr32);
- printf("GPGPU-Sim PTX: Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref));
- printf("GPGPU-Sim PTX: Texture Normalized? = %d\n", texref->normalized);
- gpu->gpgpu_ptx_sim_bindTextureToArray(texref, array);
- return g_last_cudaError = cudaSuccess;
+ CUctx_st *context = GPGPUSim_Context();
+ gpgpu_t *gpu = context->get_device()->get_gpgpu();
+ printf("GPGPU-Sim PTX: in cudaBindTextureToArray: %p %p\n", texref, array);
+ printf("GPGPU-Sim PTX: devPtr32 = %x\n", array->devPtr32);
+ printf("GPGPU-Sim PTX: Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref));
+ printf("GPGPU-Sim PTX: Texture Normalized? = %d\n", texref->normalized);
+ gpu->gpgpu_ptx_sim_bindTextureToArray(texref, array);
+ return g_last_cudaError = cudaSuccess;
}
__host__ cudaError_t CUDARTAPI cudaUnbindTexture(const struct textureReference *texref)
{
- return g_last_cudaError = cudaSuccess;
+ return g_last_cudaError = cudaSuccess;
}
__host__ cudaError_t CUDARTAPI cudaGetTextureAlignmentOffset(size_t *offset, const struct textureReference *texref)
{
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaErrorUnknown;
+ cuda_not_implemented(__my_func__,__LINE__);
+ return g_last_cudaError = cudaErrorUnknown;
}
__host__ cudaError_t CUDARTAPI cudaGetTextureReference(const struct textureReference **texref, const char *symbol)
{
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaErrorUnknown;
+ cuda_not_implemented(__my_func__,__LINE__);
+ return g_last_cudaError = cudaErrorUnknown;
}
__host__ cudaError_t CUDARTAPI cudaGetChannelDesc(struct cudaChannelFormatDesc *desc, const struct cudaArray *array)
{
- *desc = array->desc;
- return g_last_cudaError = cudaSuccess;
+ *desc = array->desc;
+ return g_last_cudaError = cudaSuccess;
}
__host__ struct cudaChannelFormatDesc CUDARTAPI cudaCreateChannelDesc(int x, int y, int z, int w, enum cudaChannelFormatKind f)
{
- struct cudaChannelFormatDesc dummy;
- dummy.x = x;
- dummy.y = y;
- dummy.z = z;
- dummy.w = w;
- dummy.f = f;
- return dummy;
+ struct cudaChannelFormatDesc dummy;
+ dummy.x = x;
+ dummy.y = y;
+ dummy.z = z;
+ dummy.w = w;
+ dummy.f = f;
+ return dummy;
}
__host__ cudaError_t CUDARTAPI cudaGetLastError(void)
{
- return g_last_cudaError;
+ return g_last_cudaError;
}
__host__ const char* CUDARTAPI cudaGetErrorString(cudaError_t error)
{
- if( g_last_cudaError == cudaSuccess )
- return "no error";
- char buf[1024];
- snprintf(buf,1024,"<<GPGPU-Sim PTX: there was an error (code = %d)>>", g_last_cudaError);
- return strdup(buf);
+ if( g_last_cudaError == cudaSuccess )
+ return "no error";
+ char buf[1024];
+ snprintf(buf,1024,"<<GPGPU-Sim PTX: there was an error (code = %d)>>", g_last_cudaError);
+ return strdup(buf);
}
__host__ cudaError_t CUDARTAPI cudaConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem, cudaStream_t stream)
{
- struct CUstream_st *s = (struct CUstream_st *)stream;
- g_cuda_launch_stack.push_back( kernel_config(gridDim,blockDim,sharedMem,s) );
- return g_last_cudaError = cudaSuccess;
+ struct CUstream_st *s = (struct CUstream_st *)stream;
+ g_cuda_launch_stack.push_back( kernel_config(gridDim,blockDim,sharedMem,s) );
+ return g_last_cudaError = cudaSuccess;
}
__host__ cudaError_t CUDARTAPI cudaSetupArgument(const void *arg, size_t size, size_t offset)
{
- gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" );
- kernel_config &config = g_cuda_launch_stack.back();
- config.set_arg(arg,size,offset);
+ gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" );
+ kernel_config &config = g_cuda_launch_stack.back();
+ config.set_arg(arg,size,offset);
+
+ struct gpgpu_ptx_sim_arg *param = (gpgpu_ptx_sim_arg*) calloc(1,sizeof(struct gpgpu_ptx_sim_arg));
+ param->m_start = arg;
+ param->m_nbytes = size;
+ param->m_offset = offset;
- struct gpgpu_ptx_sim_arg *param = (gpgpu_ptx_sim_arg*) calloc(1,sizeof(struct gpgpu_ptx_sim_arg));
- param->m_start = arg;
- param->m_nbytes = size;
- param->m_offset = offset;
-
- return g_last_cudaError = cudaSuccess;
+ return g_last_cudaError = cudaSuccess;
}
__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);
- gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" );
- kernel_config config = g_cuda_launch_stack.back();
- struct CUstream_st *stream = config.get_stream();
- printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s) on stream %u\n", hostFun,
- g_ptx_sim_mode?"functional simulation":"performance simulation", stream?stream->get_uid():0 );
- kernel_info_t *grid = gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context);
- std::string kname = grid->name();
- dim3 gridDim = config.grid_dim();
- dim3 blockDim = config.block_dim();
- printf("GPGPU-Sim PTX: pushing kernel \'%s\' to stream %u, gridDim= (%u,%u,%u) blockDim = (%u,%u,%u) \n",
- kname.c_str(), stream?stream->get_uid():0, gridDim.x,gridDim.y,gridDim.z,blockDim.x,blockDim.y,blockDim.z );
- stream_operation op(grid,g_ptx_sim_mode,stream);
- g_stream_manager->push(op);
- g_cuda_launch_stack.pop_back();
- return g_last_cudaError = cudaSuccess;
+ CUctx_st* context = GPGPUSim_Context();
+ char *mode = getenv("PTX_SIM_MODE_FUNC");
+ if( mode )
+ sscanf(mode,"%u", &g_ptx_sim_mode);
+ gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" );
+ kernel_config config = g_cuda_launch_stack.back();
+ struct CUstream_st *stream = config.get_stream();
+ printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s) on stream %u\n", hostFun,
+ g_ptx_sim_mode?"functional simulation":"performance simulation", stream?stream->get_uid():0 );
+ kernel_info_t *grid = gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context);
+ std::string kname = grid->name();
+ dim3 gridDim = config.grid_dim();
+ dim3 blockDim = config.block_dim();
+ printf("GPGPU-Sim PTX: pushing kernel \'%s\' to stream %u, gridDim= (%u,%u,%u) blockDim = (%u,%u,%u) \n",
+ kname.c_str(), stream?stream->get_uid():0, gridDim.x,gridDim.y,gridDim.z,blockDim.x,blockDim.y,blockDim.z );
+ stream_operation op(grid,g_ptx_sim_mode,stream);
+ g_stream_manager->push(op);
+ g_cuda_launch_stack.pop_back();
+ return g_last_cudaError = cudaSuccess;
}
/*******************************************************************************
-* *
-* *
-* *
-*******************************************************************************/
+ * *
+ * *
+ * *
+ *******************************************************************************/
__host__ cudaError_t CUDARTAPI cudaStreamCreate(cudaStream_t *stream)
{
- printf("GPGPU-Sim PTX: cudaStreamCreate\n");
+ printf("GPGPU-Sim PTX: cudaStreamCreate\n");
#if (CUDART_VERSION >= 3000)
- *stream = new struct CUstream_st();
- g_stream_manager->add_stream(*stream);
+ *stream = new struct CUstream_st();
+ g_stream_manager->add_stream(*stream);
#else
- *stream = 0;
- printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__);
+ *stream = 0;
+ printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__);
#endif
- return g_last_cudaError = cudaSuccess;
+ return g_last_cudaError = cudaSuccess;
}
__host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream)
{
#if (CUDART_VERSION >= 3000)
- g_stream_manager->destroy_stream(stream);
+ g_stream_manager->destroy_stream(stream);
#endif
- return g_last_cudaError = cudaSuccess;
+ return g_last_cudaError = cudaSuccess;
}
__host__ cudaError_t CUDARTAPI cudaStreamSynchronize(cudaStream_t stream)
{
#if (CUDART_VERSION >= 3000)
- if( stream == NULL )
- return g_last_cudaError = cudaErrorInvalidResourceHandle;
- stream->synchronize();
+ if( stream == NULL )
+ return g_last_cudaError = cudaErrorInvalidResourceHandle;
+ stream->synchronize();
#else
- printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__);
+ printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__);
#endif
- return g_last_cudaError = cudaSuccess;
+ return g_last_cudaError = cudaSuccess;
}
__host__ cudaError_t CUDARTAPI cudaStreamQuery(cudaStream_t stream)
{
#if (CUDART_VERSION >= 3000)
- if( stream == NULL )
- return g_last_cudaError = cudaErrorInvalidResourceHandle;
- return g_last_cudaError = stream->empty()?cudaSuccess:cudaErrorNotReady;
+ if( stream == NULL )
+ return g_last_cudaError = cudaErrorInvalidResourceHandle;
+ return g_last_cudaError = stream->empty()?cudaSuccess:cudaErrorNotReady;
#else
- printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__);
- return g_last_cudaError = cudaSuccess; // it is always success because all cuda calls are synchronous
+ printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__);
+ return g_last_cudaError = cudaSuccess; // it is always success because all cuda calls are synchronous
#endif
}
/*******************************************************************************
-* *
-* *
-* *
-*******************************************************************************/
+ * *
+ * *
+ * *
+ *******************************************************************************/
__host__ cudaError_t CUDARTAPI cudaEventCreate(cudaEvent_t *event)
{
- CUevent_st *e = new CUevent_st(false);
- g_timer_events[e->get_uid()] = e;
+ CUevent_st *e = new CUevent_st(false);
+ g_timer_events[e->get_uid()] = e;
#if CUDART_VERSION >= 3000
- *event = e;
+ *event = e;
#else
- *event = e->get_uid();
+ *event = e->get_uid();
#endif
- return g_last_cudaError = cudaSuccess;
+ return g_last_cudaError = cudaSuccess;
}
CUevent_st *get_event(cudaEvent_t event)
{
- unsigned event_uid;
+ unsigned event_uid;
#if CUDART_VERSION >= 3000
- event_uid = event->get_uid();
+ event_uid = event->get_uid();
#else
- event_uid = event;
+ event_uid = event;
#endif
- event_tracker_t::iterator e = g_timer_events.find(event_uid);
- if( e == g_timer_events.end() )
- return NULL;
- return e->second;
+ event_tracker_t::iterator e = g_timer_events.find(event_uid);
+ if( e == g_timer_events.end() )
+ return NULL;
+ return e->second;
}
__host__ cudaError_t CUDARTAPI cudaEventRecord(cudaEvent_t event, cudaStream_t stream)
{
- CUevent_st *e = get_event(event);
- if( !e ) return g_last_cudaError = cudaErrorUnknown;
- struct CUstream_st *s = (struct CUstream_st *)stream;
- stream_operation op(e,s);
- g_stream_manager->push(op);
- return g_last_cudaError = cudaSuccess;
+ CUevent_st *e = get_event(event);
+ if( !e ) return g_last_cudaError = cudaErrorUnknown;
+ struct CUstream_st *s = (struct CUstream_st *)stream;
+ stream_operation op(e,s);
+ g_stream_manager->push(op);
+ return g_last_cudaError = cudaSuccess;
}
__host__ cudaError_t CUDARTAPI cudaEventQuery(cudaEvent_t event)
{
- CUevent_st *e = get_event(event);
- if( e == NULL ) {
- return g_last_cudaError = cudaErrorInvalidValue;
- } else if( e->done() ) {
- return g_last_cudaError = cudaSuccess;
- } else {
- return g_last_cudaError = cudaErrorNotReady;
- }
+ CUevent_st *e = get_event(event);
+ if( e == NULL ) {
+ return g_last_cudaError = cudaErrorInvalidValue;
+ } else if( e->done() ) {
+ return g_last_cudaError = cudaSuccess;
+ } else {
+ return g_last_cudaError = cudaErrorNotReady;
+ }
}
__host__ cudaError_t CUDARTAPI cudaEventSynchronize(cudaEvent_t event)
{
- printf("GPGPU-Sim API: cudaEventSynchronize ** waiting for event\n");
- fflush(stdout);
- CUevent_st *e = (CUevent_st*) event;
- while( !e->done() )
- ;
- printf("GPGPU-Sim API: cudaEventSynchronize ** event detected\n");
- fflush(stdout);
- return g_last_cudaError = cudaSuccess;
+ printf("GPGPU-Sim API: cudaEventSynchronize ** waiting for event\n");
+ fflush(stdout);
+ CUevent_st *e = (CUevent_st*) event;
+ while( !e->done() )
+ ;
+ printf("GPGPU-Sim API: cudaEventSynchronize ** event detected\n");
+ fflush(stdout);
+ return g_last_cudaError = cudaSuccess;
}
__host__ cudaError_t CUDARTAPI cudaEventDestroy(cudaEvent_t event)
{
- CUevent_st *e = get_event(event);
- unsigned event_uid = e->get_uid();
- event_tracker_t::iterator pe = g_timer_events.find(event_uid);
- if( pe == g_timer_events.end() )
- return g_last_cudaError = cudaErrorInvalidValue;
- g_timer_events.erase(pe);
- return g_last_cudaError = cudaSuccess;
+ CUevent_st *e = get_event(event);
+ unsigned event_uid = e->get_uid();
+ event_tracker_t::iterator pe = g_timer_events.find(event_uid);
+ if( pe == g_timer_events.end() )
+ return g_last_cudaError = cudaErrorInvalidValue;
+ g_timer_events.erase(pe);
+ return g_last_cudaError = cudaSuccess;
}
__host__ cudaError_t CUDARTAPI cudaEventElapsedTime(float *ms, cudaEvent_t start, cudaEvent_t end)
{
- time_t elapsed_time;
- CUevent_st *s = get_event(start);
- CUevent_st *e = get_event(end);
- if( s==NULL || e==NULL )
- return g_last_cudaError = cudaErrorUnknown;
- elapsed_time = e->clock() - s->clock();
- *ms = 1000*elapsed_time;
- return g_last_cudaError = cudaSuccess;
+ time_t elapsed_time;
+ CUevent_st *s = get_event(start);
+ CUevent_st *e = get_event(end);
+ if( s==NULL || e==NULL )
+ return g_last_cudaError = cudaErrorUnknown;
+ elapsed_time = e->clock() - s->clock();
+ *ms = 1000*elapsed_time;
+ return g_last_cudaError = cudaSuccess;
}
/*******************************************************************************
-* *
-* *
-* *
-*******************************************************************************/
+ * *
+ * *
+ * *
+ *******************************************************************************/
__host__ cudaError_t CUDARTAPI cudaThreadExit(void)
{
- exit_simulation();
- return g_last_cudaError = cudaSuccess;
+ exit_simulation();
+ return g_last_cudaError = cudaSuccess;
}
__host__ cudaError_t CUDARTAPI cudaThreadSynchronize(void)
{
- //Called on host side
- synchronize();
- return g_last_cudaError = cudaSuccess;
+ //Called on host side
+ synchronize();
+ return g_last_cudaError = cudaSuccess;
};
int CUDARTAPI __cudaSynchronizeThreads(void**, void*)
{
- return cudaThreadExit();
+ return cudaThreadExit();
}
+
+
+
/*******************************************************************************
-* *
-* *
-* *
-*******************************************************************************/
+ * *
+ * *
+ * *
+ *******************************************************************************/
+
+//#include "../../cuobjdump_to_ptxplus/cuobjdump_parser.h"
+
+enum cuobjdumpSectionType {
+ PTXSECTION=0,
+ ELFSECTION
+};
+
+typedef struct cuobjdumpSectionRec {
+ enum cuobjdumpSectionType type;
+ char* arch;
+ char* identifier;
+ char* ptxfilename;
+ char* elffilename;
+ char* sassfilename;
+}cuobjdumpSection;
+
+
+
+std::list<cuobjdumpSection> cuobjdump;
+
+// sectiontype: 0 for ptx, 1 for elf
+void addCuobjdumpSection(int sectiontype){
+ cuobjdumpSection x;
+ x.type = sectiontype? ELFSECTION: PTXSECTION;
+ cuobjdump.push_front(x);
+ printf("## Adding new section %s\n", x.type==PTXSECTION?"PTX":"ELF");
+}
+
+void setCuobjdumparch(const char* arch){
+ printf("Adding arch: %s\n", arch);
+ cuobjdump.front().arch = strdup(arch);
+}
+
+void setCuobjdumpidentifier(const char* identifier){
+ printf("Adding identifier: %s\n", identifier);
+ cuobjdump.front().identifier = strdup(identifier);
+}
+
+void setCuobjdumpptxfilename(const char* filename){
+ printf("Adding ptx filename: %s\n", filename);
+ cuobjdump.front().ptxfilename = strdup(filename);
+}
+
+void setCuobjdumpelffilename(const char* filename){
+ cuobjdump.front().elffilename = strdup(filename);
+}
+
+void setCuobjdumpsassfilename(const char* filename){
+ cuobjdump.front().sassfilename = strdup(filename);
+}
+extern "C" int cuobjdump_parse();
+extern "C" FILE *cuobjdump_in;
+//! Return PTX file
+/*!
+ * This Function extract the whole PTX (for all the files) using cuobjdump
+ * */
+void extract_ptx(){
+ char command[1000];
+ char* whole_code;
+ char fname[1024]="_ptx_whole_code_WESWW";
+ //! Running CUobjdump using dynamic link to current process
+ snprintf(command,1000,"cuobjdump -ptx -elf -sass /proc/%d/exe > %s",getpid(),fname);
+ printf("Running cuobjdump using \"%s\"\n", command);
+ int result = system(command);
+ if(result) {printf("ERROR: Failed to execute: %s\n", command); exit(1);}
+
+ printf("Parsing file %s\n", fname);
+ cuobjdump_in = fopen(fname, "r");
+
+ cuobjdump_parse();
+ fclose(cuobjdump_in);
+ printf("Done parsing!!!\n");
+ // Just to stop things here
+
+ /* Hadi's old code
+ FILE* fp = fopen(fname,"r");
+ //! finding size of the file
+ int lSize= 0;
+ fseek (fp , 0 , SEEK_END);
+ lSize = ftell (fp);
+ rewind (fp);
+ //! allocate and copy the entire ptx
+ whole_code = (char*)malloc(lSize * sizeof(char));
+ fread(whole_code,1,lSize,fp);
+ fclose(fp);
+ snprintf(command,1000,"rm %s",fname);
+ //system(command);
+ return whole_code;
+ */
+}
+
+//! Return proper ptx version
+/*!
+ * This function return newest ptx version inside the argument
+ * which is is not newer than forced_max_capability
+ */
+unsigned get_best_version(std::list<cuobjdumpSection> sectionlist, CUctx_st *context){
+
+ unsigned forced_max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability();
+ unsigned max_capability=0;
-void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin )
+ std::list<cuobjdumpSection>::iterator iter;
+
+ for ( iter = sectionlist.begin();
+ iter != sectionlist.end();
+ iter++
+ ){
+ unsigned capability = 0;
+ sscanf(iter->arch,"sm_%u", &capability);
+ if (capability > max_capability &&
+ (forced_max_capability == 0 || capability <= forced_max_capability)){
+ max_capability = capability;
+ }
+ }
+ return max_capability;
+
+ /* Hadi's old code
+ char* start_entry;
+ char* arch_pointer;
+ start_entry = strstr(ptx,"Fatbin ptx code:");
+ while(start_entry){
+ arch_pointer = strstr (start_entry,"arch = ");
+ unsigned int capability = 0;
+ sscanf(arch_pointer,"arch = sm_%u",&capability);
+ if(capability > max_capability && (forced_max_capability == 0 || capability <= forced_max_capability)){
+ max_capability = capability;
+ }
+ start_entry = strstr(arch_pointer,"Fatbin ptx code:");
+ }
+ return max_capability;
+ */
+}
+/*!
+ * Return number of diffrent ptx files generated in first argument
+ * which have sm version equal to selected_capability
+ */
+unsigned get_number_of_ptx(std::list<cuobjdumpSection> sectionlist, unsigned selected_capability){
+ int result = 0;
+
+ for ( std::list<cuobjdumpSection>::iterator iter = sectionlist.begin();
+ iter != sectionlist.end();
+ iter++
+ ){
+ unsigned capability = 0;
+ sscanf(iter->arch, "sm_%u", &capability);
+ if(capability <= selected_capability &&
+ iter->type == PTXSECTION){
+ printf("Found compatible PTX section with capability %s\n", iter->arch);
+ result++;
+ }
+ }
+ return result;
+ /* Hadi's old code
+ char* start_entry;
+ char* arch_pointer;
+ start_entry = strstr(ptx,"Fatbin ptx code:");
+ while(start_entry){
+ arch_pointer = strstr (start_entry,"arch = ");
+ unsigned int capability = 0;
+ sscanf(arch_pointer,"arch = sm_%u",&capability);
+ if(capability == selected_capability)
+ result++;
+ start_entry = strstr(arch_pointer,"Fatbin ptx code:");
+ }
+ */
+
+}
+
+char* readfile (const char* filename){
+ assert (filename != NULL);
+ char str[128];
+ FILE* fp = fopen(filename,"r");
+ if (!fp) {
+ printf("ERROR: Could not open file %s for reading\n", filename);
+ assert (0);
+ }
+ //! finding size of the file
+ int filesize= 0;
+ fseek (fp , 0 , SEEK_END);
+
+ filesize = ftell (fp);
+ fseek (fp, 0, SEEK_SET);
+ //! allocate and copy the entire ptx
+ char* ret = (char*)malloc((filesize +1)* sizeof(char));
+ fread(ret,1,filesize,fp);
+ ret[filesize]='\0';
+ fclose(fp);
+ return ret;
+}
+
+cuobjdumpSection* findelfsection(std::list<cuobjdumpSection> sectionlist, unsigned selected_capability, const char* identifier){
+
+ std::list<cuobjdumpSection>::iterator iter;
+ for ( iter = sectionlist.begin();
+ iter != sectionlist.end();
+ iter++
+ ){
+ unsigned capability = 0;
+ sscanf(iter->arch,"sm_%u", &capability);
+ if(capability <= selected_capability &&
+ strcmp(identifier, iter->identifier)==0 &&
+ iter->type == ELFSECTION)
+ return &(*iter);
+ }
+ return NULL;
+}
+
+void useCuobjdump() {
+
+ CUctx_st *context = GPGPUSim_Context();
+ unsigned source_num=1;
+ char *sass, *elf;
+ extract_ptx(); //extract all the output of cuobjdump to _cuobjdump_*.*
+ unsigned selected_capability = get_best_version(cuobjdump, context); //Find max capability less than forced_max_capability
+ unsigned total_ptx_files = get_number_of_ptx(cuobjdump,selected_capability); //Count ptx files for the given capability
+
+ for ( std::list<cuobjdumpSection>::iterator iter = cuobjdump.begin();
+ iter != cuobjdump.end();
+ iter++
+ ){
+ unsigned capability = 0;
+ sscanf(iter->arch,"sm_%u", &capability);
+ if((capability <= selected_capability) && (iter->type == PTXSECTION)){
+ symbol_table *symtab;
+ char *ptxcode = readfile(iter->ptxfilename);
+ if(context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() ) {
+ cuobjdumpSection* elfsection = findelfsection(cuobjdump, selected_capability, iter->identifier);
+ assert (elfsection!= NULL);
+ //char *elfcode = readfile(elfsection->elffilename);
+ //char *sasscode = readfile(elfsection->sassfilename);
+ char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(
+ iter->ptxfilename,
+ elfsection->elffilename,
+ elfsection->sassfilename);
+ symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxplus_str,source_num);
+ printf("Adding %s with cubin handle %u\n", iter->ptxfilename, source_num);
+ //context->add_binary(symtab,total_ptx_files/*fat_cubin_handle*/);
+ //gpgpu_ptxinfo_load_from_string(ptxcode,total_ptx_files/*source_num*/);
+ context->add_binary(symtab, source_num/*fat_cubin_handle*/);
+ gpgpu_ptxinfo_load_from_string( ptxcode,total_ptx_files-source_num/*source_num*/);
+ delete[] ptxplus_str;
+ } else {
+ symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxcode, source_num/*total_ptx_files*//*source_num*/);
+ context->add_binary(symtab,source_num/*fat_cubin_handle*/);
+ gpgpu_ptxinfo_load_from_string( ptxcode, total_ptx_files-source_num/*source_num */);
+ }
+ source_num++;
+ /*!
+ * The order of files in output of cuobjdump is reverse
+ */
+ load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu());
+ load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu());
+ }
+ }
+}
+
+void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin )
{
#if (CUDART_VERSION < 2010)
- printf("GPGPU-Sim PTX: ERROR ** this version of GPGPU-Sim requires CUDA 2.1 or higher\n");
- exit(1);
+ printf("GPGPU-Sim PTX: ERROR ** this version of GPGPU-Sim requires CUDA 2.1 or higher\n");
+ exit(1);
#endif
+ CUctx_st *context = GPGPUSim_Context();
+ static unsigned next_fat_bin_handle = 1;
+ if(context->get_device()->get_gpgpu()->get_config().use_cuobjdump()) {
+ unsigned fat_cubin_handle = next_fat_bin_handle;
+ next_fat_bin_handle++;
+ printf("GPGPU-Sim PTX: __cudaRegisterFatBinary, fat_cubin_handle = %u\n", fat_cubin_handle);
+ /*!
+ * This function extracting all data from all files in first call
+ * then for next calls, only return the appropriate number
+ */
+ assert(fat_cubin_handle >= 1);
+ if(fat_cubin_handle == 1)
+ useCuobjdump();
- 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++;
- __cudaFatCudaBinary *info = (__cudaFatCudaBinary *)fatCubin;
- assert( info->version >= 3 );
- unsigned num_ptx_versions=0;
- unsigned max_capability=0;
- unsigned selected_capability=0;
- bool found=false;
- unsigned forced_max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability();
- 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 );
- if( forced_max_capability ) {
- if( capability > max_capability && capability <= forced_max_capability ) {
- found = true;
- max_capability=capability;
- selected_capability = num_ptx_versions;
- }
- } else {
- if( capability > max_capability ) {
- found = true;
- max_capability=capability;
- 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()->get_config().convert_to_ptxplus() ) {
- if (info->cubin[selected_capability].cubin ==NULL) {
- printf("GPGPU-Sim PTX: Cannot convert to ptxplus no cubin found, probably because it was compiled using newer version of cuda (>=3.0)\nGPGPU-Sim PTX: Exiting ...\n");
- exit(1);
- }
- char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_to_ptxplus(ptx, info->cubin[selected_capability].cubin, source_num++,
- context->get_device()->get_gpgpu()->get_config().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: warning -- did not find an appropriate PTX in cubin\n");
- }
- return (void**)fat_cubin_handle;
+ return (void**)fat_cubin_handle;
+ } else {
+ static unsigned source_num=1;
+ unsigned fat_cubin_handle = next_fat_bin_handle++;
+ __cudaFatCudaBinary *info = (__cudaFatCudaBinary *)fatCubin;
+ assert( info->version >= 3 );
+ unsigned num_ptx_versions=0;
+ unsigned max_capability=0;
+ unsigned selected_capability=0;
+ bool found=false;
+ unsigned forced_max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability();
+ 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 );
+ if( forced_max_capability ) {
+ if( capability > max_capability && capability <= forced_max_capability ) {
+ found = true;
+ max_capability=capability;
+ selected_capability = num_ptx_versions;
+ }
+ } else {
+ if( capability > max_capability ) {
+ found = true;
+ max_capability=capability;
+ 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()->get_config().convert_to_ptxplus() ) {
+ if (info->cubin[selected_capability].cubin ==NULL) {
+ printf("GPGPU-Sim PTX: Cannot convert to ptxplus no cubin found, probably because it was compiled using newer version of cuda (>=3.0)\nGPGPU-Sim PTX: Exiting ...\n");
+ exit(1);
+ }
+ char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_to_ptxplus(ptx, info->cubin[selected_capability].cubin, source_num++,
+ context->get_device()->get_gpgpu()->get_config().saved_converted_ptxplus(), false);
+ 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: warning -- did not find an appropriate PTX in cubin\n");
+ }
+ return (void**)fat_cubin_handle;
+ }
}
void __cudaUnregisterFatBinary(void **fatCubinHandle)
{
- ;
+ ;
+}
+
+cudaError_t cudaDeviceReset ( void ) {
+ // Should reset the simulated GPU
+ return g_last_cudaError = cudaSuccess;
+}
+cudaError_t CUDARTAPI cudaDeviceSynchronize(void){
+ // I don't know what this should do
+ return g_last_cudaError = cudaSuccess;
}
void CUDARTAPI __cudaRegisterFunction(
- void **fatCubinHandle,
- const char *hostFun,
- char *deviceFun,
- const char *deviceName,
- int thread_limit,
- uint3 *tid,
- uint3 *bid,
- dim3 *bDim,
- dim3 *gDim
- )
+ void **fatCubinHandle,
+ const char *hostFun,
+ char *deviceFun,
+ const char *deviceName,
+ int thread_limit,
+ uint3 *tid,
+ uint3 *bid,
+ dim3 *bDim,
+ dim3 *gDim
+)
{
- 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 );
+ 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(
- void **fatCubinHandle,
- char *hostVar, //pointer to...something
- char *deviceAddress, //name of variable
- const char *deviceName, //name of variable (same as above)
- int ext,
- int size,
- int constant,
- int global )
+ void **fatCubinHandle,
+ char *hostVar, //pointer to...something
+ char *deviceAddress, //name of variable
+ const char *deviceName, //name of variable (same as above)
+ int ext,
+ int size,
+ int constant,
+ int global )
{
- printf("GPGPU-Sim PTX: __cudaRegisterVar: hostVar = %p; deviceAddress = %s; deviceName = %s\n", hostVar, deviceAddress, deviceName);
- printf("GPGPU-Sim PTX: __cudaRegisterVar: Registering const memory space of %d bytes\n", size);
- fflush(stdout);
- if ( constant && !global && !ext ) {
- gpgpu_ptx_sim_register_const_variable(hostVar,deviceName,size);
- } else if ( !constant && !global && !ext ) {
- gpgpu_ptx_sim_register_global_variable(hostVar,deviceName,size);
- } else cuda_not_implemented(__my_func__,__LINE__);
+ printf("GPGPU-Sim PTX: __cudaRegisterVar: hostVar = %p; deviceAddress = %s; deviceName = %s\n", hostVar, deviceAddress, deviceName);
+ printf("GPGPU-Sim PTX: __cudaRegisterVar: Registering const memory space of %d bytes\n", size);
+ fflush(stdout);
+ if ( constant && !global && !ext ) {
+ gpgpu_ptx_sim_register_const_variable(hostVar,deviceName,size);
+ } else if ( !constant && !global && !ext ) {
+ gpgpu_ptx_sim_register_global_variable(hostVar,deviceName,size);
+ } else cuda_not_implemented(__my_func__,__LINE__);
}
void __cudaRegisterShared(
- void **fatCubinHandle,
- void **devicePtr
- )
+ void **fatCubinHandle,
+ void **devicePtr
+)
{
- // we don't do anything here
- printf("GPGPU-Sim PTX: __cudaRegisterShared\n" );
+ // we don't do anything here
+ printf("GPGPU-Sim PTX: __cudaRegisterShared\n" );
}
void CUDARTAPI __cudaRegisterSharedVar(
- void **fatCubinHandle,
- void **devicePtr,
- size_t size,
- size_t alignment,
- int storage
- )
+ void **fatCubinHandle,
+ void **devicePtr,
+ size_t size,
+ size_t alignment,
+ int storage
+)
{
- // we don't do anything here
- printf("GPGPU-Sim PTX: __cudaRegisterSharedVar\n" );
+ // we don't do anything here
+ printf("GPGPU-Sim PTX: __cudaRegisterSharedVar\n" );
}
void __cudaRegisterTexture(
- void **fatCubinHandle,
- const struct textureReference *hostVar,
- const void **deviceAddress,
- const char *deviceName,
- int dim,
- int norm,
- int ext
- ) //passes in a newly created textureReference
+ void **fatCubinHandle,
+ const struct textureReference *hostVar,
+ const void **deviceAddress,
+ const char *deviceName,
+ int dim,
+ int norm,
+ int ext
+) //passes in a newly created textureReference
{
- CUctx_st *context = GPGPUSim_Context();
- gpgpu_t *gpu = context->get_device()->get_gpgpu();
- printf("GPGPU-Sim PTX: in __cudaRegisterTexture:\n");
- gpu->gpgpu_ptx_sim_bindNameToTexture(deviceName, hostVar);
- printf("GPGPU-Sim PTX: int dim = %d\n", dim);
- printf("GPGPU-Sim PTX: int norm = %d\n", norm);
- printf("GPGPU-Sim PTX: int ext = %d\n", ext);
- printf("GPGPU-Sim PTX: Execution warning: Not finished implementing \"%s\"\n", __my_func__ );
+ CUctx_st *context = GPGPUSim_Context();
+ gpgpu_t *gpu = context->get_device()->get_gpgpu();
+ printf("GPGPU-Sim PTX: in __cudaRegisterTexture:\n");
+ gpu->gpgpu_ptx_sim_bindNameToTexture(deviceName, hostVar);
+ printf("GPGPU-Sim PTX: int dim = %d\n", dim);
+ printf("GPGPU-Sim PTX: int norm = %d\n", norm);
+ printf("GPGPU-Sim PTX: int ext = %d\n", ext);
+ printf("GPGPU-Sim PTX: Execution warning: Not finished implementing \"%s\"\n", __my_func__ );
}
#ifndef OPENGL_SUPPORT
@@ -1260,15 +1542,15 @@ typedef unsigned long GLuint;
cudaError_t cudaGLRegisterBufferObject(GLuint bufferObj)
{
- printf("GPGPU-Sim PTX: Execution warning: ignoring call to \"%s\"\n", __my_func__ );
- return g_last_cudaError = cudaSuccess;
+ printf("GPGPU-Sim PTX: Execution warning: ignoring call to \"%s\"\n", __my_func__ );
+ return g_last_cudaError = cudaSuccess;
}
struct glbmap_entry {
- GLuint m_bufferObj;
- void *m_devPtr;
- size_t m_size;
- struct glbmap_entry *m_next;
+ GLuint m_bufferObj;
+ void *m_devPtr;
+ size_t m_size;
+ struct glbmap_entry *m_next;
};
typedef struct glbmap_entry glbmap_entry_t;
@@ -1277,182 +1559,182 @@ glbmap_entry_t* g_glbmap = NULL;
cudaError_t cudaGLMapBufferObject(void** devPtr, GLuint bufferObj)
{
#ifdef OPENGL_SUPPORT
- GLint buffer_size=0;
- CUctx_st* ctx = GPGPUSim_Context();
+ GLint buffer_size=0;
+ CUctx_st* ctx = GPGPUSim_Context();
- glbmap_entry_t *p = g_glbmap;
- while ( p && p->m_bufferObj != bufferObj )
- p = p->m_next;
- if ( p == NULL ) {
- glBindBuffer(GL_ARRAY_BUFFER,bufferObj);
- glGetBufferParameteriv(GL_ARRAY_BUFFER,GL_BUFFER_SIZE,&buffer_size);
- assert( buffer_size != 0 );
- *devPtr = ctx->get_device()->get_gpgpu()->gpu_malloc(buffer_size);
+ glbmap_entry_t *p = g_glbmap;
+ while ( p && p->m_bufferObj != bufferObj )
+ p = p->m_next;
+ if ( p == NULL ) {
+ glBindBuffer(GL_ARRAY_BUFFER,bufferObj);
+ glGetBufferParameteriv(GL_ARRAY_BUFFER,GL_BUFFER_SIZE,&buffer_size);
+ assert( buffer_size != 0 );
+ *devPtr = ctx->get_device()->get_gpgpu()->gpu_malloc(buffer_size);
- // create entry and insert to front of list
- glbmap_entry_t *n = (glbmap_entry_t *) calloc(1,sizeof(glbmap_entry_t));
- n->m_next = g_glbmap;
- g_glbmap = n;
+ // create entry and insert to front of list
+ glbmap_entry_t *n = (glbmap_entry_t *) calloc(1,sizeof(glbmap_entry_t));
+ n->m_next = g_glbmap;
+ g_glbmap = n;
- // initialize entry
- n->m_bufferObj = bufferObj;
- n->m_devPtr = *devPtr;
- n->m_size = buffer_size;
+ // initialize entry
+ n->m_bufferObj = bufferObj;
+ n->m_devPtr = *devPtr;
+ n->m_size = buffer_size;
- p = n;
- } else {
- buffer_size = p->m_size;
- *devPtr = p->m_devPtr;
- }
+ p = n;
+ } else {
+ buffer_size = p->m_size;
+ *devPtr = p->m_devPtr;
+ }
- if ( *devPtr ) {
- char *data = (char *) calloc(p->m_size,1);
- glGetBufferSubData(GL_ARRAY_BUFFER,0,buffer_size,data);
- memcpy_to_gpu( (size_t) *devPtr, data, buffer_size );
- free(data);
- printf("GPGPU-Sim PTX: cudaGLMapBufferObject %zu bytes starting at 0x%llx..\n", (size_t)buffer_size,
- (unsigned long long) *devPtr);
- return g_last_cudaError = cudaSuccess;
- } else {
- return g_last_cudaError = cudaErrorMemoryAllocation;
- }
+ if ( *devPtr ) {
+ char *data = (char *) calloc(p->m_size,1);
+ glGetBufferSubData(GL_ARRAY_BUFFER,0,buffer_size,data);
+ memcpy_to_gpu( (size_t) *devPtr, data, buffer_size );
+ free(data);
+ printf("GPGPU-Sim PTX: cudaGLMapBufferObject %zu bytes starting at 0x%llx..\n", (size_t)buffer_size,
+ (unsigned long long) *devPtr);
+ return g_last_cudaError = cudaSuccess;
+ } else {
+ return g_last_cudaError = cudaErrorMemoryAllocation;
+ }
- return g_last_cudaError = cudaSuccess;
+ return g_last_cudaError = cudaSuccess;
#else
- fflush(stdout);
- fflush(stderr);
- printf("GPGPU-Sim PTX: GPGPU-Sim support for OpenGL integration disabled -- exiting\n");
- fflush(stdout);
- exit(50);
+ fflush(stdout);
+ fflush(stderr);
+ printf("GPGPU-Sim PTX: GPGPU-Sim support for OpenGL integration disabled -- exiting\n");
+ fflush(stdout);
+ exit(50);
#endif
}
cudaError_t cudaGLUnmapBufferObject(GLuint bufferObj)
{
#ifdef OPENGL_SUPPORT
- glbmap_entry_t *p = g_glbmap;
- while ( p && p->m_bufferObj != bufferObj )
- p = p->m_next;
- if ( p == NULL )
- return g_last_cudaError = cudaErrorUnknown;
+ glbmap_entry_t *p = g_glbmap;
+ while ( p && p->m_bufferObj != bufferObj )
+ p = p->m_next;
+ if ( p == NULL )
+ return g_last_cudaError = cudaErrorUnknown;
- char *data = (char *) calloc(p->m_size,1);
- memcpy_from_gpu( data,(size_t)p->m_devPtr,p->m_size );
- glBufferSubData(GL_ARRAY_BUFFER,0,p->m_size,data);
- free(data);
+ char *data = (char *) calloc(p->m_size,1);
+ memcpy_from_gpu( data,(size_t)p->m_devPtr,p->m_size );
+ glBufferSubData(GL_ARRAY_BUFFER,0,p->m_size,data);
+ free(data);
- return g_last_cudaError = cudaSuccess;
+ return g_last_cudaError = cudaSuccess;
#else
- fflush(stdout);
- fflush(stderr);
- printf("GPGPU-Sim PTX: support for OpenGL integration disabled -- exiting\n");
- fflush(stdout);
- exit(50);
+ fflush(stdout);
+ fflush(stderr);
+ printf("GPGPU-Sim PTX: support for OpenGL integration disabled -- exiting\n");
+ fflush(stdout);
+ exit(50);
#endif
}
cudaError_t cudaGLUnregisterBufferObject(GLuint bufferObj)
{
- printf("GPGPU-Sim PTX: Execution warning: ignoring call to \"%s\"\n", __my_func__ );
- return g_last_cudaError = cudaSuccess;
+ printf("GPGPU-Sim PTX: Execution warning: ignoring call to \"%s\"\n", __my_func__ );
+ return g_last_cudaError = cudaSuccess;
}
#if (CUDART_VERSION >= 2010)
cudaError_t CUDARTAPI cudaHostAlloc(void **pHost, size_t bytes, unsigned int flags)
{
- *pHost = malloc(bytes);
- if( *pHost )
- return g_last_cudaError = cudaSuccess;
- else
- return g_last_cudaError = cudaErrorMemoryAllocation;
+ *pHost = malloc(bytes);
+ if( *pHost )
+ return g_last_cudaError = cudaSuccess;
+ else
+ return g_last_cudaError = cudaErrorMemoryAllocation;
}
cudaError_t CUDARTAPI cudaHostGetDevicePointer(void **pDevice, void *pHost, unsigned int flags)
{
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaErrorUnknown;
+ cuda_not_implemented(__my_func__,__LINE__);
+ return g_last_cudaError = cudaErrorUnknown;
}
cudaError_t CUDARTAPI cudaSetValidDevices(int *device_arr, int len)
{
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaErrorUnknown;
+ cuda_not_implemented(__my_func__,__LINE__);
+ return g_last_cudaError = cudaErrorUnknown;
}
cudaError_t CUDARTAPI cudaSetDeviceFlags( int flags )
{
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaErrorUnknown;
+ cuda_not_implemented(__my_func__,__LINE__);
+ return g_last_cudaError = cudaErrorUnknown;
}
cudaError_t CUDARTAPI cudaFuncGetAttributes(struct cudaFuncAttributes *attr, const char *hostFun )
{
- CUctx_st *context = GPGPUSim_Context();
- function_info *entry = context->get_kernel(hostFun);
- if( entry ) {
- const struct gpgpu_ptx_sim_kernel_info *kinfo = entry->get_kernel_info();
- attr->sharedSizeBytes = kinfo->smem;
- attr->constSizeBytes = kinfo->cmem;
- attr->localSizeBytes = kinfo->lmem;
- attr->numRegs = kinfo->regs;
- attr->maxThreadsPerBlock = 0; // from pragmas?
+ CUctx_st *context = GPGPUSim_Context();
+ function_info *entry = context->get_kernel(hostFun);
+ if( entry ) {
+ const struct gpgpu_ptx_sim_kernel_info *kinfo = entry->get_kernel_info();
+ attr->sharedSizeBytes = kinfo->smem;
+ attr->constSizeBytes = kinfo->cmem;
+ attr->localSizeBytes = kinfo->lmem;
+ attr->numRegs = kinfo->regs;
+ attr->maxThreadsPerBlock = 0; // from pragmas?
#if CUDART_VERSION >= 3000
- attr->ptxVersion = kinfo->ptx_version;
- attr->binaryVersion = kinfo->sm_target;
+ attr->ptxVersion = kinfo->ptx_version;
+ attr->binaryVersion = kinfo->sm_target;
#endif
- }
- return g_last_cudaError = cudaSuccess;
+ }
+ return g_last_cudaError = cudaSuccess;
}
cudaError_t CUDARTAPI cudaEventCreateWithFlags(cudaEvent_t *event, int flags)
{
- CUevent_st *e = new CUevent_st(flags==cudaEventBlockingSync);
- g_timer_events[e->get_uid()] = e;
+ CUevent_st *e = new CUevent_st(flags==cudaEventBlockingSync);
+ g_timer_events[e->get_uid()] = e;
#if CUDART_VERSION >= 3000
- *event = e;
+ *event = e;
#else
- *event = e->get_uid();
+ *event = e->get_uid();
#endif
- return g_last_cudaError = cudaSuccess;
+ return g_last_cudaError = cudaSuccess;
}
cudaError_t CUDARTAPI cudaDriverGetVersion(int *driverVersion)
{
- *driverVersion = CUDART_VERSION;
- return g_last_cudaError = cudaErrorUnknown;
+ *driverVersion = CUDART_VERSION;
+ return g_last_cudaError = cudaErrorUnknown;
}
cudaError_t CUDARTAPI cudaRuntimeGetVersion(int *runtimeVersion)
{
- *runtimeVersion = CUDART_VERSION;
- return g_last_cudaError = cudaErrorUnknown;
+ *runtimeVersion = CUDART_VERSION;
+ return g_last_cudaError = cudaErrorUnknown;
}
#endif
cudaError_t CUDARTAPI cudaGLSetGLDevice(int device)
{
- printf("GPGPU-Sim PTX: Execution warning: ignoring call to \"%s\"\n", __my_func__ );
- return g_last_cudaError = cudaErrorUnknown;
+ printf("GPGPU-Sim PTX: Execution warning: ignoring call to \"%s\"\n", __my_func__ );
+ return g_last_cudaError = cudaErrorUnknown;
}
typedef void* HGPUNV;
cudaError_t CUDARTAPI cudaWGLGetDevice(int *device, HGPUNV hGpu)
{
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaErrorUnknown;
+ cuda_not_implemented(__my_func__,__LINE__);
+ return g_last_cudaError = cudaErrorUnknown;
}
void CUDARTAPI __cudaMutexOperation(int lock)
{
- cuda_not_implemented(__my_func__,__LINE__);
+ cuda_not_implemented(__my_func__,__LINE__);
}
void CUDARTAPI __cudaTextureFetch(const void *tex, void *index, int integer, void *val)
{
- cuda_not_implemented(__my_func__,__LINE__);
+ cuda_not_implemented(__my_func__,__LINE__);
}
}
@@ -1461,18 +1743,18 @@ namespace cuda_math {
void CUDARTAPI __cudaMutexOperation(int lock)
{
- cuda_not_implemented(__my_func__,__LINE__);
+ cuda_not_implemented(__my_func__,__LINE__);
}
void CUDARTAPI __cudaTextureFetch(const void *tex, void *index, int integer, void *val)
{
- cuda_not_implemented(__my_func__,__LINE__);
+ cuda_not_implemented(__my_func__,__LINE__);
}
int CUDARTAPI __cudaSynchronizeThreads(void**, void*)
{
- //TODO This function should syncronize if we support Asyn kernel calls
- return g_last_cudaError = cudaSuccess;
+ //TODO This function should syncronize if we support Asyn kernel calls
+ return g_last_cudaError = cudaSuccess;
}
}
@@ -1492,106 +1774,106 @@ extern "C" FILE *ptxinfo_in;
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);
- int ng_bytes=0;
- symbol_table::iterator g=symtab->global_iterator_begin();
+ printf( "GPGPU-Sim PTX: loading globals with explicit initializers... \n" );
+ fflush(stdout);
+ int ng_bytes=0;
+ symbol_table::iterator g=symtab->global_iterator_begin();
- for ( ; g!=symtab->global_iterator_end(); g++) {
- symbol *global = *g;
- if ( global->has_initializer() ) {
- printf( "GPGPU-Sim PTX: initializing '%s' ... ", global->name().c_str() );
- unsigned addr=global->get_address();
- const type_info *type = global->type();
- type_info_key ti=type->get_key();
- size_t size;
- int t;
- ti.type_decode(size,t);
- int nbytes = size/8;
- int offset=0;
- std::list<operand_info> init_list = global->get_initializer();
- for ( std::list<operand_info>::iterator i=init_list.begin(); i!=init_list.end(); i++ ) {
- 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
- gpu->get_global_memory()->write(addr+offset,nbytes,&value,NULL,NULL); // assuming little endian here
- offset+=nbytes;
- ng_bytes+=nbytes;
- }
- printf(" wrote %u bytes\n", offset );
- }
- }
- printf( "GPGPU-Sim PTX: finished loading globals (%u bytes total).\n", ng_bytes );
- fflush(stdout);
- return ng_bytes;
+ for ( ; g!=symtab->global_iterator_end(); g++) {
+ symbol *global = *g;
+ if ( global->has_initializer() ) {
+ printf( "GPGPU-Sim PTX: initializing '%s' ... ", global->name().c_str() );
+ unsigned addr=global->get_address();
+ const type_info *type = global->type();
+ type_info_key ti=type->get_key();
+ size_t size;
+ int t;
+ ti.type_decode(size,t);
+ int nbytes = size/8;
+ int offset=0;
+ std::list<operand_info> init_list = global->get_initializer();
+ for ( std::list<operand_info>::iterator i=init_list.begin(); i!=init_list.end(); i++ ) {
+ 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
+ gpu->get_global_memory()->write(addr+offset,nbytes,&value,NULL,NULL); // assuming little endian here
+ offset+=nbytes;
+ ng_bytes+=nbytes;
+ }
+ printf(" wrote %u bytes\n", offset );
+ }
+ }
+ printf( "GPGPU-Sim PTX: finished loading globals (%u bytes total).\n", ng_bytes );
+ fflush(stdout);
+ return ng_bytes;
}
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);
- int nc_bytes = 0;
- symbol_table::iterator g=symtab->const_iterator_begin();
+ printf( "GPGPU-Sim PTX: loading constants with explicit initializers... " );
+ fflush(stdout);
+ int nc_bytes = 0;
+ symbol_table::iterator g=symtab->const_iterator_begin();
- for ( ; g!=symtab->const_iterator_end(); g++) {
- symbol *constant = *g;
- if ( constant->is_const() && constant->has_initializer() ) {
+ for ( ; g!=symtab->const_iterator_end(); g++) {
+ symbol *constant = *g;
+ if ( constant->is_const() && constant->has_initializer() ) {
- // get the constant element data size
- int basic_type;
- size_t num_bits;
- constant->type()->get_key().type_decode(num_bits,basic_type);
+ // get the constant element data size
+ int basic_type;
+ size_t num_bits;
+ constant->type()->get_key().type_decode(num_bits,basic_type);
- std::list<operand_info> init_list = constant->get_initializer();
- int nbytes_written = 0;
- for ( std::list<operand_info>::iterator i=init_list.begin(); i!=init_list.end(); i++ ) {
- operand_info op = *i;
- ptx_reg_t value = op.get_literal_value();
- int nbytes = num_bits/8;
- switch ( op.get_type() ) {
- case int_t: assert(nbytes >= 1); break;
- case float_op_t: assert(nbytes == 4); break;
- case double_op_t: assert(nbytes >= 4); break; // account for double DEMOTING
- default:
- abort();
- }
- unsigned addr=constant->get_address() + nbytes_written;
- assert( addr+nbytes < min_gaddr );
+ std::list<operand_info> init_list = constant->get_initializer();
+ int nbytes_written = 0;
+ for ( std::list<operand_info>::iterator i=init_list.begin(); i!=init_list.end(); i++ ) {
+ operand_info op = *i;
+ ptx_reg_t value = op.get_literal_value();
+ int nbytes = num_bits/8;
+ switch ( op.get_type() ) {
+ case int_t: assert(nbytes >= 1); break;
+ case float_op_t: assert(nbytes == 4); break;
+ case double_op_t: assert(nbytes >= 4); break; // account for double DEMOTING
+ default:
+ abort();
+ }
+ unsigned addr=constant->get_address() + nbytes_written;
+ assert( addr+nbytes < min_gaddr );
- 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;
- }
- }
- }
- printf( " done.\n");
- fflush(stdout);
- return nc_bytes;
+ 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;
+ }
+ }
+ }
+ printf( " done.\n");
+ fflush(stdout);
+ return nc_bytes;
}
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 )
+ gpgpu_ptx_sim_arg_list_t args,
+ struct dim3 gridDim,
+ struct dim3 blockDim,
+ CUctx_st* context )
{
- function_info *entry = context->get_kernel(hostFun);
- kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry);
- if( entry == NULL ) {
- printf("GPGPU-Sim PTX: ERROR launching kernel -- no PTX implementation found\n");
- abort();
- }
- unsigned argcount=args.size();
- unsigned argn=1;
- for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) {
- entry->add_param_data(argcount-argn,&(*a));
- argn++;
- }
+ function_info *entry = context->get_kernel(hostFun);
+ kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry);
+ if( entry == NULL ) {
+ printf("GPGPU-Sim PTX: ERROR launching kernel -- no PTX implementation found for %p\n", hostFun);
+ abort();
+ }
+ unsigned argcount=args.size();
+ unsigned argn=1;
+ for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) {
+ entry->add_param_data(argcount-argn,&(*a));
+ argn++;
+ }
- entry->finalize(result->get_param_memory());
- g_ptx_kernel_count++;
- fflush(stdout);
+ entry->finalize(result->get_param_memory());
+ g_ptx_kernel_count++;
+ fflush(stdout);
- return result;
+ return result;
}
diff --git a/libopencl/Makefile b/libopencl/Makefile
index cc3221b..3502ed3 100644
--- a/libopencl/Makefile
+++ b/libopencl/Makefile
@@ -117,73 +117,74 @@ clean:
rm -f bin/nvopencl_wrapper
# DO NOT DELETE
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/stdlib.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/features.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/sys/cdefs.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/bits/wordsize.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/gnu/stubs.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/gnu/stubs-64.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/sys/types.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/bits/types.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/bits/typesizes.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/time.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/endian.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/bits/endian.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/bits/byteswap.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/sys/select.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/bits/select.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/bits/sigset.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/bits/time.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/sys/sysmacros.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/bits/pthreadtypes.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/alloca.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/stdio.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/libio.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/_G_config.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/wchar.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/bits/stdio_lim.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/bits/sys_errlist.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/string.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/assert.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/abstract_hardware_model.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/cuda-sim/cuda-sim.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/abstract_hardware_model.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/cuda-sim/ptx_loader.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/cuda-sim/ptx_ir.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/cuda-sim/ptx.tab.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/cuda-sim/ptx_sim.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/tr1_hash_map.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/cuda-sim/opcodes.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/cuda-sim/opcodes.def
-../build/debug/libopencl/opencl_runtime_api.o: ../src/cuda-sim/memory.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/option_parser.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/gpgpusim_entrypoint.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/abstract_hardware_model.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/gpgpu-sim/gpu-sim.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/gpgpu-sim/addrdec.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/gpgpu-sim/shader.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/math.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/bits/huge_val.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/bits/mathdef.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/bits/mathcalls.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/cuda-sim/ptx.tab.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/gpgpu-sim/delayqueue.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/intersim/statwraper.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/gpgpu-sim/gpu-misc.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/gpgpu-sim/stack.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/gpgpu-sim/dram.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/zlib.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/zconf.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/unistd.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/bits/posix_opt.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/bits/confname.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/getopt.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/gpgpu-sim/scoreboard.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/gpgpu-sim/mem_fetch.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/gpgpu-sim/stats.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/gpgpu-sim/gpu-cache.h
-../build/debug/libopencl/opencl_runtime_api.o: ../src/gpgpu-sim/shader.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/CL/cl.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/CL/cl_platform.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/stdint.h
-../build/debug/libopencl/opencl_runtime_api.o: /usr/include/bits/wchar.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/stdlib.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/features.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/sys/cdefs.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/bits/wordsize.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/gnu/stubs.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/gnu/stubs-64.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/sys/types.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/bits/types.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/bits/typesizes.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/time.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/endian.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/bits/endian.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/bits/byteswap.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/sys/select.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/bits/select.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/bits/sigset.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/bits/time.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/sys/sysmacros.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/bits/pthreadtypes.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/alloca.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/stdio.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/libio.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/_G_config.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/wchar.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/bits/stdio_lim.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/bits/sys_errlist.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/string.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/xlocale.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/assert.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/abstract_hardware_model.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/cuda-sim/cuda-sim.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/abstract_hardware_model.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/cuda-sim/ptx_loader.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/cuda-sim/ptx_ir.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/cuda-sim/ptx.tab.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/cuda-sim/ptx_sim.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/tr1_hash_map.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/cuda-sim/opcodes.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/cuda-sim/opcodes.def
+../build/release/libopencl/opencl_runtime_api.o: ../src/cuda-sim/memory.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/option_parser.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/gpgpusim_entrypoint.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/abstract_hardware_model.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/gpgpu-sim/gpu-sim.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/gpgpu-sim/addrdec.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/gpgpu-sim/shader.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/math.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/bits/huge_val.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/bits/huge_valf.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/bits/huge_vall.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/bits/inf.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/bits/nan.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/bits/mathdef.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/bits/mathcalls.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/cuda-sim/ptx.tab.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/gpgpu-sim/delayqueue.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/intersim/statwraper.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/gpgpu-sim/gpu-misc.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/gpgpu-sim/stack.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/gpgpu-sim/dram.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/zlib.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/zconf.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/unistd.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/bits/posix_opt.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/bits/confname.h
+../build/release/libopencl/opencl_runtime_api.o: /usr/include/getopt.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/gpgpu-sim/scoreboard.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/gpgpu-sim/mem_fetch.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/gpgpu-sim/stats.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/gpgpu-sim/gpu-cache.h
+../build/release/libopencl/opencl_runtime_api.o: ../src/gpgpu-sim/shader.h
diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc
index 9e476eb..6457edd 100644
--- a/src/abstract_hardware_model.cc
+++ b/src/abstract_hardware_model.cc
@@ -49,9 +49,18 @@ void move_warp( warp_inst_t *&dst, warp_inst_t *&src )
void gpgpu_functional_sim_config::reg_options(class OptionParser * opp)
{
+ option_parser_register(opp, "-gpgpu_ptx_use_cuobjdump", OPT_BOOL,
+ &m_ptx_use_cuobjdump,
+ "Use cuobjdump to extract ptx and sass from binaries",
+#if (CUDART_VERSION >= 4000)
+ "1"
+#else
+ "0"
+#endif
+ );
option_parser_register(opp, "-gpgpu_ptx_convert_to_ptxplus", OPT_BOOL,
&m_ptx_convert_to_ptxplus,
- "Convert embedded ptx to ptxplus",
+ "Convert SASS (native ISA) to ptxplus and run ptxplus",
"0");
option_parser_register(opp, "-gpgpu_ptx_save_converted_ptxplus", OPT_BOOL,
&m_ptx_save_converted_ptxplus,
diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h
index a6cc811..4bfdda9 100644
--- a/src/abstract_hardware_model.h
+++ b/src/abstract_hardware_model.h
@@ -337,6 +337,7 @@ public:
unsigned get_forced_max_capability() const { return m_ptx_force_max_capability; }
bool convert_to_ptxplus() const { return m_ptx_convert_to_ptxplus; }
+ bool use_cuobjdump() const { return m_ptx_use_cuobjdump; }
bool saved_converted_ptxplus() const { return m_ptx_save_converted_ptxplus; }
int get_ptx_inst_debug_to_file() const { return g_ptx_inst_debug_to_file; }
@@ -348,6 +349,7 @@ private:
// PTX options
int m_ptx_convert_to_ptxplus;
int m_ptx_save_converted_ptxplus;
+ int m_ptx_use_cuobjdump;
unsigned m_ptx_force_max_capability;
int g_ptx_inst_debug_to_file;
diff --git a/src/cuda-sim/Makefile b/src/cuda-sim/Makefile
index a475ca2..a70ee5f 100644
--- a/src/cuda-sim/Makefile
+++ b/src/cuda-sim/Makefile
@@ -159,13 +159,15 @@ $(OUTPUT_DIR)/lex.ptx_.o: ptx.tab.c
../../build/release/cuda-sim/cuda_device_printf.o: /usr/include/gnu/stubs-64.h
../../build/release/cuda-sim/cuda_device_printf.o: ptx.tab.h ptx_sim.h
../../build/release/cuda-sim/cuda_device_printf.o: /usr/include/stdlib.h
+../../build/release/cuda-sim/cuda_device_printf.o: /usr/include/bits/waitflags.h
+../../build/release/cuda-sim/cuda_device_printf.o: /usr/include/bits/waitstatus.h
+../../build/release/cuda-sim/cuda_device_printf.o: /usr/include/endian.h
+../../build/release/cuda-sim/cuda_device_printf.o: /usr/include/bits/endian.h
+../../build/release/cuda-sim/cuda_device_printf.o: /usr/include/bits/byteswap.h
../../build/release/cuda-sim/cuda_device_printf.o: /usr/include/sys/types.h
../../build/release/cuda-sim/cuda_device_printf.o: /usr/include/bits/types.h
../../build/release/cuda-sim/cuda_device_printf.o: /usr/include/bits/typesizes.h
../../build/release/cuda-sim/cuda_device_printf.o: /usr/include/time.h
-../../build/release/cuda-sim/cuda_device_printf.o: /usr/include/endian.h
-../../build/release/cuda-sim/cuda_device_printf.o: /usr/include/bits/endian.h
-../../build/release/cuda-sim/cuda_device_printf.o: /usr/include/bits/byteswap.h
../../build/release/cuda-sim/cuda_device_printf.o: /usr/include/sys/select.h
../../build/release/cuda-sim/cuda_device_printf.o: /usr/include/bits/select.h
../../build/release/cuda-sim/cuda_device_printf.o: /usr/include/bits/sigset.h
@@ -193,13 +195,15 @@ $(OUTPUT_DIR)/lex.ptx_.o: ptx.tab.c
../../build/release/cuda-sim/cuda-sim.o: /usr/include/bits/wordsize.h
../../build/release/cuda-sim/cuda-sim.o: /usr/include/gnu/stubs.h
../../build/release/cuda-sim/cuda-sim.o: /usr/include/gnu/stubs-64.h
+../../build/release/cuda-sim/cuda-sim.o: /usr/include/bits/waitflags.h
+../../build/release/cuda-sim/cuda-sim.o: /usr/include/bits/waitstatus.h
+../../build/release/cuda-sim/cuda-sim.o: /usr/include/endian.h
+../../build/release/cuda-sim/cuda-sim.o: /usr/include/bits/endian.h
+../../build/release/cuda-sim/cuda-sim.o: /usr/include/bits/byteswap.h
../../build/release/cuda-sim/cuda-sim.o: /usr/include/sys/types.h
../../build/release/cuda-sim/cuda-sim.o: /usr/include/bits/types.h
../../build/release/cuda-sim/cuda-sim.o: /usr/include/bits/typesizes.h
../../build/release/cuda-sim/cuda-sim.o: /usr/include/time.h
-../../build/release/cuda-sim/cuda-sim.o: /usr/include/endian.h
-../../build/release/cuda-sim/cuda-sim.o: /usr/include/bits/endian.h
-../../build/release/cuda-sim/cuda-sim.o: /usr/include/bits/byteswap.h
../../build/release/cuda-sim/cuda-sim.o: /usr/include/sys/select.h
../../build/release/cuda-sim/cuda-sim.o: /usr/include/bits/select.h
../../build/release/cuda-sim/cuda-sim.o: /usr/include/bits/sigset.h
@@ -242,6 +246,7 @@ $(OUTPUT_DIR)/lex.ptx_.o: ptx.tab.c
../../build/release/cuda-sim/cuda-sim.o: /usr/include/zconf.h
../../build/release/cuda-sim/cuda-sim.o: /usr/include/unistd.h
../../build/release/cuda-sim/cuda-sim.o: /usr/include/bits/posix_opt.h
+../../build/release/cuda-sim/cuda-sim.o: /usr/include/bits/environments.h
../../build/release/cuda-sim/cuda-sim.o: /usr/include/bits/confname.h
../../build/release/cuda-sim/cuda-sim.o: /usr/include/getopt.h
../../build/release/cuda-sim/cuda-sim.o: ../gpgpu-sim/scoreboard.h
@@ -261,13 +266,15 @@ $(OUTPUT_DIR)/lex.ptx_.o: ptx.tab.c
../../build/release/cuda-sim/instructions.o: /usr/include/gnu/stubs-64.h
../../build/release/cuda-sim/instructions.o: ptx.tab.h ptx_sim.h
../../build/release/cuda-sim/instructions.o: /usr/include/stdlib.h
+../../build/release/cuda-sim/instructions.o: /usr/include/bits/waitflags.h
+../../build/release/cuda-sim/instructions.o: /usr/include/bits/waitstatus.h
+../../build/release/cuda-sim/instructions.o: /usr/include/endian.h
+../../build/release/cuda-sim/instructions.o: /usr/include/bits/endian.h
+../../build/release/cuda-sim/instructions.o: /usr/include/bits/byteswap.h
../../build/release/cuda-sim/instructions.o: /usr/include/sys/types.h
../../build/release/cuda-sim/instructions.o: /usr/include/bits/types.h
../../build/release/cuda-sim/instructions.o: /usr/include/bits/typesizes.h
../../build/release/cuda-sim/instructions.o: /usr/include/time.h
-../../build/release/cuda-sim/instructions.o: /usr/include/endian.h
-../../build/release/cuda-sim/instructions.o: /usr/include/bits/endian.h
-../../build/release/cuda-sim/instructions.o: /usr/include/bits/byteswap.h
../../build/release/cuda-sim/instructions.o: /usr/include/sys/select.h
../../build/release/cuda-sim/instructions.o: /usr/include/bits/select.h
../../build/release/cuda-sim/instructions.o: /usr/include/bits/sigset.h
@@ -311,6 +318,7 @@ $(OUTPUT_DIR)/lex.ptx_.o: ptx.tab.c
../../build/release/cuda-sim/instructions.o: /usr/include/zconf.h
../../build/release/cuda-sim/instructions.o: /usr/include/unistd.h
../../build/release/cuda-sim/instructions.o: /usr/include/bits/posix_opt.h
+../../build/release/cuda-sim/instructions.o: /usr/include/bits/environments.h
../../build/release/cuda-sim/instructions.o: /usr/include/bits/confname.h
../../build/release/cuda-sim/instructions.o: /usr/include/getopt.h
../../build/release/cuda-sim/instructions.o: ../gpgpu-sim/scoreboard.h
@@ -336,11 +344,13 @@ $(OUTPUT_DIR)/lex.ptx_.o: ptx.tab.c
../../build/release/cuda-sim/memory.o: /usr/include/bits/stdio_lim.h
../../build/release/cuda-sim/memory.o: /usr/include/bits/sys_errlist.h
../../build/release/cuda-sim/memory.o: /usr/include/stdlib.h
-../../build/release/cuda-sim/memory.o: /usr/include/sys/types.h
-../../build/release/cuda-sim/memory.o: /usr/include/time.h
+../../build/release/cuda-sim/memory.o: /usr/include/bits/waitflags.h
+../../build/release/cuda-sim/memory.o: /usr/include/bits/waitstatus.h
../../build/release/cuda-sim/memory.o: /usr/include/endian.h
../../build/release/cuda-sim/memory.o: /usr/include/bits/endian.h
../../build/release/cuda-sim/memory.o: /usr/include/bits/byteswap.h
+../../build/release/cuda-sim/memory.o: /usr/include/sys/types.h
+../../build/release/cuda-sim/memory.o: /usr/include/time.h
../../build/release/cuda-sim/memory.o: /usr/include/sys/select.h
../../build/release/cuda-sim/memory.o: /usr/include/bits/select.h
../../build/release/cuda-sim/memory.o: /usr/include/bits/sigset.h
@@ -358,13 +368,15 @@ $(OUTPUT_DIR)/lex.ptx_.o: ptx.tab.c
../../build/release/cuda-sim/ptx_ir.o: /usr/include/gnu/stubs.h
../../build/release/cuda-sim/ptx_ir.o: /usr/include/gnu/stubs-64.h ptx.tab.h
../../build/release/cuda-sim/ptx_ir.o: ptx_sim.h /usr/include/stdlib.h
+../../build/release/cuda-sim/ptx_ir.o: /usr/include/bits/waitflags.h
+../../build/release/cuda-sim/ptx_ir.o: /usr/include/bits/waitstatus.h
+../../build/release/cuda-sim/ptx_ir.o: /usr/include/endian.h
+../../build/release/cuda-sim/ptx_ir.o: /usr/include/bits/endian.h
+../../build/release/cuda-sim/ptx_ir.o: /usr/include/bits/byteswap.h
../../build/release/cuda-sim/ptx_ir.o: /usr/include/sys/types.h
../../build/release/cuda-sim/ptx_ir.o: /usr/include/bits/types.h
../../build/release/cuda-sim/ptx_ir.o: /usr/include/bits/typesizes.h
../../build/release/cuda-sim/ptx_ir.o: /usr/include/time.h
-../../build/release/cuda-sim/ptx_ir.o: /usr/include/endian.h
-../../build/release/cuda-sim/ptx_ir.o: /usr/include/bits/endian.h
-../../build/release/cuda-sim/ptx_ir.o: /usr/include/bits/byteswap.h
../../build/release/cuda-sim/ptx_ir.o: /usr/include/sys/select.h
../../build/release/cuda-sim/ptx_ir.o: /usr/include/bits/select.h
../../build/release/cuda-sim/ptx_ir.o: /usr/include/bits/sigset.h
@@ -393,13 +405,15 @@ $(OUTPUT_DIR)/lex.ptx_.o: ptx.tab.c
../../build/release/cuda-sim/ptx_loader.o: /usr/include/gnu/stubs-64.h
../../build/release/cuda-sim/ptx_loader.o: ptx.tab.h ptx_sim.h
../../build/release/cuda-sim/ptx_loader.o: /usr/include/stdlib.h
+../../build/release/cuda-sim/ptx_loader.o: /usr/include/bits/waitflags.h
+../../build/release/cuda-sim/ptx_loader.o: /usr/include/bits/waitstatus.h
+../../build/release/cuda-sim/ptx_loader.o: /usr/include/endian.h
+../../build/release/cuda-sim/ptx_loader.o: /usr/include/bits/endian.h
+../../build/release/cuda-sim/ptx_loader.o: /usr/include/bits/byteswap.h
../../build/release/cuda-sim/ptx_loader.o: /usr/include/sys/types.h
../../build/release/cuda-sim/ptx_loader.o: /usr/include/bits/types.h
../../build/release/cuda-sim/ptx_loader.o: /usr/include/bits/typesizes.h
../../build/release/cuda-sim/ptx_loader.o: /usr/include/time.h
-../../build/release/cuda-sim/ptx_loader.o: /usr/include/endian.h
-../../build/release/cuda-sim/ptx_loader.o: /usr/include/bits/endian.h
-../../build/release/cuda-sim/ptx_loader.o: /usr/include/bits/byteswap.h
../../build/release/cuda-sim/ptx_loader.o: /usr/include/sys/select.h
../../build/release/cuda-sim/ptx_loader.o: /usr/include/bits/select.h
../../build/release/cuda-sim/ptx_loader.o: /usr/include/bits/sigset.h
@@ -433,13 +447,15 @@ $(OUTPUT_DIR)/lex.ptx_.o: ptx.tab.c
../../build/release/cuda-sim/ptx_parser.o: /usr/include/gnu/stubs-64.h
../../build/release/cuda-sim/ptx_parser.o: ptx.tab.h ptx_sim.h
../../build/release/cuda-sim/ptx_parser.o: /usr/include/stdlib.h
+../../build/release/cuda-sim/ptx_parser.o: /usr/include/bits/waitflags.h
+../../build/release/cuda-sim/ptx_parser.o: /usr/include/bits/waitstatus.h
+../../build/release/cuda-sim/ptx_parser.o: /usr/include/endian.h
+../../build/release/cuda-sim/ptx_parser.o: /usr/include/bits/endian.h
+../../build/release/cuda-sim/ptx_parser.o: /usr/include/bits/byteswap.h
../../build/release/cuda-sim/ptx_parser.o: /usr/include/sys/types.h
../../build/release/cuda-sim/ptx_parser.o: /usr/include/bits/types.h
../../build/release/cuda-sim/ptx_parser.o: /usr/include/bits/typesizes.h
../../build/release/cuda-sim/ptx_parser.o: /usr/include/time.h
-../../build/release/cuda-sim/ptx_parser.o: /usr/include/endian.h
-../../build/release/cuda-sim/ptx_parser.o: /usr/include/bits/endian.h
-../../build/release/cuda-sim/ptx_parser.o: /usr/include/bits/byteswap.h
../../build/release/cuda-sim/ptx_parser.o: /usr/include/sys/select.h
../../build/release/cuda-sim/ptx_parser.o: /usr/include/bits/select.h
../../build/release/cuda-sim/ptx_parser.o: /usr/include/bits/sigset.h
@@ -465,13 +481,15 @@ $(OUTPUT_DIR)/lex.ptx_.o: ptx.tab.c
../../build/release/cuda-sim/ptx_sim.o: /usr/include/bits/wordsize.h
../../build/release/cuda-sim/ptx_sim.o: /usr/include/gnu/stubs.h
../../build/release/cuda-sim/ptx_sim.o: /usr/include/gnu/stubs-64.h
+../../build/release/cuda-sim/ptx_sim.o: /usr/include/bits/waitflags.h
+../../build/release/cuda-sim/ptx_sim.o: /usr/include/bits/waitstatus.h
+../../build/release/cuda-sim/ptx_sim.o: /usr/include/endian.h
+../../build/release/cuda-sim/ptx_sim.o: /usr/include/bits/endian.h
+../../build/release/cuda-sim/ptx_sim.o: /usr/include/bits/byteswap.h
../../build/release/cuda-sim/ptx_sim.o: /usr/include/sys/types.h
../../build/release/cuda-sim/ptx_sim.o: /usr/include/bits/types.h
../../build/release/cuda-sim/ptx_sim.o: /usr/include/bits/typesizes.h
../../build/release/cuda-sim/ptx_sim.o: /usr/include/time.h
-../../build/release/cuda-sim/ptx_sim.o: /usr/include/endian.h
-../../build/release/cuda-sim/ptx_sim.o: /usr/include/bits/endian.h
-../../build/release/cuda-sim/ptx_sim.o: /usr/include/bits/byteswap.h
../../build/release/cuda-sim/ptx_sim.o: /usr/include/sys/select.h
../../build/release/cuda-sim/ptx_sim.o: /usr/include/bits/select.h
../../build/release/cuda-sim/ptx_sim.o: /usr/include/bits/sigset.h
@@ -513,6 +531,7 @@ $(OUTPUT_DIR)/lex.ptx_.o: ptx.tab.c
../../build/release/cuda-sim/ptx_sim.o: /usr/include/zconf.h
../../build/release/cuda-sim/ptx_sim.o: /usr/include/unistd.h
../../build/release/cuda-sim/ptx_sim.o: /usr/include/bits/posix_opt.h
+../../build/release/cuda-sim/ptx_sim.o: /usr/include/bits/environments.h
../../build/release/cuda-sim/ptx_sim.o: /usr/include/bits/confname.h
../../build/release/cuda-sim/ptx_sim.o: /usr/include/getopt.h
../../build/release/cuda-sim/ptx_sim.o: ../gpgpu-sim/scoreboard.h
@@ -530,13 +549,15 @@ $(OUTPUT_DIR)/lex.ptx_.o: ptx.tab.c
../../build/release/cuda-sim/ptx-stats.o: /usr/include/gnu/stubs-64.h
../../build/release/cuda-sim/ptx-stats.o: ptx.tab.h ptx_sim.h
../../build/release/cuda-sim/ptx-stats.o: /usr/include/stdlib.h
+../../build/release/cuda-sim/ptx-stats.o: /usr/include/bits/waitflags.h
+../../build/release/cuda-sim/ptx-stats.o: /usr/include/bits/waitstatus.h
+../../build/release/cuda-sim/ptx-stats.o: /usr/include/endian.h
+../../build/release/cuda-sim/ptx-stats.o: /usr/include/bits/endian.h
+../../build/release/cuda-sim/ptx-stats.o: /usr/include/bits/byteswap.h
../../build/release/cuda-sim/ptx-stats.o: /usr/include/sys/types.h
../../build/release/cuda-sim/ptx-stats.o: /usr/include/bits/types.h
../../build/release/cuda-sim/ptx-stats.o: /usr/include/bits/typesizes.h
../../build/release/cuda-sim/ptx-stats.o: /usr/include/time.h
-../../build/release/cuda-sim/ptx-stats.o: /usr/include/endian.h
-../../build/release/cuda-sim/ptx-stats.o: /usr/include/bits/endian.h
-../../build/release/cuda-sim/ptx-stats.o: /usr/include/bits/byteswap.h
../../build/release/cuda-sim/ptx-stats.o: /usr/include/sys/select.h
../../build/release/cuda-sim/ptx-stats.o: /usr/include/bits/select.h
../../build/release/cuda-sim/ptx-stats.o: /usr/include/bits/sigset.h
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index d8c2c00..7d8d254 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -43,6 +43,7 @@
#include "ptx_loader.h"
#include "ptx_parser.h"
#include "../gpgpu-sim/gpu-sim.h"
+#include "ptx_sim.h"
#include "../gpgpusim_entrypoint.h"
#include "decuda_pred_table/decuda_pred_table.h"
#include "../stream_manager.h"
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index 550a076..4537bd0 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -80,11 +80,12 @@ ptx_reg_t ptx_thread_info::get_reg( const symbol *reg )
const std::string &name = reg->name();
unsigned call_uid = m_callstack.back().m_call_uid;
ptx_reg_t uninit_reg;
- uninit_reg.u32 = 0xDEADBEEF;
+ uninit_reg.u32 = 0x0;
set_reg(reg, uninit_reg); // give it a value since we are going to warn the user anyway
std::string file_loc = get_location();
if( !unfound_register_warned ) {
- printf("GPGPU-Sim PTX: WARNING (%s) ** reading undefined register \'%s\' (cuid:%u). Setting to 0XDEADBEEF.\n",
+ printf("GPGPU-Sim PTX: WARNING (%s) ** reading undefined register \'%s\' (cuid:%u). Setting to 0X00000000. This is okay if you are simulating the native ISA"
+ "\n",
file_loc.c_str(), name.c_str(), call_uid );
unfound_register_warned = true;
}
@@ -818,13 +819,21 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread )
const operand_info &src2 = pI->src2(); // b
// Get operand values
- src1_data = thread->get_operand_value(src1, dst, to_type, thread, 1); // a
- src2_data = thread->get_operand_value(src2, dst, to_type, thread, 1); // b
+ if (dst.get_symbol()->type()){
+ src1_data = thread->get_operand_value(src1, dst, to_type, thread, 1); // a
+ src2_data = thread->get_operand_value(src2, dst, to_type, thread, 1); // b
+ } else {
+ //This is the case whent he first argument (dest) is '_'
+ src1_data = thread->get_operand_value(src1, src1, to_type, thread, 1); // a
+ src2_data = thread->get_operand_value(src2, src1, to_type, thread, 1); // b
+ }
// Copy value pointed to in operand 'a' into register 'd'
// (i.e. copy src1_data to dst)
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'
+ if (dst.get_symbol()->type()){
+ thread->set_operand_value(dst, data, to_type, thread, pI); // Write value into register 'd'
+ }
// Get the atomic operation to be performed
unsigned m_atomic_spec = pI->get_atomic();
@@ -1080,7 +1089,12 @@ void atom_impl( const ptx_instruction *pI, ptx_thread_info *thread )
const operand_info &src1 = pI->src1();
const operand_info &dst = pI->dst();
unsigned i_type = pI->get_type();
- ptx_reg_t src1_data = thread->get_operand_value(src1, dst, i_type, thread, 1);
+ ptx_reg_t src1_data;
+ if (dst.get_symbol()->type()){
+ src1_data = thread->get_operand_value(src1, dst, i_type, thread, 1);
+ } else {
+ src1_data = thread->get_operand_value(src1, src1, i_type, thread, 1);
+ }
memory_space_t space = pI->get_space();
@@ -1944,10 +1958,10 @@ void ex2_impl( const ptx_instruction *pI, ptx_thread_info *thread )
}
void exit_impl( const ptx_instruction *pI, ptx_thread_info *thread )
-{
+{
+ thread->set_done();
thread->exitCore();
thread->registerExit();
- thread->set_done();
}
void mad_def( const ptx_instruction *pI, ptx_thread_info *thread );
@@ -2755,9 +2769,9 @@ void ret_impl( const ptx_instruction *pI, ptx_thread_info *thread )
{
bool empty = thread->callstack_pop();
if( empty ) {
+ thread->set_done();
thread->exitCore();
thread->registerExit();
- thread->set_done();
}
}
@@ -2766,9 +2780,9 @@ void retp_impl( const ptx_instruction *pI, ptx_thread_info *thread )
{
bool empty = thread->callstack_pop_plus();
if( empty ) {
+ thread->set_done();
thread->exitCore();
thread->registerExit();
- thread->set_done();
}
}
diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l
index 70d4c23..8c330fc 100644
--- a/src/cuda-sim/ptx.l
+++ b/src/cuda-sim/ptx.l
@@ -381,8 +381,9 @@ breakaddr TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE;
[^\"]* TC; ptx_lval.string_value = strdup(yytext);
}
-<INITIAL,NOT_OPCODE>. TC; ptx_error((const char*)NULL);
+<*>\t@@DWARF.*\n
+<INITIAL,NOT_OPCODE>. TC; ptx_error((const char*)NULL);
%%
extern int g_error_detected;
diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y
index f0b67fb..430fc00 100644
--- a/src/cuda-sim/ptx.y
+++ b/src/cuda-sim/ptx.y
@@ -476,7 +476,9 @@ operand: IDENTIFIER { add_scalar_operand( $1 ); }
| tex_operand
| IDENTIFIER PLUS INT_OPERAND { add_address_operand($1,$3); }
| IDENTIFIER LO_OPTION { add_scalar_operand( $1 ); change_operand_lohi(1);}
+ | MINUS IDENTIFIER LO_OPTION { add_scalar_operand( $2 ); change_operand_lohi(1); change_operand_neg();}
| IDENTIFIER HI_OPTION { add_scalar_operand( $1 ); change_operand_lohi(2);}
+ | MINUS IDENTIFIER HI_OPTION { add_scalar_operand( $2 ); change_operand_lohi(2); change_operand_neg();}
| IDENTIFIER PIPE IDENTIFIER { add_2vector_operand($1,$3); change_double_operand_type(-1);}
| IDENTIFIER PIPE IDENTIFIER LO_OPTION { add_2vector_operand($1,$3); change_double_operand_type(-1); change_operand_lohi(1);}
| IDENTIFIER PIPE IDENTIFIER HI_OPTION { add_2vector_operand($1,$3); change_double_operand_type(-1); change_operand_lohi(2);}
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index 61a357e..d3d87f7 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -29,6 +29,7 @@
#include "ptx_ir.h"
#include "cuda-sim.h"
#include "ptx_parser.h"
+#include <unistd.h>
#include <dirent.h>
#include <fstream>
@@ -45,7 +46,7 @@ bool g_override_embedded_ptx = false;
extern "C" int ptx_parse();
extern "C" int ptx__scan_string(const char*);
-const char *g_ptxinfo_filename = NULL;
+extern "C" const char *g_ptxinfo_filename = NULL;
extern "C" int ptxinfo_parse();
extern "C" int ptxinfo_debug;
extern "C" FILE *ptxinfo_in;
@@ -89,8 +90,9 @@ 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, bool save_converted )
+char* gpgpu_ptx_sim_convert_ptx_to_ptxplus(const char *ptx_str, const char *cubin_str, unsigned source_num, bool save_converted, bool cuobjdump)
{
+
printf("GPGPU-Sim PTX: converting EMBEDDED .ptx file to ptxplus \n");
// Extract ptx to a file
@@ -104,49 +106,60 @@ char* gpgpu_ptx_sim_convert_ptx_to_ptxplus(const char *ptx_str, const char *cubi
fprintf(ptxfile,"%s",ptx_str);
fclose(ptxfile);
- // Extract cubin to a file
- char fname_cubin[1024];
- snprintf(fname_cubin,1024,"_cubin_XXXXXX");
- int fd2=mkstemp(fname_cubin);
- close(fd2);
-
- printf("GPGPU-Sim PTX: extracting embedded cubin to temporary file \"%s\"\n", fname_cubin);
- FILE *cubinfile = fopen(fname_cubin,"w");
- fprintf(cubinfile,"%s",cubin_str);
- fclose(cubinfile);
+ char fname_ptxplus[1024];
+ snprintf(fname_ptxplus,1024,"_ptxplus_XXXXXX");
+ int fd4=mkstemp(fname_ptxplus);
+ close(fd4);
- // Run decuda
+ // Extract cubin to a file
char fname_decuda[1024];
- snprintf(fname_decuda,1024,"_decuda_XXXXXX");
- int fd3=mkstemp(fname_decuda);
- close(fd3);
+ char fname_cubin[1024];
+ if( !cuobjdump){
- char decuda_commandline[1024];
- snprintf(decuda_commandline,1024,"$DECUDA_INSTALL_PATH/decuda.py -o %s %s", fname_decuda, fname_cubin);
+ snprintf(fname_cubin,1024,"_cubin_XXXXXX");
+ int fd2=mkstemp(fname_cubin);
+ close(fd2);
- printf("GPGPU-Sim PTX: calling decuda on cubin file, decuda output file = \"%s\"\n", fname_decuda);
- int decuda_result = system(decuda_commandline);
- if( decuda_result != 0 ) {
- printf("GPGPU-Sim PTX: ERROR ** while calling decuda (b) %d\n", decuda_result);
- printf(" Ensure env variable DECUDA_INSTALL_PATH is set and points to decuda base directory.\n");
- exit(1);
- }
+ printf("GPGPU-Sim PTX: extracting embedded cubin to temporary file \"%s\"\n", fname_cubin);
+ FILE *cubinfile = fopen(fname_cubin,"w");
+ fprintf(cubinfile,"%s",cubin_str);
+ fclose(cubinfile);
- // Run decuda_to_ptxplus
- char fname_ptxplus[1024];
- snprintf(fname_ptxplus,1024,"_ptxplus_XXXXXX");
- int fd4=mkstemp(fname_ptxplus);
- close(fd4);
- char d2pp_commandline[1024];
- snprintf(d2pp_commandline,1024,"$D2PP_INSTALL_PATH/decuda_to_ptxplus %s %s %s %s > /dev/null", fname_decuda, fname_ptx, fname_cubin, fname_ptxplus);
+ snprintf(fname_decuda,1024,"_decuda_XXXXXX");
+ int fd3=mkstemp(fname_decuda);
+ close(fd3);
- printf("GPGPU-Sim PTX: calling decuda_to_ptxplus, ptxplus output file = \"%s\"\n", fname_ptxplus);
- int d2pp_result = system(d2pp_commandline);
- if( d2pp_result != 0 ) {
- printf("GPGPU-Sim PTX: ERROR ** while calling decuda_to_ptxplus %d\n", d2pp_result);
- printf(" Ensure env variable D2PP_INSTALL_PATH is set and points to decuda_to_ptxplus base directory.\n");
- exit(1);
+ // Run decuda
+ char decuda_commandline[1024];
+ snprintf(decuda_commandline,1024,"$DECUDA_INSTALL_PATH/decuda.py -o %s %s", fname_decuda, fname_cubin);
+ printf("GPGPU-Sim PTX: calling decuda on cubin file, decuda output file = \"%s\"\n", fname_decuda);
+ int decuda_result = system(decuda_commandline);
+ if( decuda_result != 0 ) {
+ printf("GPGPU-Sim PTX: ERROR ** while calling decuda (b) %d\n", decuda_result);
+ printf(" Ensure env variable DECUDA_INSTALL_PATH is set and points to decuda base directory.\n");
+ exit(1);
+ }
+ // Run decuda_to_ptxplus
+ char d2pp_commandline[1024];
+ snprintf(d2pp_commandline,1024,"$D2PP_INSTALL_PATH/decuda_to_ptxplus %s %s %s %s > /dev/null", fname_decuda, fname_ptx, fname_cubin, fname_ptxplus);
+ printf("GPGPU-Sim PTX: calling decuda_to_ptxplus, ptxplus output file = \"%s\"\n", fname_ptxplus);
+ int d2pp_result = system(d2pp_commandline);
+ if( d2pp_result != 0 ) {
+ printf("GPGPU-Sim PTX: ERROR ** while calling decuda_to_ptxplus %d\n", d2pp_result);
+ printf(" Ensure env variable D2PP_INSTALL_PATH is set and points to decuda_to_ptxplus base directory.\n");
+ exit(1);
+ }
+ }
+ else {
+ // Run cuobjdump_to_ptxplus
+ char commandline[1024];
+ int result;
+ snprintf(commandline, 1024, "$GPGPUSIM_ROOT/cuobjdump_to_ptxplus/cuobjdump_to_ptxplus /proc/%d/exe %s", getpid(), fname_ptxplus);
+ fflush(stdout);
+ printf("GPGPU-Sim PTX: calling cuobjdump_to_ptxplus on exe file, output file = \"%s\"\n", fname_ptxplus);
+ result = system(commandline);
+ if(result){printf("GPGPU-Sim PTX: ERROR ** could not execute %s\n", commandline); exit(1);}
}
// Get ptxplus from file
@@ -173,7 +186,12 @@ char* gpgpu_ptx_sim_convert_ptx_to_ptxplus(const char *ptx_str, const char *cubi
// Remove temporary files
char rm_commandline[1024];
- snprintf(rm_commandline,1024,"rm -f %s %s %s %s", fname_ptx, fname_cubin, fname_decuda, fname_ptxplus);
+ if (!cuobjdump) {
+ snprintf(rm_commandline,1024,"rm -f %s %s %s %s", fname_ptx, fname_cubin, fname_decuda, fname_ptxplus);
+ }
+ else {
+ snprintf(rm_commandline,1024,"rm -f %s %s", fname_ptx, fname_ptxplus);
+ }
printf("GPGPU-Sim PTX: removing temporary files using \"%s\"\n", rm_commandline);
int rm_result = system(rm_commandline);
if( rm_result != 0 ) {
@@ -186,6 +204,92 @@ char* gpgpu_ptx_sim_convert_ptx_to_ptxplus(const char *ptx_str, const char *cubi
return ptxplus_str;
}
+char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const char *ptxfilename, const char *elffilename, const char *sassfilename)
+{
+
+ printf("GPGPU-Sim PTX: converting EMBEDDED .ptx file to ptxplus \n");
+
+/*
+ // Extract ptx to a file
+ char fname_ptx[1024];
+ snprintf(fname_ptx,1024,"_ptx_XXXXXX");
+ int fd1=mkstemp(fname_ptx);
+ close(fd1);
+
+ printf("GPGPU-Sim PTX: extracting embedded .ptx to temporary file \"%s\"\n", fname_ptx);
+ FILE *ptxfile = fopen(fname_ptx,"w");
+ fprintf(ptxfile,"%s",ptx_str);
+ fclose(ptxfile);
+
+ char fname_sass[1024];
+ snprintf(fname_sass,1024,"_sass_XXXXXX");
+ int fd2=mkstemp(fname_sass);
+ close(fd2);
+
+ printf("GPGPU-Sim PTX: extracting .sass to temporary file \"%s\"\n", fname_sass);
+ FILE *sassfile = fopen(fname_sass,"w");
+ fprintf(sassfile,"%s",sass_str);
+ fclose(sassfile);
+
+ char fname_elf[1024];
+ snprintf(fname_elf,1024,"_elf_XXXXXX");
+ int fd3=mkstemp(fname_elf);
+ close(fd3);
+
+ printf("GPGPU-Sim PTX: extracting .elf to temporary file \"%s\"\n", fname_elf);
+ FILE *elffile = fopen(fname_elf,"w");
+ fprintf(elffile,"%s",elf_str);
+ fclose(elffile);
+*/
+ char fname_ptxplus[1024];
+ snprintf(fname_ptxplus,1024,"_ptxplus_XXXXXX");
+ int fd4=mkstemp(fname_ptxplus);
+ close(fd4);
+
+ // Extract cubin to a file
+ char fname_decuda[1024];
+ char fname_cubin[1024];
+
+
+ // Run cuobjdump_to_ptxplus
+ char commandline[1024];
+ int result;
+ snprintf(commandline, 1024, "$GPGPUSIM_ROOT/cuobjdump_to_ptxplus/cuobjdump_to_ptxplus %s %s %s %s", ptxfilename, sassfilename, elffilename, fname_ptxplus);
+ fflush(stdout);
+ printf("GPGPU-Sim PTX: calling cuobjdump_to_ptxplus\ncommandline: %s\n", commandline);
+ result = system(commandline);
+ if(result){printf("GPGPU-Sim PTX: ERROR ** could not execute %s\n", commandline); exit(1);}
+
+
+ // Get ptxplus from file
+ std::ifstream fileStream(fname_ptxplus, std::ios::in);
+ std::string text, line;
+ while(getline(fileStream,line)) {
+ text += (line + "\n");
+ }
+ fileStream.close();
+
+ char* ptxplus_str = new char [strlen(text.c_str())+1];
+ strcpy(ptxplus_str, text.c_str());
+
+ // Remove temporary files
+ /*
+ char rm_commandline[1024];
+
+ snprintf(rm_commandline,1024,"rm -f %s", fname_ptxplus);
+
+ printf("GPGPU-Sim PTX: removing temporary files using \"%s\"\n", rm_commandline);
+ int rm_result = system(rm_commandline);
+ if( rm_result != 0 ) {
+ printf("GPGPU-Sim PTX: ERROR ** while removing temporary files %d\n", rm_result);
+ exit(1);
+ }
+ */
+ 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, unsigned source_num )
{
@@ -237,6 +341,7 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
close(fd);
char commandline2[4096];
snprintf(commandline2,4096,"cat %s | sed 's/.version 1.5/.version 1.4/' | sed 's/, texmode_independent//' | sed 's/\\(\\.extern \\.const\\[1\\] .b8 \\w\\+\\)\\[\\]/\\1\\[1\\]/' | sed 's/const\\[.\\]/const\\[0\\]/g' > %s", fname, fname2);
+ printf("Running: %s\n", commandline2);
int result = system(commandline2);
if( result != 0 ) {
printf("GPGPU-Sim PTX: ERROR ** while loading PTX (a) %d\n", result);
@@ -252,10 +357,10 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
extra_flags[0]=0;
#if CUDART_VERSION >= 3000
- snprintf(extra_flags,1024,"--gpu-name=sm_20");
+ snprintf(extra_flags,1024,"--gpu-name=sm_13");
#endif
- snprintf(commandline,1024,"ptxas %s -v %s --output-file /dev/null 2> %s",
+ 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);
result = system(commandline);
diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h
index 0266755..aa98a91 100644
--- a/src/cuda-sim/ptx_loader.h
+++ b/src/cuda-sim/ptx_loader.h
@@ -32,6 +32,7 @@ extern bool g_override_embedded_ptx;
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 );
+char* gpgpu_ptx_sim_convert_ptx_to_ptxplus(const char *ptx_str, const char *cubin_str, unsigned source_num, bool save_converted, bool cuobjdump );
+char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const char *ptx_str, const char *sass_str, const char *elf_str);
#endif
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index 3db984d..308405a 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -43,6 +43,7 @@ const char *g_filename;
unsigned g_max_regs_per_thread = 0;
// the program intermediate representation...
+static symbol_table *g_global_allfiles_symbol_table = NULL;
static symbol_table *g_global_symbol_table = NULL;
std::map<std::string,symbol_table*> g_sym_name_to_symbol_table;
static symbol_table *g_current_symbol_table = NULL;
@@ -105,7 +106,13 @@ void read_parser_environment_variables()
symbol_table *init_parser( const char *ptx_filename )
{
g_filename = strdup(ptx_filename);
- g_global_symbol_table = g_current_symbol_table = new symbol_table("global",0,NULL);
+ if (g_global_allfiles_symbol_table == NULL) {
+ g_global_allfiles_symbol_table = new symbol_table("global_allfiles", 0, NULL);
+ g_global_symbol_table = g_current_symbol_table = g_global_allfiles_symbol_table;
+ }
+ else {
+ g_global_symbol_table = g_current_symbol_table = new symbol_table("global",0,g_global_allfiles_symbol_table);
+ }
ptx_lineno = 1;
#define DEF(X,Y) g_ptx_token_decode[X] = Y;
@@ -767,7 +774,7 @@ void add_scalar_operand( const char *identifier )
DPRINTF("add_scalar_operand");
const symbol *s = g_current_symbol_table->lookup(identifier);
if ( s == NULL ) {
- if ( g_opcode == BRA_OP || g_opcode == CALLP_OP || g_opcode == BREAKADDR_OP) {
+ if ( g_opcode == BRA_OP || g_opcode == CALLP_OP) {
// forward branch target...
s = g_current_symbol_table->add_variable(identifier,NULL,0,g_filename,ptx_lineno);
} else {
diff --git a/src/gpgpu-sim/dram.cc b/src/gpgpu-sim/dram.cc
index 104ec66..c073f0d 100644
--- a/src/gpgpu-sim/dram.cc
+++ b/src/gpgpu-sim/dram.cc
@@ -56,13 +56,24 @@ dram_t::dram_t( unsigned int partition_id, const struct memory_config *config, m
rw = READ; //read mode is default
+ bkgrp = (bankgrp_t**) calloc(sizeof(bankgrp_t*), m_config->nbkgrp);
+ bkgrp[0] = (bankgrp_t*) calloc(sizeof(bank_t), m_config->nbkgrp);
+ for (unsigned i=1; i<m_config->nbkgrp; i++) {
+ bkgrp[i] = bkgrp[0] + i;
+ }
+ for (unsigned i=0; i<m_config->nbkgrp; i++) {
+ bkgrp[i]->CCDLc = 0;
+ bkgrp[i]->RTPLc = 0;
+ }
+
bk = (bank_t**) calloc(sizeof(bank_t*),m_config->nbk);
bk[0] = (bank_t*) calloc(sizeof(bank_t),m_config->nbk);
for (unsigned i=1;i<m_config->nbk;i++)
bk[i] = bk[0] + i;
- for (unsigned i=0;i<m_config->nbk;i++)
+ for (unsigned i=0;i<m_config->nbk;i++) {
bk[i]->state = BANK_IDLE;
-
+ bk[i]->bkgrpindex = i/(m_config->nbk/m_config->nbkgrp);
+ }
prio = 0;
rwq = new fifo_pipeline<dram_req_t>("rwq",m_config->CL,m_config->CL+1);
mrqq = new fifo_pipeline<dram_req_t>("mrqq",0,2);
@@ -248,10 +259,12 @@ void dram_t::cycle()
// check if any bank is ready to issue a new read
for (unsigned i=0;i<m_config->nbk;i++) {
unsigned j = (i + prio) % m_config->nbk;
+ unsigned grp = j>>m_config->bk_tag_length;
if (bk[j]->mrq) { //if currently servicing a memory request
bk[j]->mrq->data->set_status(IN_PARTITION_DRAM,gpu_sim_cycle+gpu_tot_sim_cycle);
// correct row activated for a READ
if ( !issued && !CCDc && !bk[j]->RCDc &&
+ !(bkgrp[grp]->CCDLc) &&
(bk[j]->curr_row == bk[j]->mrq->row) &&
(bk[j]->mrq->rw == READ) && (WTRc == 0 ) &&
(bk[j]->state == BANK_ACTIVE) &&
@@ -263,8 +276,10 @@ void dram_t::cycle()
rwq->push(bk[j]->mrq);
bk[j]->mrq->txbytes += m_config->BL * m_config->busW * m_config->gpu_n_mem_per_ctrlr; //16 bytes
CCDc = m_config->tCCD;
+ bkgrp[grp]->CCDLc = m_config->tCCDL;
RTWc = m_config->tRTW;
bk[j]->RTPc = m_config->BL/2;
+ bkgrp[grp]->RTPLc = m_config->tRTPL;
issued = true;
n_rd++;
bwutil+= m_config->BL/2;
@@ -283,6 +298,7 @@ void dram_t::cycle()
} else
// correct row activated for a WRITE
if ( !issued && !CCDc && !bk[j]->RCDWRc &&
+ !(bkgrp[grp]->CCDLc) &&
(bk[j]->curr_row == bk[j]->mrq->row) &&
(bk[j]->mrq->rw == WRITE) && (RTWc == 0 ) &&
(bk[j]->state == BANK_ACTIVE) &&
@@ -295,6 +311,7 @@ void dram_t::cycle()
bk[j]->mrq->txbytes += m_config->BL * m_config->busW * m_config->gpu_n_mem_per_ctrlr; /*16 bytes*/
CCDc = m_config->tCCD;
+ bkgrp[grp]->CCDLc = m_config->tCCDL;
WTRc = m_config->tWTR;
bk[j]->WTPc = m_config->tWTP;
issued = true;
@@ -342,7 +359,9 @@ void dram_t::cycle()
if ( (!issued) &&
(bk[j]->curr_row != bk[j]->mrq->row) &&
(bk[j]->state == BANK_ACTIVE) &&
- (!bk[j]->RASc && !bk[j]->WTPc && !bk[j]->RTPc) ) {
+ (!bk[j]->RASc && !bk[j]->WTPc &&
+ !bk[j]->RTPc &&
+ !bkgrp[grp]->RTPLc) ) {
// make the bank idle again
bk[j]->state = BANK_IDLE;
bk[j]->RPc = m_config->tRP;
@@ -389,6 +408,10 @@ void dram_t::cycle()
DEC2ZERO(bk[j]->WTPc);
DEC2ZERO(bk[j]->RTPc);
}
+ for (unsigned j=0; j<m_config->nbkgrp; j++) {
+ DEC2ZERO(bkgrp[j]->CCDLc);
+ DEC2ZERO(bkgrp[j]->RTPLc);
+ }
#ifdef DRAM_VISUALIZE
visualize();
diff --git a/src/gpgpu-sim/dram.h b/src/gpgpu-sim/dram.h
index 96b912b..ad773b3 100644
--- a/src/gpgpu-sim/dram.h
+++ b/src/gpgpu-sim/dram.h
@@ -58,6 +58,12 @@ public:
class mem_fetch * data;
};
+struct bankgrp_t
+{
+ unsigned int CCDLc;
+ unsigned int RTPLc;
+};
+
struct bank_t
{
unsigned int RCDc;
@@ -77,6 +83,8 @@ struct bank_t
unsigned int n_access;
unsigned int n_writes;
unsigned int n_idle;
+
+ unsigned int bkgrpindex;
};
struct mem_fetch;
@@ -110,6 +118,8 @@ private:
const struct memory_config *m_config;
+ bankgrp_t **bkgrp;
+
bank_t **bk;
unsigned int prio;
diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc
index 51b832c..e71818b 100644
--- a/src/gpgpu-sim/gpu-sim.cc
+++ b/src/gpgpu-sim/gpu-sim.cc
@@ -123,8 +123,8 @@ void memory_config::reg_options(class OptionParser * opp)
"Burst length of each DRAM request (default = 4 DDR cycle)",
"4");
option_parser_register(opp, "-gpgpu_dram_timing_opt", OPT_CSTR, &gpgpu_dram_timing_opt,
- "DRAM timing parameters = {nbk:tCCD:tRRD:tRCD:tRAS:tRP:tRC:CL:WL:tCDLR:tWR}",
- "4:2:8:12:21:13:34:9:4:5:13");
+ "DRAM timing parameters = {nbk:nbkgrp:tCCD:tCCDL:tRRD:tRCD:tRAS:tRP:tRTPL:tRC:CL:WL:tCDLR:tWR}",
+ "4:1:2:0:8:12:21:13:0:34:9:4:5:13");
option_parser_register(opp, "-rop_latency", OPT_UINT32, &rop_latency,
"ROP queue latency (default 85)",
"85");
diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h
index 31def4b..fe66d04 100644
--- a/src/gpgpu-sim/gpu-sim.h
+++ b/src/gpgpu-sim/gpu-sim.h
@@ -72,8 +72,14 @@ struct memory_config {
void init()
{
assert(gpgpu_dram_timing_opt);
- sscanf(gpgpu_dram_timing_opt,"%d:%d:%d:%d:%d:%d:%d:%d:%d:%d:%d",
- &nbk,&tCCD,&tRRD,&tRCD,&tRAS,&tRP,&tRC,&CL,&WL,&tCDLR,&tWR);
+ sscanf(gpgpu_dram_timing_opt,"%d:%d:%d:%d:%d:%d:%d:%d:%d:%d:%d:%d:%d:%d",
+ &nbk,&nbkgrp,&tCCD,&tCCDL,&tRRD,&tRCD,&tRAS,&tRP,&tRTPL,&tRC,&CL,&WL,&tCDLR,&tWR);
+ int nbkt = nbk/nbkgrp;
+ unsigned i;
+ for (i=0; nbkt>0; i++) {
+ nbkt = nbkt>>1;
+ }
+ bk_tag_length = i;
tRCDWR = tRCD-(WL+1);
tRTW = (CL+(BL/2)+2-WL);
tWTR = (WL+(BL/2)+tCDLR);
@@ -101,6 +107,10 @@ struct memory_config {
unsigned dram_latency;
// DRAM parameters
+
+ unsigned tCCDL; //column to column delay when bank groups are enabled
+ unsigned tRTPL; //read to precharge delay when bank groups are enabled for GDDR5 this is identical to RTPS, if for other DRAM this is different, you will need to split them in two
+
unsigned tCCD; //column to column delay
unsigned tRRD; //minimal time required between activation of rows in different banks
unsigned tRCD; //row to column delay - time required to activate a row before a read
@@ -119,6 +129,9 @@ struct memory_config {
unsigned tWTP; //time to switch from write to precharge in the same bank
unsigned busW;
+ unsigned nbkgrp; // number of bank groups (has to be power of 2)
+ unsigned bk_tag_length; //number of bits that define a bank inside a bank group
+
unsigned nbk;
linear_to_raw_address_translation m_address_mapping;
@@ -237,7 +250,7 @@ public:
const gpgpu_sim_config &get_config() const { return m_config; }
void gpu_print_stat() const;
void dump_pipeline( int mask, int s, int m ) const;
-
+
//The next three functions added to be used by the functional simulation function
//! Get shader core configuration
@@ -270,6 +283,7 @@ private:
void shader_print_runtime_stat( FILE *fout );
void shader_print_l1_miss_stat( FILE *fout ) const;
void visualizer_printstat();
+ void print_shader_cycle_distro( FILE *fout ) const;
void gpgpu_debug();
@@ -311,7 +325,7 @@ public:
unsigned long long gpu_sim_insn;
unsigned long long gpu_tot_sim_insn;
unsigned long long gpu_sim_insn_last_update;
- unsigned gpu_sim_insn_last_update_sid;
+ unsigned gpu_sim_insn_last_update_sid;
};
#endif
diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc
index 4c90667..527fdcb 100644
--- a/src/gpgpu-sim/shader.cc
+++ b/src/gpgpu-sim/shader.cc
@@ -1736,7 +1736,22 @@ void barrier_set_t::dump() const
void shader_core_ctx::warp_exit( unsigned warp_id )
{
- m_barriers.warp_exit( warp_id );
+ bool done = true;
+ for ( unsigned i = warp_id*get_config()->warp_size;
+ i < (warp_id+1)*get_config()->warp_size;
+ i++ ) {
+
+// if(this->m_thread[i]->m_functional_model_thread_state && this->m_thread[i].m_functional_model_thread_state->donecycle()==0) {
+// done = false;
+// }
+
+
+ if (m_thread[i] && !m_thread[i]->is_done()) done = false;
+ }
+ //if (m_warp[warp_id].get_n_completed() == get_config()->warp_size)
+ //if (this->m_simt_stack[warp_id]->get_num_entries() == 0)
+ if (done)
+ m_barriers.warp_exit( warp_id );
}
bool shader_core_ctx::warp_waiting_at_barrier( unsigned warp_id ) const
diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc
index e18b088..add9a22 100644
--- a/src/gpgpusim_entrypoint.cc
+++ b/src/gpgpusim_entrypoint.cc
@@ -238,15 +238,8 @@ int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid )
return 0;
}
-//! Functional simulation of OpenCL
-/*!
- * This function call the CUDA PTX functional simulator
- */
int gpgpu_opencl_ptx_sim_main_func( kernel_info_t *grid )
{
- //calling the CUDA PTX simulator, sending the kernel by reference and a flag set to true,
- //the flag used by the function to distinguish OpenCL calls from the CUDA simulation calls which
- //it is needed by the called function to not register the exit the exit of OpenCL kernel as it doesn't register entering in the first place as the CUDA kernels does
- gpgpu_cuda_ptx_sim_main_func( *grid, true );
- return 0;
+ printf("GPGPU-Sim PTX API: OpenCL functional-only simulation not yet implemented (use performance simulation)\n");
+ exit(1);
}
generated by cgit v1.3 (git 2.53.0) at 2026-07-10 21:23:17 +0000