From 46aad91327a265c2fea2cfe629cc38eadb629200 Mon Sep 17 00:00:00 2001
From: speverel <speverelle8@gmail.com>
Date: Thu, 2 Jun 2016 11:28:15 -0700
Subject: Added handling of .cc option for arithmetic instructions. NOTE: Only
 made changes to parse instructions. Carry functionality NOT fully
 implemented; .cc instructions function like their unmodified ueqivelents.
 Also modified GTX750Ti config to model L1 data cache as simply not being used
 for global loads (instead of not existing at all). Changed ptxinfo parsing to
 avoid crashing when info includes texture information.

---
 src/cuda-sim/cuda-sim.cc     | 2 +-
 src/cuda-sim/instructions.cc | 2 +-
 src/cuda-sim/opcodes.def     | 2 +-
 src/cuda-sim/ptx.l           | 3 ++-
 src/cuda-sim/ptx.y           | 2 ++
 src/cuda-sim/ptx_ir.cc       | 2 ++
 src/cuda-sim/ptxinfo.l       | 1 +
 src/cuda-sim/ptxinfo.y       | 2 ++
 8 files changed, 12 insertions(+), 4 deletions(-)

(limited to 'src')

diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 715be98..15417d1 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:
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index cf7f04a..4dd5ed8 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -2421,7 +2421,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..95ab74c 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;
@@ -249,6 +249,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;
diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y
index 79faddf..82abcbb 100644
--- a/src/cuda-sim/ptx.y
+++ b/src/cuda-sim/ptx.y
@@ -103,6 +103,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
@@ -438,6 +439,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); }
diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc
index 751b3f4..8f9c3d2 100644
--- a/src/cuda-sim/ptx_ir.cc
+++ b/src/cuda-sim/ptx_ir.cc
@@ -1167,6 +1167,8 @@ ptx_instruction::ptx_instruction( int opcode,
       case HALF_OPTION:
          m_inst_size = 4; // bytes
          break;
+	  case EXTP_OPTION:
+		 break;
       default:
          assert(0);
          break;
diff --git a/src/cuda-sim/ptxinfo.l b/src/cuda-sim/ptxinfo.l
index 99ee1fc..f9b6846 100644
--- a/src/cuda-sim/ptxinfo.l
+++ b/src/cuda-sim/ptxinfo.l
@@ -59,6 +59,7 @@ unsigned ptxinfo_col = 0;
 "gmem"		TC; return GMEM;
 "line"		TC; return LINE;
 "for"		TC; return FOR;
+"textures"	TC; return TEXTURES;
 
 [_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..faa33eb 100644
--- a/src/cuda-sim/ptxinfo.y
+++ b/src/cuda-sim/ptxinfo.y
@@ -54,6 +54,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 %token LINE
 %token <string_value> WARNING
 %token FOR
+%token TEXTURES
 
 %{
 	#include <stdlib.h>
@@ -104,6 +105,7 @@ 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; }
-- 
cgit v1.3


From 068e34728a0706cd8671b816eea466491ab2db06 Mon Sep 17 00:00:00 2001
From: speverel <speverelle8@gmail.com>
Date: Thu, 2 Jun 2016 13:29:46 -0700
Subject: Updated parser and config file to support compute versions up to 5.2.
 Full support is NOT claimed; however, it has been tested to work on a number
 of CUDA version 7.5 benchmarks such as matrix multiply and simpleMultiGPU.

---
 configs/GeForceGTX750Ti/gpgpusim.config | 2 +-
 src/cuda-sim/ptx.l                      | 1 +
 src/cuda-sim/ptx.y                      | 2 ++
 3 files changed, 4 insertions(+), 1 deletion(-)

(limited to 'src')

diff --git a/configs/GeForceGTX750Ti/gpgpusim.config b/configs/GeForceGTX750Ti/gpgpusim.config
index be3ae3c..7920763 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)
diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l
index 95ab74c..66ff48f 100644
--- a/src/cuda-sim/ptx.l
+++ b/src/cuda-sim/ptx.l
@@ -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 */
diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y
index 82abcbb..fca94db 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
@@ -243,6 +244,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"); }
 	;
 
-- 
cgit v1.3


From 71ccce6074de5d7eef3fbe2cda1dbca008549f07 Mon Sep 17 00:00:00 2001
From: speverel <speverelle8@gmail.com>
Date: Fri, 3 Jun 2016 15:24:40 -0700
Subject: Added support for %laneid SFR. Also added a notice clarifying that
 power modeling for GTX750Ti is currently completely untested and should not
 be considered supported.

---
 configs/GeForceGTX750Ti/gpgpusim.config        | 4 +++-
 configs/GeForceGTX750Ti/gpuwattch_gtx750Ti.xml | 1 +
 src/cuda-sim/ptx_sim.cc                        | 2 +-
 3 files changed, 5 insertions(+), 2 deletions(-)

(limited to 'src')

diff --git a/configs/GeForceGTX750Ti/gpgpusim.config b/configs/GeForceGTX750Ti/gpgpusim.config
index 7920763..8b030b6 100644
--- a/configs/GeForceGTX750Ti/gpgpusim.config
+++ b/configs/GeForceGTX750Ti/gpgpusim.config
@@ -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/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;
-- 
cgit v1.3


From 8ed1522aede92bcafc60f69d7e03e6d48c44a86c Mon Sep 17 00:00:00 2001
From: speverel <speverelle8@gmail.com>
Date: Mon, 6 Jun 2016 14:48:13 -0700
Subject: Added support for BFE (Bit field extract) instruction.

---
 src/cuda-sim/instructions.cc | 77 +++++++++++++++++++++++++++++++++++++++++++-
 1 file changed, 76 insertions(+), 1 deletion(-)

(limited to 'src')

diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index 4dd5ed8..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); }
 
