Merge pull request #26 from sspenst/dev

Additional support for CUDA 7.5 

15 files changed, 390 insertions, 28 deletions

diff --git a/configs/GeForceGTX750Ti/gpgpusim.config b/configs/GeForceGTX750Ti/gpgpusim.config
index afd3825..8b030b6 100644
--- a/configs/GeForceGTX750Ti/gpgpusim.config
+++ b/configs/GeForceGTX750Ti/gpgpusim.config
@@ -1,7 +1,7 @@
 # functional simulator specification
 -gpgpu_ptx_instruction_classification 0
 -gpgpu_ptx_sim_mode 0
--gpgpu_ptx_force_max_capability 20 
+-gpgpu_ptx_force_max_capability 52 
 
 
 # SASS execution (only supported with CUDA >= 4.0)
@@ -42,8 +42,8 @@
 -ptx_opcode_latency_dp 8,19,8,8,330
 -ptx_opcode_initiation_dp 8,16,8,8,130
 
--gpgpu_cache:dl1 none
-# 32:128:4,L:L:m:N:H,A:32:8,8
+-gpgpu_cache:dl1 32:128:4,L:L:m:N:H,A:32:8,8
+-gmem_skip_L1D 1
 -gpgpu_shmem_size 65536
 
 # The alternative configuration for fermi in case cudaFuncCachePreferL1 is selected
@@ -119,8 +119,10 @@
 -visualizer_enabled 0
 
 # power model configs
+# power modeling is currently disabled for GTX750Ti. The gpuwattch_gtx750ti.xml file present is only a placeholder and has NOT been tested for accuracy.
+# proper modeling of Maxwell power is planned, but should be considered completely unsupported at present.
 -power_simulation_enabled 0
--gpuwattch_xml_file gpuwattch_gtx480.xml
+-gpuwattch_xml_file gpuwattch_gtx750Ti.xml
 
 # tracing functionality
 #-trace_enabled 1
diff --git a/configs/GeForceGTX750Ti/gpuwattch_gtx750Ti.xml b/configs/GeForceGTX750Ti/gpuwattch_gtx750Ti.xml
index 304e0fd..e2b2324 100755
--- a/configs/GeForceGTX750Ti/gpuwattch_gtx750Ti.xml
+++ b/configs/GeForceGTX750Ti/gpuwattch_gtx750Ti.xml
@@ -1,3 +1,4 @@
+<!-- THIS CONFIG IS CURRENTLY PRESENT ONLY AS A PLACEHOLDER. IT HAS NOT BEEN TESTED FOR ACCURACY VERSUS GTX750TI HARDWARE. ANY USE OF THIS CONFIG WILL LIKELY GIVE HIGHLY INACCURATE RESULTS. -->
 <?xml version="1.0" ?>
 <component id="root" name="root">
 	<component id="system" name="system">
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 910bebd..e8a0e91 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -109,6 +109,7 @@
 #include <stdarg.h>
 #include <iostream>
 #include <string>
+#include <regex>
 #include <sstream>
 #include <fstream>
 #ifdef OPENGL_SUPPORT
@@ -133,6 +134,7 @@
 #include "../src/cuda-sim/ptx_parser.h"
 #include "../src/gpgpusim_entrypoint.h"
 #include "../src/stream_manager.h"
+#include "../src/abstract_hardware_model.h"
 
 #include <pthread.h>
 #include <semaphore.h>
@@ -319,9 +321,9 @@ class _cuda_device_id *GPGPUSim_Init()
 
 		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->major = 5;
+		prop->minor = 2;
+		prop->totalGlobalMem = 0x80000000 /* 2 GB */;
 		prop->memPitch = 0;
 		prop->maxThreadsPerBlock = 512;
 		prop->maxThreadsDim[0] = 512;
@@ -438,6 +440,10 @@ extern "C" {
  *                                                                              *
  *                                                                              *
  *******************************************************************************/
+cudaError_t cudaPeekAtLastError(void)
+{
+	return g_last_cudaError;
+}
 
 __host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size) 
 {
@@ -532,6 +538,22 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t cou
 		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 if ( kind == cudaMemcpyDefault ) {
+		if ((size_t)src >= GLOBAL_HEAP_START) {
+			if ((size_t)dst >= GLOBAL_HEAP_START)
+				g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,0) ); // device to device
+			else
+				g_stream_manager->push( stream_operation((size_t)src,dst,count,0) ); // device to host
+		}
+		else {
+			if ((size_t)dst >= GLOBAL_HEAP_START)
+				g_stream_manager->push( stream_operation(src,(size_t)dst,count,0) );
+			else {
+				printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported transfer: host to host\n");
+				abort();
+			}
+		}
+	}
 	else {
 		printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported cudaMemcpyKind\n");
 		abort();
@@ -966,6 +988,10 @@ __host__ cudaError_t CUDARTAPI cudaStreamCreate(cudaStream_t *stream)
 	return g_last_cudaError = cudaSuccess;
 }
 
