From 041192e93f61bf27c7f7928d87afe5623cc08728 Mon Sep 17 00:00:00 2001 From: speverel Date: Wed, 1 Jun 2016 11:26:47 -0700 Subject: Added GTX750Ti config files. Also modified shader.h to support larger number of warps per CTA (64), which is necessary for Maxwell boards like the GTX750Ti. --- src/gpgpu-sim/shader.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index 1aa468b..fe00770 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -267,7 +267,7 @@ private: inline unsigned hw_tid_from_wid(unsigned wid, unsigned warp_size, unsigned i){return wid * warp_size + i;}; inline unsigned wid_from_hw_tid(unsigned tid, unsigned warp_size){return tid/warp_size;}; -const unsigned WARP_PER_CTA_MAX = 48; +const unsigned WARP_PER_CTA_MAX = 64; typedef std::bitset warp_set_t; int register_bank(int regnum, int wid, unsigned num_banks, unsigned bank_warp_shift); -- cgit v1.3 From 88e1bcdaa3271fb79c999389c82733736154e48a Mon Sep 17 00:00:00 2001 From: speverel Date: Wed, 1 Jun 2016 11:33:01 -0700 Subject: Added new GTX750Ti config files. --- configs/GeForceGTX750Ti/config_fermi_islip.icnt | 70 +++ configs/GeForceGTX750Ti/gpgpusim.config | 128 ++++++ configs/GeForceGTX750Ti/gpuwattch_gtx750Ti.xml | 538 ++++++++++++++++++++++++ 3 files changed, 736 insertions(+) create mode 100644 configs/GeForceGTX750Ti/config_fermi_islip.icnt create mode 100644 configs/GeForceGTX750Ti/gpgpusim.config create mode 100755 configs/GeForceGTX750Ti/gpuwattch_gtx750Ti.xml diff --git a/configs/GeForceGTX750Ti/config_fermi_islip.icnt b/configs/GeForceGTX750Ti/config_fermi_islip.icnt new file mode 100644 index 0000000..069ca02 --- /dev/null +++ b/configs/GeForceGTX750Ti/config_fermi_islip.icnt @@ -0,0 +1,70 @@ +//21*1 fly with 32 flits per packet under gpgpusim injection mode +use_map = 0; +flit_size = 32; + +// currently we do not use this, see subnets below +network_count = 2; + +// Topology +topology = fly; +k = 7; +n = 1; + +// Routing + +routing_function = dest_tag; + +// Flow control + +num_vcs = 1; +vc_buf_size = 8; + +wait_for_tail_credit = 0; + +// Router architecture + +vc_allocator = islip; //separable_input_first; +sw_allocator = islip; //separable_input_first; +alloc_iters = 1; + +credit_delay = 0; +routing_delay = 0; +vc_alloc_delay = 1; +sw_alloc_delay = 1; + +input_speedup = 2; +output_speedup = 1; +internal_speedup = 1.0; + +// Traffic, GPGPU-Sim does not use this + +traffic = uniform; +packet_size ={{1,2,3,4},{10,20}}; +packet_size_rate={{1,1,1,1},{2,1}}; + +// Simulation - Don't change + +sim_type = gpgpusim; +//sim_type = latency; +injection_rate = 0.1; + +subnets = 2; + +// Always use read and write no matter following line +//use_read_write = 1; + + +read_request_subnet = 0; +read_reply_subnet = 1; +write_request_subnet = 0; +write_reply_subnet = 1; + +read_request_begin_vc = 0; +read_request_end_vc = 0; +write_request_begin_vc = 0; +write_request_end_vc = 0; +read_reply_begin_vc = 0; +read_reply_end_vc = 0; +write_reply_begin_vc = 0; +write_reply_end_vc = 0; + diff --git a/configs/GeForceGTX750Ti/gpgpusim.config b/configs/GeForceGTX750Ti/gpgpusim.config new file mode 100644 index 0000000..afd3825 --- /dev/null +++ b/configs/GeForceGTX750Ti/gpgpusim.config @@ -0,0 +1,128 @@ +# functional simulator specification +-gpgpu_ptx_instruction_classification 0 +-gpgpu_ptx_sim_mode 0 +-gpgpu_ptx_force_max_capability 20 + + +# SASS execution (only supported with CUDA >= 4.0) +-gpgpu_ptx_convert_to_ptxplus 0 +-gpgpu_ptx_save_converted_ptxplus 0 + +# high level architecture configuration +-gpgpu_n_clusters 5 +-gpgpu_n_cores_per_cluster 1 +-gpgpu_n_mem 2 +-gpgpu_n_sub_partition_per_mchannel 1 + +# Maxwell clock domains +#-gpgpu_clock_domains ::: +# In Maxwell, each pipeline has 32 execution units, so the Core clock is just the reported number. +-gpgpu_clock_domains 1137.0:1137.0:1137.0:2700.0 + +# shader core pipeline config +-gpgpu_shader_registers 65536 + +# This implies a maximum of 64 warps/SM +-gpgpu_shader_core_pipeline 2048:32 +-gpgpu_shader_cta 8 +-gpgpu_simd_model 1 + +# Pipeline widths and number of FUs +# ID_OC_SP,ID_OC_SFU,ID_OC_MEM,OC_EX_SP,OC_EX_SFU,OC_EX_MEM,EX_WB +-gpgpu_pipeline_widths 2,1,1,2,1,1,2 +-gpgpu_num_sp_units 8 +-gpgpu_num_sfu_units 32 + +# Instruction latencies and initiation intervals +# "ADD,MAX,MUL,MAD,DIV" +-ptx_opcode_latency_int 6,12,13,13,210 +-ptx_opcode_initiation_int 1,1,1,1,4 +-ptx_opcode_latency_fp 6,12,6,6,374 +-ptx_opcode_initiation_fp 1,1,1,1,4 +-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_shmem_size 65536 + +# The alternative configuration for fermi in case cudaFuncCachePreferL1 is selected +#-gpgpu_cache:dl1 64:128:6,L:L:m:N:H,A:32:8,8 +#-gpgpu_shmem_size 16384 + +# 1024 sets, each 128 bytes 8-way for each memory sub partition. This gives 2MB L2 cache +-gpgpu_cache:dl2 1024:128:8,L:B:m:W:L,A:32:4,4:0,32 +-gpgpu_cache:dl2_texture_only 0 + +-gpgpu_cache:il1 4:128:4,L:R:f:N:L,A:2:32,4 +-gpgpu_tex_cache:l1 16:128:32,L:R:m:N:L,F:128:4,128:2 +-gpgpu_const_cache:l1 64:64:2,L:R:f:N:L,A:2:32,4 + +# enable operand collector +-gpgpu_operand_collector_num_units_sp 6 +-gpgpu_operand_collector_num_units_sfu 8 +-gpgpu_operand_collector_num_in_ports_sp 2 +-gpgpu_operand_collector_num_out_ports_sp 2 +-gpgpu_num_reg_banks 16 + +# shared memory bankconflict detection +-gpgpu_shmem_num_banks 32 +-gpgpu_shmem_limited_broadcast 0 +-gpgpu_shmem_warp_parts 1 + +-gpgpu_max_insn_issue_per_warp 2 + +# interconnection +-network_mode 1 +-inter_config_file config_fermi_islip.icnt + +# memory partition latency config +-rop_latency 150 +-dram_latency 130 + +# dram model config +-gpgpu_dram_scheduler 1 +# The DRAM return queue and the scheduler queue together should provide buffer +# to sustain the memory level parallelism to tolerate DRAM latency +# To allow 100% DRAM utility, there should at least be enough buffer to sustain +# the minimum DRAM latency (100 core cycles). I.e. +# Total buffer space required = 100 x 924MHz / 700MHz = 132 +-gpgpu_frfcfs_dram_sched_queue_size 16 +-gpgpu_dram_return_queue_size 300 + +# for Maxwell, bus width is 128bits, this is 64 bytes (32 bytes at each DRAM chip) per memory partition +-gpgpu_n_mem_per_ctrlr 2 +-gpgpu_dram_buswidth 32 +-gpgpu_dram_burst_length 8 +-dram_data_command_freq_ratio 4 # GDDR5 is QDR +-gpgpu_mem_address_mask 1 +-gpgpu_mem_addr_mapping dramid@8;00000000.00000000.00000000.00000000.0000RRRR.RRRRRRRR.BBBCCCCB.CCSSSSSS + +# GDDR5 timing from hynix H5GQ1H24AFR +# to disable bank groups, set nbkgrp to 1 and tCCDL and tRTPL to 0 +-gpgpu_dram_timing_opt "nbk=16:CCD=2:RRD=6:RCD=12:RAS=28:RP=12:RC=40: + CL=12:WL=4:CDLR=5:WR=12:nbkgrp=4:CCDL=3:RTPL=2" + +# Maxwell has four schedulers per core +-gpgpu_num_sched_per_core 4 +# Two Level Scheduler with active and pending pools +#-gpgpu_scheduler two_level_active:6:0:1 +# Loose round robbin scheduler +#-gpgpu_scheduler lrr +# Greedy then oldest scheduler +-gpgpu_scheduler gto + +# stat collection +-gpgpu_memlatency_stat 14 +-gpgpu_runtime_stat 500 +-enable_ptx_file_line_stats 1 +-visualizer_enabled 0 + +# power model configs +-power_simulation_enabled 0 +-gpuwattch_xml_file gpuwattch_gtx480.xml + +# tracing functionality +#-trace_enabled 1 +#-trace_components WARP_SCHEDULER,SCOREBOARD +#-trace_sampling_core 0 diff --git a/configs/GeForceGTX750Ti/gpuwattch_gtx750Ti.xml b/configs/GeForceGTX750Ti/gpuwattch_gtx750Ti.xml new file mode 100755 index 0000000..304e0fd --- /dev/null +++ b/configs/GeForceGTX750Ti/gpuwattch_gtx750Ti.xml @@ -0,0 +1,538 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + -- cgit v1.3 From 4cbb9730252dadcde68322cedbe004676420e9e6 Mon Sep 17 00:00:00 2001 From: speverel 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. --- configs/GeForceGTX750Ti/gpgpusim.config | 4 ++-- libcuda/cuda_runtime_api.cc | 8 +++++++- 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 ++ 10 files changed, 21 insertions(+), 7 deletions(-) diff --git a/configs/GeForceGTX750Ti/gpgpusim.config b/configs/GeForceGTX750Ti/gpgpusim.config index afd3825..be3ae3c 100644 --- a/configs/GeForceGTX750Ti/gpgpusim.config +++ b/configs/GeForceGTX750Ti/gpgpusim.config @@ -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 diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 910bebd..e2626d2 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -109,6 +109,7 @@ #include #include #include +#include #include #include #ifdef OPENGL_SUPPORT @@ -1816,10 +1817,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/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 WARNING %token FOR +%token TEXTURES %{ #include @@ -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 46aad91327a265c2fea2cfe629cc38eadb629200 Mon Sep 17 00:00:00 2001 From: speverel 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. --- configs/GeForceGTX750Ti/gpgpusim.config | 4 ++-- libcuda/cuda_runtime_api.cc | 8 +++++++- 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 ++ 10 files changed, 21 insertions(+), 7 deletions(-) diff --git a/configs/GeForceGTX750Ti/gpgpusim.config b/configs/GeForceGTX750Ti/gpgpusim.config index afd3825..be3ae3c 100644 --- a/configs/GeForceGTX750Ti/gpgpusim.config +++ b/configs/GeForceGTX750Ti/gpgpusim.config @@ -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 diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 910bebd..e2626d2 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -109,6 +109,7 @@ #include #include #include +#include #include #include #ifdef OPENGL_SUPPORT @@ -1816,10 +1817,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/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 WARNING %token FOR +%token TEXTURES %{ #include @@ -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 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(-) 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 a2c48c472fc2ab159b826f3af1d1503c23eaed39 Mon Sep 17 00:00:00 2001 From: sspenst Date: Fri, 3 Jun 2016 14:10:20 -0700 Subject: Added mergeSectionList to combine any PTX files that remain after pruning into a single file. --- libcuda/cuda_runtime_api.cc | 44 ++++++++++++++++++++++++++++++++++++++++++++ libcuda/cuobjdump.y | 2 +- 2 files changed, 45 insertions(+), 1 deletion(-) diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index e2626d2..8049f43 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1496,6 +1496,49 @@ std::list pruneSectionList(std::list cuobj return prunedList; } +//! Merge all remaining PTX sections in the section list into one PTX file +// NOTE: this function needs some tweaking to deal with repeated fucntion/variable declarations/definitions +std::list mergeSectionList(std::list cuobjdumpSectionList){ + char *ptxcode = ""; + std::list::iterator old_iter; + cuobjdumpPTXSection* old_ptxsection = NULL; + cuobjdumpPTXSection* ptxsection; + std::list mergedList; + + for ( std::list::iterator iter = cuobjdumpSectionList.begin(); + iter != cuobjdumpSectionList.end(); + iter++){ + if((ptxsection=dynamic_cast(*iter)) != NULL){ + // 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 + else { + mergedList.push_back(*iter); + } + } + + // Store the final PTX section + mergedList.push_front(*old_iter); + + return mergedList; +} + //! Within the section list, find the ELF section corresponding to a given identifier cuobjdumpELFSection* findELFSectionInList(std::list sectionlist, const std::string identifier){ @@ -1559,6 +1602,7 @@ void cuobjdumpInit(){ CUctx_st *context = GPGPUSim_Context(); extract_code_using_cuobjdump(); //extract all the output of cuobjdump to _cuobjdump_*.* cuobjdumpSectionList = pruneSectionList(cuobjdumpSectionList, context); + cuobjdumpSectionList = mergeSectionList(cuobjdumpSectionList); } std::map fatbinmap; diff --git a/libcuda/cuobjdump.y b/libcuda/cuobjdump.y index dce7e3d..9d61f25 100644 --- a/libcuda/cuobjdump.y +++ b/libcuda/cuobjdump.y @@ -114,7 +114,7 @@ compressedkeyword : H_COMPRESSED emptylines ptxcode : ptxcode PTXLINE {fprintf(ptxfile, "%s", $2);} | ; -elfcode : elfcode ELFLINE {printf(elffile, "%s", $2);} +elfcode : elfcode ELFLINE {fprintf(elffile, "%s", $2);} | ; sasscode : sasscode SASSLINE {fprintf(sassfile, "%s", $2);} -- cgit v1.3 From 71ccce6074de5d7eef3fbe2cda1dbca008549f07 Mon Sep 17 00:00:00 2001 From: speverel 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(-) 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 @@ + 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 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(-) 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 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(-) 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 #include @@ -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 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 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(-) 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 function_decl_header %type 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 41dbdd842f315645cc4cd05e3b6263b982d1be2e Mon Sep 17 00:00:00 2001 From: sspenst Date: Tue, 7 Jun 2016 16:43:09 -0700 Subject: Updated the PTX section merging to be able to deal with multiple identifiers --- libcuda/cuda_runtime_api.cc | 43 ++++++++++++++++++++++++++++++++++++------- 1 file changed, 36 insertions(+), 7 deletions(-) diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 018b387..862e2b9 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1513,9 +1513,8 @@ std::list pruneSectionList(std::list cuobj return prunedList; } -//! Merge all remaining PTX sections in the section list into one PTX file -// NOTE: this function needs some tweaking to deal with repeated fucntion/variable declarations/definitions -std::list mergeSectionList(std::list cuobjdumpSectionList){ +//! Merge all PTX sections that have a specific identifier into one file +std::list mergeMatchingSections(std::list cuobjdumpSectionList, std::string identifier){ char *ptxcode = ""; std::list::iterator old_iter; cuobjdumpPTXSection* old_ptxsection = NULL; @@ -1525,7 +1524,8 @@ std::list mergeSectionList(std::list cuobj for ( std::list::iterator iter = cuobjdumpSectionList.begin(); iter != cuobjdumpSectionList.end(); iter++){ - if((ptxsection=dynamic_cast(*iter)) != NULL){ + if((ptxsection=dynamic_cast(*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()); @@ -1544,18 +1544,47 @@ std::list mergeSectionList(std::list cuobj old_iter = iter; old_ptxsection = ptxsection; } - // Store all non-PTX sections + // Store all non-PTX sections and PTX sections with non-matching identifiers else { mergedList.push_back(*iter); } } // Store the final PTX section - mergedList.push_front(*old_iter); + mergedList.push_back(*old_iter); return mergedList; } +//! Merge any PTX sections with matching identifiers +std::list mergeSections(std::list cuobjdumpSectionList){ + std::vector identifier; + cuobjdumpPTXSection* ptxsection; + + // Add all identifiers present in PTX sections to a vector + for ( std::list::iterator iter = cuobjdumpSectionList.begin(); + iter != cuobjdumpSectionList.end(); + iter++){ + if((ptxsection=dynamic_cast(*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::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 sectionlist, const std::string identifier){ @@ -1619,7 +1648,7 @@ void cuobjdumpInit(){ CUctx_st *context = GPGPUSim_Context(); extract_code_using_cuobjdump(); //extract all the output of cuobjdump to _cuobjdump_*.* cuobjdumpSectionList = pruneSectionList(cuobjdumpSectionList, context); - cuobjdumpSectionList = mergeSectionList(cuobjdumpSectionList); + cuobjdumpSectionList = mergeSections(cuobjdumpSectionList); } std::map fatbinmap; -- cgit v1.3 From 5e9b6521e625d30f9f3efbb292fd22bdab5c884a Mon Sep 17 00:00:00 2001 From: sspenst Date: Thu, 9 Jun 2016 13:38:13 -0700 Subject: Added a basic implementation of cudaStreamCreateWithFlags --- libcuda/cuda_runtime_api.cc | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 862e2b9..77456d4 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -984,6 +984,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) { + cudaStreamCreate(stream); +} + __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream) { #if (CUDART_VERSION >= 3000) -- cgit v1.3 From adabbc7070704d1be79465786d59e170030c7b0d Mon Sep 17 00:00:00 2001 From: sspenst Date: Thu, 9 Jun 2016 14:23:28 -0700 Subject: Forgot to add a return value --- libcuda/cuda_runtime_api.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 77456d4..4cc002d 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -985,7 +985,7 @@ __host__ cudaError_t CUDARTAPI cudaStreamCreate(cudaStream_t *stream) } __host__ __device__ cudaError_t CUDARTAPI cudaStreamCreateWithFlags(cudaStream_t *stream, unsigned int flags) { - cudaStreamCreate(stream); + return cudaStreamCreate(stream); } __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream) -- cgit v1.3 From ac11fd57b93e22ee5a50ebf8f6d6b4d6dadbe3cb Mon Sep 17 00:00:00 2001 From: speverel Date: Thu, 9 Jun 2016 14:58:38 -0700 Subject: Modified runtime API to support synchronization with the null stream. --- libcuda/cuda_runtime_api.cc | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 77456d4..350ba0b 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -440,6 +440,10 @@ extern "C" { * * * * *******************************************************************************/ +cudaError_t cudaPeekAtLastError(void) +{ + return g_last_cudaError; +} __host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size) { @@ -1000,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__); -- cgit v1.3 From 7356669b66425ee73ba61466f37fe25b886a3641 Mon Sep 17 00:00:00 2001 From: sspenst Date: Fri, 10 Jun 2016 14:32:00 -0700 Subject: Vastly improved runtime by relating symbol tables to file names --- libcuda/cuda_runtime_api.cc | 14 ++++++++++---- 1 file changed, 10 insertions(+), 4 deletions(-) diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 5558c4b..e8a0e91 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1662,6 +1662,7 @@ void cuobjdumpInit(){ std::map fatbinmap; std::mapfatbin_registered; +std::map name_symtab; //! Keep track of the association between filename and cubin handle void cuobjdumpRegisterFatBinary(unsigned int handle, char* filename){ @@ -1674,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::iterator iter = cuobjdumpSectionList.begin(); @@ -1682,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); @@ -1717,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 } -- cgit v1.3 From 547ce656018a6439e658be3c283721a961b0217f Mon Sep 17 00:00:00 2001 From: sspenst 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(-) 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 get_duplicate(); + int g_ptxinfo_error_detected; static char *g_ptxinfo_kname = NULL; static struct gpgpu_ptx_sim_info g_ptxinfo; +static std::map 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 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 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 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::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 WARNING %token FOR %token TEXTURES +%token DUPLICATE +%token FUNCTION +%token VARIABLE +%token FATAL %{ #include #include - + 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 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(-) 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; { @@ -390,7 +392,9 @@ breakaddr TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE; } { "*/" 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 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(-) 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; { -- cgit v1.3