-- 
cgit v1.3


From 7aeadc95cc50d266f93cdb3ada1c192d9b5a1046 Mon Sep 17 00:00:00 2001
From: speverel <speverelle8@gmail.com>
Date: Tue, 7 Jun 2016 14:24:24 -0700
Subject: Added support for cudaMemcpyDefault flag in cudaMemcpy. Also
 increased the maximum allowable memory to 2GB and the compute version to 5.2.

---
 libcuda/cuda_runtime_api.cc   | 23 ++++++++++++++++++++---
 src/abstract_hardware_model.h |  2 +-
 2 files changed, 21 insertions(+), 4 deletions(-)

(limited to 'src')

diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 8049f43..018b387 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -134,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>
@@ -320,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;
@@ -533,6 +534,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();
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)
-- 
cgit v1.3


From 587853a81f6fa6088b7f3d93fc8862a8b2610da7 Mon Sep 17 00:00:00 2001
From: sspenst <spencerspenst@gmail.com>
Date: Tue, 7 Jun 2016 16:22:45 -0700
Subject: The ptx parser now recognizes the NC option for ld.global, however
 this option is not actually implemented

---
 libcuda/cuobjdump.y    | 3 +--
 src/cuda-sim/ptx.l     | 2 ++
 src/cuda-sim/ptx.y     | 2 ++
 src/cuda-sim/ptx_ir.cc | 2 ++
 4 files changed, 7 insertions(+), 2 deletions(-)

(limited to 'src')

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/cuda-sim/ptx.l b/src/cuda-sim/ptx.l
index 66ff48f..a44177b 100644
--- a/src/cuda-sim/ptx.l
+++ b/src/cuda-sim/ptx.l
@@ -326,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;
diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y
index fca94db..4de39d1 100644
--- a/src/cuda-sim/ptx.y
+++ b/src/cuda-sim/ptx.y
@@ -193,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
@@ -449,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 8f9c3d2..2eccabc 100644
--- a/src/cuda-sim/ptx_ir.cc
+++ b/src/cuda-sim/ptx_ir.cc
@@ -1169,6 +1169,8 @@ ptx_instruction::ptx_instruction( int opcode,
          break;
 	  case EXTP_OPTION:
 		 break;
+	  case NC_OPTION:
+		 break;
       default:
          assert(0);
          break;
-- 
cgit v1.3


From 547ce656018a6439e658be3c283721a961b0217f Mon Sep 17 00:00:00 2001
From: sspenst <spencerspenst@gmail.com>
Date: Mon, 13 Jun 2016 01:08:36 -0700
Subject: If ptxas notices any duplicate errors, they now automatically get
 resolved and the program continues with the duplicate function/variable
 declarations removed.

---
 src/cuda-sim/cuda-sim.cc   |  19 +++++++
 src/cuda-sim/ptx_loader.cc | 126 +++++++++++++++++++++++++++++++++++++++++++--
 src/cuda-sim/ptxinfo.l     |   4 ++
 src/cuda-sim/ptxinfo.y     |  16 +++++-
 4 files changed, 160 insertions(+), 5 deletions(-)

(limited to 'src')

diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 15417d1..fba3a59 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -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/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/ptxinfo.l b/src/cuda-sim/ptxinfo.l
index f9b6846..33c2748 100644
--- a/src/cuda-sim/ptxinfo.l
+++ b/src/cuda-sim/ptxinfo.l
@@ -60,6 +60,10 @@ unsigned ptxinfo_col = 0;
 "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 faa33eb..37092f4 100644
--- a/src/cuda-sim/ptxinfo.y
+++ b/src/cuda-sim/ptxinfo.y
@@ -55,11 +55,15 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 %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);
@@ -71,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* );
 %}
 
 %%