+__host__ __device__ cudaError_t CUDARTAPI cudaStreamCreateWithFlags(cudaStream_t *stream, unsigned int flags) {
+	return cudaStreamCreate(stream);
+}
+
 __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream)
 {
 #if (CUDART_VERSION >= 3000)
@@ -978,7 +1004,8 @@ __host__ cudaError_t CUDARTAPI cudaStreamSynchronize(cudaStream_t stream)
 {
 #if (CUDART_VERSION >= 3000)
 	if( stream == NULL )
-		return g_last_cudaError = cudaErrorInvalidResourceHandle;
+		synchronize();
+		return g_last_cudaError = cudaSuccess;
 	stream->synchronize();
 #else
 	printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__);
@@ -1495,6 +1522,78 @@ std::list<cuobjdumpSection*> pruneSectionList(std::list<cuobjdumpSection*> cuobj
 	return prunedList;
 }
 
+//! Merge all PTX sections that have a specific identifier into one file
+std::list<cuobjdumpSection*> mergeMatchingSections(std::list<cuobjdumpSection*> cuobjdumpSectionList, std::string identifier){
+	char *ptxcode = "";
+	std::list<cuobjdumpSection*>::iterator old_iter;
+	cuobjdumpPTXSection* old_ptxsection = NULL;
+	cuobjdumpPTXSection* ptxsection;
+	std::list<cuobjdumpSection*> mergedList;
+
+	for (	std::list<cuobjdumpSection*>::iterator iter = cuobjdumpSectionList.begin();
+			iter != cuobjdumpSectionList.end();
+			iter++){
+		if((ptxsection=dynamic_cast<cuobjdumpPTXSection*>(*iter)) != NULL &&
+			strcmp(ptxsection->getIdentifier().c_str(), identifier.c_str()) == 0){
+			// Read and remove the last PTX section
+			if (old_ptxsection != NULL) {
+				ptxcode = readfile(old_ptxsection->getPTXfilename());
+				// remove ptx file?
+				delete *old_iter;
+			}
+
+			// Append all the PTX from the last PTX section into the current PTX section
+			// Add 50 to ptxcode to ignore the information regarding version/target/address_size
+			if (strlen(ptxcode) >= 50) {
+				FILE *ptxfile = fopen((ptxsection->getPTXfilename()).c_str(), "a");
+				fprintf(ptxfile, "%s", ptxcode + 50);
+				fclose(ptxfile);
+			}
+
+			old_iter = iter;
+			old_ptxsection = ptxsection;
+		}
+		// Store all non-PTX sections and PTX sections with non-matching identifiers
+		else {
+			mergedList.push_back(*iter);
+		}
+	}
+
+	// Store the final PTX section
+	mergedList.push_back(*old_iter);
+
+	return mergedList;
+}
+
+//! Merge any PTX sections with matching identifiers
+std::list<cuobjdumpSection*> mergeSections(std::list<cuobjdumpSection*> cuobjdumpSectionList){
+	std::vector<std::string> identifier;
+	cuobjdumpPTXSection* ptxsection;
+
+	// Add all identifiers present in PTX sections to a vector
+	for (	std::list<cuobjdumpSection*>::iterator iter = cuobjdumpSectionList.begin();
+			iter != cuobjdumpSectionList.end();
+			iter++){
+		if((ptxsection=dynamic_cast<cuobjdumpPTXSection*>(*iter)) != NULL){
+			std::string current_id = ptxsection->getIdentifier();
+
+			// If we haven't yet seen a given identifier, add it to the vector
+			if (std::find(identifier.begin(), identifier.end(), current_id) == identifier.end()) {
+				identifier.push_back(current_id);
+			}
+		}
+	}
+
+	// Call mergeMatchingSections on all identifiers in the vector
+	for (	std::vector<std::string>::iterator iter = identifier.begin();
+			iter != identifier.end();
+			iter++) {
+		cuobjdumpSectionList = mergeMatchingSections(cuobjdumpSectionList, *iter);
+	}
+
+	return cuobjdumpSectionList;
+}
+
 
 //! Within the section list, find the ELF section corresponding to a given identifier
 cuobjdumpELFSection* findELFSectionInList(std::list<cuobjdumpSection*> sectionlist, const std::string identifier){
@@ -1558,10 +1657,12 @@ void cuobjdumpInit(){
 	CUctx_st *context = GPGPUSim_Context();
 	extract_code_using_cuobjdump(); //extract all the output of cuobjdump to _cuobjdump_*.*
 	cuobjdumpSectionList = pruneSectionList(cuobjdumpSectionList, context);
+	cuobjdumpSectionList = mergeSections(cuobjdumpSectionList);
 }
 
 std::map<int, std::string> fatbinmap;
 std::map<int, bool>fatbin_registered;
+std::map<std::string, symbol_table*> name_symtab;
 
 //! Keep track of the association between filename and cubin handle
 void cuobjdumpRegisterFatBinary(unsigned int handle, char* filename){
@@ -1574,6 +1675,13 @@ void cuobjdumpParseBinary(unsigned int handle){
 	if(fatbin_registered[handle]) return;
 	fatbin_registered[handle] = true;
 	CUctx_st *context = GPGPUSim_Context();
+	std::string fname = fatbinmap[handle];
+
+	if (name_symtab.find(fname) != name_symtab.end()) {
+		symbol_table *symtab = name_symtab[fname];
+		context->add_binary(symtab, handle);
+		return;
+	}
 
 	unsigned max_capability = 0;
 	for (	std::list<cuobjdumpSection*>::iterator iter = cuobjdumpSectionList.begin();
@@ -1582,16 +1690,13 @@ void cuobjdumpParseBinary(unsigned int handle){
 		unsigned capability = (*iter)->getArch();
 		if (capability > max_capability) max_capability = capability;
 	}
-
 	if (max_capability > 20) printf("WARNING: No guarantee that PTX will be parsed for SM version %u\n", max_capability);
 
-	std::string fname = fatbinmap[handle];
 	cuobjdumpPTXSection* ptx = findPTXSection(fname);
-
 	symbol_table *symtab;
 	char *ptxcode;
 	const char *override_ptx_name = getenv("PTX_SIM_KERNELFILE"); 
-   if (override_ptx_name == NULL or getenv("PTX_SIM_USE_PTX_FILE") == NULL) {
+	if (override_ptx_name == NULL or getenv("PTX_SIM_USE_PTX_FILE") == NULL) {
 		ptxcode = readfile(ptx->getPTXfilename());
 	} else {
 		printf("GPGPU-Sim PTX: overriding embedded ptx with '%s' (PTX_SIM_USE_PTX_FILE is set)\n", override_ptx_name);
@@ -1617,6 +1722,7 @@ void cuobjdumpParseBinary(unsigned int handle){
 	}
 	load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu());
 	load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu());
+	name_symtab[fname] = symtab;
 
 	//TODO: Remove temporarily files as per configurations
 }
@@ -1816,10 +1922,15 @@ void __cudaRegisterTexture(
 		int ext
 ) //passes in a newly created textureReference
 {
+	std::string devStr (deviceName);
+	#if (CUDART_VERSION > 4020)
+	if (devStr.size() > 2 && devStr.data()[0] == ':' && devStr.data()[1] == ':')
+		devStr = devStr.replace(0, 2, "");
+	#endif
 	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, dim, norm, ext);
+	gpu->gpgpu_ptx_sim_bindNameToTexture(devStr.data(), hostVar, dim, norm, ext);
 	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);
diff --git a/libcuda/cuobjdump.y b/libcuda/cuobjdump.y
index 9d61f25..31760f7 100644
--- a/libcuda/cuobjdump.y
+++ b/libcuda/cuobjdump.y
@@ -82,8 +82,7 @@ section :	PTXHEADER {
 				snprintf(filename, 1024, "_cuobjdump_%d.elf", elfserial);
 				elffile = fopen(filename, "w");
 				setCuobjdumpelffilename(filename);
-			} headerinfo identifier{
-			} elfcode {
+			} headerinfo compressedkeyword identifier elfcode {
 				fclose(elffile);
 				snprintf(filename, 1024, "_cuobjdump_%d.sass", elfserial++);
 				sassfile = fopen(filename, "w");
diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h
index ba4ea29..b29f918 100644
--- a/src/abstract_hardware_model.h
+++ b/src/abstract_hardware_model.h
@@ -334,7 +334,7 @@ protected:
     std::deque<simt_stack_entry> m_stack;
 };
 
-#define GLOBAL_HEAP_START 0x80000000
+#define GLOBAL_HEAP_START 0x703E20000
    // start allocating from this address (lower values used for allocating globals in .ptx file)
 #define SHARED_MEM_SIZE_MAX (64*1024)
 #define LOCAL_MEM_SIZE_MAX (8*1024)
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 715be98..fba3a59 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -719,7 +719,7 @@ void ptx_instruction::set_opcode_and_latency()
 		   break;
 	   }
 	   break;
-   case MAD_OP: case MADP_OP:
+   case MAD_OP: case MADC_OP:
 	   //MAD latency
 	   switch(get_type()){
 	   case F32_TYPE:
@@ -1864,10 +1864,14 @@ unsigned translate_pc_to_ptxlineno(unsigned pc)
 
 // ptxinfo parser
 
+extern std::map<unsigned,const char*> get_duplicate();
+
 int g_ptxinfo_error_detected;
 
 static char *g_ptxinfo_kname = NULL;
 static struct gpgpu_ptx_sim_info g_ptxinfo;
+static std::map<unsigned,const char*> g_duplicate;
+static const char *g_last_dup_type;
 
 const char *get_ptxinfo_kname() 
 { 
@@ -1897,6 +1901,21 @@ struct gpgpu_ptx_sim_info get_ptxinfo()
     return g_ptxinfo;
 }
 
+std::map<unsigned,const char*> get_duplicate()
+{
+	return g_duplicate;
+}
+
+void ptxinfo_linenum( unsigned linenum )
+{
+	g_duplicate[linenum] = g_last_dup_type;
+}
+
+void ptxinfo_dup_type( const char *dup_type )
+{
+	g_last_dup_type = dup_type;
+}
+
 void ptxinfo_function(const char *fname )
 {
     clear_ptxinfo();
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index cf7f04a..02ce01c 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -1336,7 +1336,82 @@ void bar_impl( const ptx_instruction *pIin, ptx_thread_info *thread )
    thread->m_last_dram_callback.instruction = pIin;
 }
 
-void bfe_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); }
+void bfe_impl( const ptx_instruction *pI, ptx_thread_info *thread ) 
+{
+	unsigned i_type = pI->get_type();
+	unsigned msb = (i_type == U32_TYPE || i_type == S32_TYPE) ? 31 : 63;
+    const operand_info &dst  = pI->dst();
+    const operand_info &src1 = pI->src1();
+    const operand_info &src2 = pI->src2();
+    const operand_info &src3 = pI->src3();
+    ptx_reg_t a = thread->get_operand_value(src1, dst, i_type, thread, 1);
+    ptx_reg_t b = thread->get_operand_value(src2, dst, i_type, thread, 1);
+    ptx_reg_t c = thread->get_operand_value(src3, dst, i_type, thread, 1);
+	unsigned pos = b.u32 & 0xFF;
+	unsigned len = c.u32 & 0xFF;
+	unsigned d = 0;
+	switch (i_type)
+	{
+		case U32_TYPE:
+		{
+			unsigned mask;
+			d = a.u32 >> pos;
+			mask = 0xFFFFFFFF >> (32 - len);
+			d &= mask;
+			break;
+		}
+		case U64_TYPE:
+		{
+			unsigned long mask;
+			d = a.u64 >> pos;	
+			mask = 0xFFFFFFFFFFFFFFFF >> (64 - len);
+			d &= mask;
+			break;
+		}
+		case S32_TYPE:
+		{
+			unsigned mask;
+			unsigned min = MY_MIN_I(pos + len - 1, msb);
+			unsigned sbit = len == 0 ? 0 : (a.s32 >> min) & 0x1;
+			d = a.s32 >> pos;
+			if (sbit > 0)
+			{
+				mask = 0xFFFFFFFF << len;
+				d |= mask;
+			}
+			else
+			{
+				mask = 0xFFFFFFFF >> (32 - len);
+				d &= mask;
+			}
+			break;
+		}
+		case S64_TYPE:
+		{
+			unsigned long mask;
+			unsigned min = MY_MIN_I(pos + len - 1, msb);
+			unsigned sbit = len == 0 ? 0 : (a.s64 >> min) & 0x1;
+			d = a.s64 >> pos;
+			if (sbit > 0)
+			{
+				mask = 0xFFFFFFFFFFFFFFFF << len;
+				d |= mask;
+			}
+			else
+			{
+				mask = 0xFFFFFFFFFFFFFFFF >> (64 - len);
+				d &= mask;
+			}
+			break;
+		}
+		default:
+		printf("Operand type not supported for BFE instruction.\n");
+		abort();
+		return;
+	}
+    thread->set_operand_value(dst,d, i_type, thread, pI);
+}
+
 void bfi_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); }
 void bfind_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); }
 