@@ -82,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
@@ -89,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
@@ -110,6 +118,10 @@ info: 	  USED INT_OPERAND REGS { ptxinfo_regs($2); }
 
 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); }
+	;
+
 %%
 
 
-- 
cgit v1.3


From 281798191f9bc37a75592d34a5e38cc5d6c41b6d Mon Sep 17 00:00:00 2001
From: speverel <speverelle8@gmail.com>
Date: Thu, 16 Jun 2016 15:51:17 -0700
Subject: Added the ability to inject arbitrary PTX instructions.              
           This will be used to add custom instructions in the future; the
 imaginary instructions 'spr' and 'ama' have been added as samples.

---
 src/cuda-sim/instructions.cc | 10 ++++++++++
 src/cuda-sim/opcodes.def     |  2 ++
 src/cuda-sim/ptx.l           |  6 +++++-
 3 files changed, 17 insertions(+), 1 deletion(-)

(limited to 'src')

diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index 02ce01c..922e14a 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -816,6 +816,11 @@ void add_impl( const ptx_instruction *pI, ptx_thread_info *thread )
 
 void addc_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); }
 
+void ama_impl(  const ptx_instruction *pI, ptx_thread_info *thread )
+{
+	printf("AMA instruction found.\n");	
+}
+
 void and_impl( const ptx_instruction *pI, ptx_thread_info *thread ) 
 { 
    ptx_reg_t src1_data, src2_data, data;
@@ -3698,6 +3703,11 @@ void slct_impl( const ptx_instruction *pI, ptx_thread_info *thread )
    thread->set_operand_value(dst,d, i_type, thread, pI);
 }
 
+void spr_impl( const ptx_instruction *pI, ptx_thread_info *thread )
+{
+	printf("SPR instruction found.\n");
+}
+
 void sqrt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) 
 { 
    ptx_reg_t a, d;
diff --git a/src/cuda-sim/opcodes.def b/src/cuda-sim/opcodes.def
index 874acc7..33ee0ca 100644
--- a/src/cuda-sim/opcodes.def
+++ b/src/cuda-sim/opcodes.def
@@ -41,6 +41,7 @@ OP_DEF(ABS_OP,abs_impl,"abs",1,1)
 OP_DEF(ADD_OP,add_impl,"add",1,1)
 OP_DEF(ADDP_OP,addp_impl,"addp",1,1)
 OP_DEF(ADDC_OP,addc_impl,"addc",1,1)
+OP_DEF(AMA_OP,ama_impl,"ama",1,2)
 OP_DEF(AND_OP,and_impl,"and",1,1)
 OP_DEF(ANDN_OP,andn_impl,"andn",1,1)
 OP_DEF(ATOM_OP,atom_impl,"atom",1,3)
@@ -101,6 +102,7 @@ OP_DEF(SHL_OP,shl_impl,"shl",1,1)
 OP_DEF(SHR_OP,shr_impl,"shr",1,1)
 OP_DEF(SIN_OP,sin_impl,"sin",1,4)
 OP_DEF(SLCT_OP,slct_impl,"slct",1,1)
+OP_DEF(SPR_OP,spr_impl,"spr",1,1)
 OP_DEF(SQRT_OP,sqrt_impl,"sqrt",1,4)
 OP_DEF(SSY_OP,ssy_impl,"ssy",0,3)
 OP_DEF(ST_OP,st_impl,"st",0,5)
diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l
index a44177b..026270a 100644
--- a/src/cuda-sim/ptx.l
+++ b/src/cuda-sim/ptx.l
@@ -145,6 +145,8 @@ xor     TC; ptx_lval.int_value = XOR_OP; return OPCODE;
 nop     TC; ptx_lval.int_value = NOP_OP; return OPCODE;
 break  TC; ptx_lval.int_value = BREAK_OP; return OPCODE;
 breakaddr  TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE;
+spr TC; ptx_lval.int_value = SPR_OP; return OPCODE;
+ama TC; ptx_lval.int_value = AMA_OP; return OPCODE;
 
 <INITIAL,NOT_OPCODE,IN_INST,IN_FUNC_DECL>{
 
@@ -390,7 +392,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; 
 }
-- 
cgit v1.3


From f7c57e76c086ce417626f37ffc91097c839c687d Mon Sep 17 00:00:00 2001
From: sspenst <spencerspenst@gmail.com>
Date: Mon, 4 Jul 2016 09:43:36 -0700
Subject: Reverted part of the previous commit so that our new changes related
 to DNNs can be done in a different branch

---
 src/cuda-sim/instructions.cc | 10 ----------
 src/cuda-sim/opcodes.def     |  2 --
 src/cuda-sim/ptx.l           |  2 --
 3 files changed, 14 deletions(-)

(limited to 'src')

diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index 922e14a..02ce01c 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -816,11 +816,6 @@ void add_impl( const ptx_instruction *pI, ptx_thread_info *thread )
 
 void addc_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); }
 
-void ama_impl(  const ptx_instruction *pI, ptx_thread_info *thread )
-{
-	printf("AMA instruction found.\n");	
-}
-
 void and_impl( const ptx_instruction *pI, ptx_thread_info *thread ) 
 { 
    ptx_reg_t src1_data, src2_data, data;
@@ -3703,11 +3698,6 @@ void slct_impl( const ptx_instruction *pI, ptx_thread_info *thread )
    thread->set_operand_value(dst,d, i_type, thread, pI);
 }
 
-void spr_impl( const ptx_instruction *pI, ptx_thread_info *thread )
-{
-	printf("SPR instruction found.\n");
-}
-
 void sqrt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) 
 { 
    ptx_reg_t a, d;
diff --git a/src/cuda-sim/opcodes.def b/src/cuda-sim/opcodes.def
index 33ee0ca..874acc7 100644
--- a/src/cuda-sim/opcodes.def
+++ b/src/cuda-sim/opcodes.def
@@ -41,7 +41,6 @@ OP_DEF(ABS_OP,abs_impl,"abs",1,1)
 OP_DEF(ADD_OP,add_impl,"add",1,1)
 OP_DEF(ADDP_OP,addp_impl,"addp",1,1)
 OP_DEF(ADDC_OP,addc_impl,"addc",1,1)
-OP_DEF(AMA_OP,ama_impl,"ama",1,2)
 OP_DEF(AND_OP,and_impl,"and",1,1)
 OP_DEF(ANDN_OP,andn_impl,"andn",1,1)
 OP_DEF(ATOM_OP,atom_impl,"atom",1,3)
@@ -102,7 +101,6 @@ OP_DEF(SHL_OP,shl_impl,"shl",1,1)
 OP_DEF(SHR_OP,shr_impl,"shr",1,1)
 OP_DEF(SIN_OP,sin_impl,"sin",1,4)
 OP_DEF(SLCT_OP,slct_impl,"slct",1,1)
-OP_DEF(SPR_OP,spr_impl,"spr",1,1)
 OP_DEF(SQRT_OP,sqrt_impl,"sqrt",1,4)
 OP_DEF(SSY_OP,ssy_impl,"ssy",0,3)
 OP_DEF(ST_OP,st_impl,"st",0,5)
diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l
index 026270a..b8ce497 100644
--- a/src/cuda-sim/ptx.l
+++ b/src/cuda-sim/ptx.l
@@ -145,8 +145,6 @@ xor     TC; ptx_lval.int_value = XOR_OP; return OPCODE;
 nop     TC; ptx_lval.int_value = NOP_OP; return OPCODE;
 break  TC; ptx_lval.int_value = BREAK_OP; return OPCODE;
 breakaddr  TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE;
-spr TC; ptx_lval.int_value = SPR_OP; return OPCODE;
-ama TC; ptx_lval.int_value = AMA_OP; return OPCODE;
 
 <INITIAL,NOT_OPCODE,IN_INST,IN_FUNC_DECL>{
 
-- 
cgit v1.3