@@ -2421,7 +2496,7 @@ void mad_impl( const ptx_instruction *pI, ptx_thread_info *thread )
    mad_def(pI, thread, false);
 }
 
-void madp_impl( const ptx_instruction *pI, ptx_thread_info *thread ) 
+void madc_impl( const ptx_instruction *pI, ptx_thread_info *thread ) 
 {
    mad_def(pI, thread, true);
 }
diff --git a/src/cuda-sim/opcodes.def b/src/cuda-sim/opcodes.def
index 7aaa14f..874acc7 100644
--- a/src/cuda-sim/opcodes.def
+++ b/src/cuda-sim/opcodes.def
@@ -69,7 +69,7 @@ OP_DEF(LDU_OP,ldu_impl,"ldu",1,5)
 OP_DEF(LG2_OP,lg2_impl,"lg2",1,4)
 OP_DEF(MAD24_OP,mad24_impl,"mad24",1,2)
 OP_DEF(MAD_OP,mad_impl,"mad",1,2)
-OP_DEF(MADP_OP,madp_impl,"madp",1,2)
+OP_DEF(MADC_OP,madc_impl,"madc",1,2)
 OP_DEF(MAX_OP,max_impl,"max",1,1)
 OP_DEF(MEMBAR_OP,membar_impl,"membar",1,3)
 OP_DEF(MIN_OP,min_impl,"min",1,1)
diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l
index dfed936..b8ce497 100644
--- a/src/cuda-sim/ptx.l
+++ b/src/cuda-sim/ptx.l
@@ -86,7 +86,7 @@ ldu     TC; ptx_lval.int_value = LDU_OP; return OPCODE;
 lg2	TC; ptx_lval.int_value = LG2_OP; return OPCODE;
 mad24   TC; ptx_lval.int_value = MAD24_OP; return OPCODE;
 mad     TC; ptx_lval.int_value = MAD_OP; return OPCODE;
-madp    TC; ptx_lval.int_value = MADP_OP; return OPCODE;
+madc	TC; ptx_lval.int_value = MADC_OP; return OPCODE;
 max     TC; ptx_lval.int_value = MAX_OP; return OPCODE;
 membar  TC; ptx_lval.int_value = MEMBAR_OP; return OPCODE;
 min     TC; ptx_lval.int_value = MIN_OP; return OPCODE;
@@ -181,6 +181,7 @@ breakaddr  TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE;
 \.version TC; return VERSION_DIRECTIVE;
 \.visible TC; return VISIBLE_DIRECTIVE;
 \.address_size TC; return ADDRESS_SIZE_DIRECTIVE;
+\.weak	TC; return WEAK_DIRECTIVE;
 
 \.constptr TC; return CONSTPTR_DIRECTIVE; /* Ptx plus directive for pointer to constant memory */
 \.ptr TC; return PTR_DIRECTIVE; /* Added for new OpenCL genrated code */
@@ -249,6 +250,7 @@ breakaddr  TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE;
 \.v4	TC; return V4_TYPE;
 
 \.half  TC; return HALF_OPTION; /* ptxplus */
+\.cc	TC; return EXTP_OPTION; /* extended precision option */
 
 \.equ	TC; return EQU_OPTION;
 \.neu	TC; return NEU_OPTION;
@@ -324,6 +326,8 @@ breakaddr  TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE;
 \.wb	TC; return WB_OPTION;
 \.wt	TC; return WT_OPTION;
 
+\.nc	TC; return NC_OPTION;
+
 \.popc	TC; return ATOMIC_POPC;
 \.and	TC; return ATOMIC_AND;
 \.or	TC; return ATOMIC_OR;
@@ -386,7 +390,9 @@ breakaddr  TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE;
 }
 <IN_COMMENT>{
 "*/"      BEGIN(INITIAL);
-[^*\n]+   // eat comment in chunks
+"CPTX_BEGIN" printf("BEGINNING CUSTOM PTX.\n"); BEGIN(INITIAL);
+[^C*\n]+   // eat comment in chunks
+"C"
 "*"       // eat the lone star
 \n        TC; 
 }
diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y
index 79faddf..4de39d1 100644
--- a/src/cuda-sim/ptx.y
+++ b/src/cuda-sim/ptx.y
@@ -47,6 +47,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 %token  PTR_DIRECTIVE
 %token  ENTRY_DIRECTIVE
 %token  EXTERN_DIRECTIVE
+%token  WEAK_DIRECTIVE
 %token  FILE_DIRECTIVE
 %token  FUNC_DIRECTIVE
 %token  GLOBAL_DIRECTIVE
@@ -103,6 +104,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 %token  COMMA
 %token  PRED
 %token  HALF_OPTION
+%token  EXTP_OPTION
 %token  EQ_OPTION
 %token  NE_OPTION
 %token  LT_OPTION
@@ -191,6 +193,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 %token  CV_OPTION;
 %token  WB_OPTION;
 %token  WT_OPTION;
+%token	NC_OPTION;
 
 %type <int_value> function_decl_header
 %type <ptr_value> function_decl
@@ -242,6 +245,7 @@ function_decl_header: ENTRY_DIRECTIVE { $$ = 1; g_func_decl=1; func_header(".ent
 	| VISIBLE_DIRECTIVE ENTRY_DIRECTIVE { $$ = 1; g_func_decl=1; func_header(".entry"); }
 	| FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; func_header(".func"); }
 	| VISIBLE_DIRECTIVE FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; func_header(".func"); }
+	| WEAK_DIRECTIVE FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; func_header(".func"); }
 	| EXTERN_DIRECTIVE FUNC_DIRECTIVE { $$ = 2; g_func_decl=1; func_header(".func"); }
 	;
 
@@ -438,6 +442,7 @@ option: type_spec
 	| atomic_operation_spec ;
 	| TO_OPTION { add_option(TO_OPTION); }
 	| HALF_OPTION { add_option(HALF_OPTION); }
+	| EXTP_OPTION { add_option(EXTP_OPTION); }
 	| CA_OPTION { add_option(CA_OPTION); }
 	| CG_OPTION { add_option(CG_OPTION); }
 	| CS_OPTION { add_option(CS_OPTION); }
@@ -445,6 +450,7 @@ option: type_spec
 	| CV_OPTION { add_option(CV_OPTION); }
 	| WB_OPTION { add_option(WB_OPTION); }
 	| WT_OPTION { add_option(WT_OPTION); }
+	| NC_OPTION { add_option(NC_OPTION); }
 	;
 
 atomic_operation_spec: ATOMIC_AND { add_option(ATOMIC_AND); } 
diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc
index 751b3f4..2eccabc 100644
--- a/src/cuda-sim/ptx_ir.cc
+++ b/src/cuda-sim/ptx_ir.cc
@@ -1167,6 +1167,10 @@ ptx_instruction::ptx_instruction( int opcode,
       case HALF_OPTION:
          m_inst_size = 4; // bytes
          break;
+	  case EXTP_OPTION:
+		 break;
+	  case NC_OPTION:
+		 break;
       default:
          assert(0);
          break;
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index f7bf70e..9bb5008 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -46,6 +46,8 @@ bool g_override_embedded_ptx = false;
 extern int ptx_parse();
 extern int ptx__scan_string(const char*);
 
+extern std::map<unsigned,const char*> get_duplicate();
+
 const char *g_ptxinfo_filename;
 extern int ptxinfo_parse();
 extern int ptxinfo_debug;
@@ -182,6 +184,109 @@ symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source
     return symtab;
 }
 
+void fix_duplicate_errors(char fname2[1024]) {
+	char tempfile[1024] = "_temp_ptx";
+	char commandline[1024];
+
+	// change the name of the ptx file to _temp_ptx
+	snprintf(commandline,1024,"mv %s %s",fname2,tempfile);
+	printf("Running: %s\n", commandline);
+	int result = system(commandline);
+	if (result != 0) {
+		printf("GPGPU-Sim PTX: ERROR ** while changing filename from %s to %s", fname2, tempfile);
+		exit(1);
+	}
+
+	// store all of the ptx into a char array
+	FILE *ptxsource = fopen(tempfile,"r");
+	fseek(ptxsource, 0, SEEK_END);
+	long filesize = ftell(ptxsource);
+	rewind(ptxsource);
+	char *ptxdata = (char*)malloc((filesize+1)*sizeof(char));
+	fread(ptxdata, filesize, 1, ptxsource);
+	fclose(ptxsource);
+
+	FILE *ptxdest = fopen(fname2,"w");
+	std::map<unsigned,const char*> duplicate = get_duplicate();
+	unsigned offset;
+	unsigned oldlinenum = 1;
+	unsigned linenum;
+	char *startptr = ptxdata;
+	char *funcptr;
+	char *tempptr = ptxdata - 1;
+	char *lineptr = ptxdata - 1;
+
+	// recreate the ptx file without duplications
+	for (	std::map<unsigned,const char*>::iterator iter = duplicate.begin();
+			iter != duplicate.end();
+			iter++){
+		// find the line of the next error
+		linenum = iter->first;
+		for (int i = oldlinenum; i < linenum; i++) {
+			lineptr = strchr(lineptr + 1, '\n');
+		}
+		
+		// find the end of the current section to be copied over
+		// then find the start of the next section that will be copied
+		if (strcmp("function", iter->second) == 0) {
+			// get location of most recent .func
+			while (tempptr < lineptr && tempptr != NULL) {
+				funcptr = tempptr;
+				tempptr = strstr(funcptr + 1, ".func");
+			}
+
+			// get the start of the previous line
+			offset = 0;
+			while (*(funcptr - offset) != '\n') offset++;
+
+			fwrite(startptr, sizeof(char), funcptr - offset + 1 - startptr, ptxdest);
+
+			//find next location of startptr
+			if (*(lineptr + 3) == ';') {
+				// for function definitions
+				startptr = lineptr + 5;
+			} else if (*(lineptr + 3) == '{') {
+				// for functions enclosed with curly brackets
+				offset = 5;
+				unsigned bracket = 1;
+				while (bracket != 0) {
+					if (*(lineptr + offset) == '{') bracket++;
+					else if (*(lineptr + offset) == '}') bracket--;
+					offset++;
+				}
+				startptr = lineptr + offset + 1;
+			} else {
+				printf("GPGPU-Sim PTX: ERROR ** Unrecognized function format\n");
+				abort();
+			}
+		} else if (strcmp("variable", iter->second) == 0) {
+			fwrite(startptr, sizeof(char), (int)(lineptr + 1 - startptr), ptxdest);
+			
+			//find next location of startptr
+			offset = 1;
+			while (*(lineptr + offset) != '\n') offset++;
+			startptr = lineptr + offset + 1;
+		} else {
+			printf("GPGPU-Sim PTX: ERROR ** Unsupported duplicate type: %s\n", iter->second);
+		}
+
+		oldlinenum = linenum;
+	}
+	// copy over the rest of the file
+	fwrite(startptr, sizeof(char), ptxdata + filesize - startptr, ptxdest);
+
+	// cleanup
+	free(ptxdata);
+	fclose(ptxdest);
+	snprintf(commandline,1024,"rm -f %s",tempfile);
+	printf("Running: %s\n", commandline);
+	result = system(commandline);
+	if (result != 0) {
+		printf("GPGPU-Sim PTX: ERROR ** while deleting %s", tempfile);
+		exit(1);
+	}
+}
+
 void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version )
 {
     char fname[1024];
@@ -225,14 +330,29 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
     printf("GPGPU-Sim PTX: generating ptxinfo using \"%s\"\n", commandline);
     result = system(commandline);
     if( result != 0 ) {
-       printf("GPGPU-Sim PTX: ERROR ** while loading PTX (b) %d\n", result);
-       printf("               Ensure ptxas is in your path.\n");
-       exit(1);
+    	// 65280 = duplicate errors
+    	if (result == 65280) {
+    		ptxinfo_in = fopen(tempfile_ptxinfo,"r");
+		g_ptxinfo_filename = tempfile_ptxinfo;
+		ptxinfo_parse();
+
+    		fix_duplicate_errors(fname2);
+    		snprintf(commandline,1024,"$CUDA_INSTALL_PATH/bin/ptxas %s -v %s --output-file  /dev/null 2> %s",
+    			 extra_flags, fname2, tempfile_ptxinfo);
+    		printf("GPGPU-Sim PTX: regenerating ptxinfo using \"%s\"\n", commandline);
+    		result = system(commandline);
+	}
+	if (result != 0) {
+		printf("GPGPU-Sim PTX: ERROR ** while loading PTX (b) %d\n", result);
+		printf("               Ensure ptxas is in your path.\n");
+		exit(1);
+	}
     }
 
     ptxinfo_in = fopen(tempfile_ptxinfo,"r");
     g_ptxinfo_filename = tempfile_ptxinfo;
     ptxinfo_parse();
+
     if( ! g_save_embedded_ptx ) {
         snprintf(commandline,1024,"rm -f %s %s %s", fname, fname2, tempfile_ptxinfo);
         printf("GPGPU-Sim PTX: removing ptxinfo using \"%s\"\n", commandline);
diff --git a/src/cuda-sim/ptx_sim.cc b/src/cuda-sim/ptx_sim.cc
index 9b32e68..09844ae 100644
--- a/src/cuda-sim/ptx_sim.cc
+++ b/src/cuda-sim/ptx_sim.cc
@@ -240,7 +240,7 @@ unsigned ptx_thread_info::get_builtin( int builtin_id, unsigned dim_mod )
    }
    case GRIDID_REG:
       return m_gridid;
-   case LANEID_REG: feature_not_implemented( "%laneid" ); return 0;
+   case LANEID_REG: return get_hw_tid() % m_core->get_warp_size();
    case LANEMASK_EQ_REG: feature_not_implemented( "%lanemask_eq" ); return 0;
    case LANEMASK_LE_REG: feature_not_implemented( "%lanemask_le" ); return 0;
    case LANEMASK_LT_REG: feature_not_implemented( "%lanemask_lt" ); return 0;
diff --git a/src/cuda-sim/ptxinfo.l b/src/cuda-sim/ptxinfo.l
index 99ee1fc..33c2748 100644
--- a/src/cuda-sim/ptxinfo.l
+++ b/src/cuda-sim/ptxinfo.l
@@ -59,6 +59,11 @@ unsigned ptxinfo_col = 0;
 "gmem"		TC; return GMEM;
 "line"		TC; return LINE;
 "for"		TC; return FOR;
+"textures"	TC; return TEXTURES;
+"error   : Duplicate definition of"	TC; return DUPLICATE;
+"function"	TC; ptxinfo_lval.string_value = strdup(yytext); return FUNCTION;
+"variable"	TC; ptxinfo_lval.string_value = strdup(yytext); return VARIABLE;
+"fatal   : Ptx assembly aborted due to errors"	TC; return FATAL;
 
 [_A-Za-z$%][_0-9A-Za-z$]*  TC; ptxinfo_lval.string_value = strdup(yytext); return IDENTIFIER;
 [-]{0,1}[0-9]+	 TC; ptxinfo_lval.int_value =  atoi(yytext); return INT_OPERAND;
diff --git a/src/cuda-sim/ptxinfo.y b/src/cuda-sim/ptxinfo.y
index 294412d..37092f4 100644
--- a/src/cuda-sim/ptxinfo.y
+++ b/src/cuda-sim/ptxinfo.y
@@ -54,11 +54,16 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 %token LINE
 %token <string_value> WARNING
 %token FOR
+%token TEXTURES
+%token DUPLICATE
+%token <string_value> FUNCTION
+%token <string_value> VARIABLE
+%token FATAL
 
 %{
 	#include <stdlib.h>
 	#include <string.h>
-
+	
 	static unsigned g_declared;
 	static unsigned g_system;
 	int ptxinfo_lex(void);
@@ -70,6 +75,8 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 	void ptxinfo_smem( unsigned declared, unsigned system );
 	void ptxinfo_cmem( unsigned nbytes, unsigned bank );
 	int ptxinfo_error(const char*);
+	void ptxinfo_linenum( unsigned );
+	void ptxinfo_dup_type( const char* );
 %}
 
 %%
@@ -81,6 +88,8 @@ input:	/* empty */
 line: 	HEADER INFO COLON line_info
 	| HEADER IDENTIFIER COMMA LINE INT_OPERAND SEMICOLON WARNING
 	| HEADER WARNING { printf("GPGPU-Sim: ptxas %s\n", $2); }
+	| HEADER IDENTIFIER COMMA LINE INT_OPERAND SEMICOLON DUPLICATE duplicate { ptxinfo_linenum($5); }
+	| HEADER FATAL
 	;
 
 line_info: function_name
@@ -88,7 +97,7 @@ line_info: function_name
 	;
 
 function_name:	FUNC QUOTE IDENTIFIER QUOTE { ptxinfo_function($3); }
-	|  FUNC QUOTE IDENTIFIER QUOTE FOR QUOTE IDENTIFIER QUOTE {ptxinfo_function($3); }
+	|  FUNC QUOTE IDENTIFIER QUOTE FOR QUOTE IDENTIFIER QUOTE { ptxinfo_function($3); }
 	;
 	
 function_info: info
@@ -104,10 +113,15 @@ info: 	  USED INT_OPERAND REGS { ptxinfo_regs($2); }
 	| INT_OPERAND BYTES SMEM { ptxinfo_smem($1,0); }
 	| INT_OPERAND BYTES CMEM { ptxinfo_cmem($1,0); }
 	| INT_OPERAND REGS { ptxinfo_regs($1); }
+	| INT_OPERAND TEXTURES {}
 	;
 
 tuple: INT_OPERAND PLUS INT_OPERAND BYTES { g_declared=$1; g_system=$3; }
 
+duplicate:	FUNCTION QUOTE IDENTIFIER QUOTE { ptxinfo_dup_type($1); }
+	| VARIABLE QUOTE IDENTIFIER QUOTE { ptxinfo_dup_type($1); }
+	;
+
 %%