diff options
Diffstat (limited to 'src/trace-driven')
13 files changed, 8 insertions, 888 deletions
diff --git a/src/trace-driven/kepler_opcode.h b/src/trace-driven/ISA_Def/kepler_opcode.h index 4aa8e0f..675ea6c 100644 --- a/src/trace-driven/kepler_opcode.h +++ b/src/trace-driven/ISA_Def/kepler_opcode.h @@ -4,7 +4,6 @@ #ifndef KEPLER_OPCODE_H #define KEPLER_OPCODE_H -#include "../abstract_hardware_model.h" #include "trace_opcode.h" #include <unordered_map> #include <string> diff --git a/src/trace-driven/pascal_opcode.h b/src/trace-driven/ISA_Def/pascal_opcode.h index 2cacb28..66a0841 100644 --- a/src/trace-driven/pascal_opcode.h +++ b/src/trace-driven/ISA_Def/pascal_opcode.h @@ -4,7 +4,6 @@ #ifndef PASCAL_OPCODE_H #define PASCAL_OPCODE_H -#include "../abstract_hardware_model.h" #include "trace_opcode.h" #include <unordered_map> #include <string> diff --git a/src/trace-driven/trace_opcode.h b/src/trace-driven/ISA_Def/trace_opcode.h index 3492fd3..ed147fc 100644 --- a/src/trace-driven/trace_opcode.h +++ b/src/trace-driven/ISA_Def/trace_opcode.h @@ -4,7 +4,7 @@ #ifndef TRACE_OPCODE_H #define TRACE_OPCODE_H -#include "../abstract_hardware_model.h" +#include "../../abstract_hardware_model.h" #include <unordered_map> #include <string> diff --git a/src/trace-driven/turing_opcode.h b/src/trace-driven/ISA_Def/turing_opcode.h index f872070..4df44d9 100644 --- a/src/trace-driven/turing_opcode.h +++ b/src/trace-driven/ISA_Def/turing_opcode.h @@ -4,7 +4,6 @@ #ifndef TURING_OPCODE_H #define TURING_OPCODE_H -#include "../abstract_hardware_model.h" #include "trace_opcode.h" #include <unordered_map> #include <string> diff --git a/src/trace-driven/volta_opcode.h b/src/trace-driven/ISA_Def/volta_opcode.h index 3d03201..03d50b7 100644 --- a/src/trace-driven/volta_opcode.h +++ b/src/trace-driven/ISA_Def/volta_opcode.h @@ -4,7 +4,6 @@ #ifndef VOLTA_OPCODE_H #define VOLTA_OPCODE_H -#include "../abstract_hardware_model.h" #include "trace_opcode.h" #include <unordered_map> #include <string> diff --git a/src/trace-driven/gpgpusim_trace_driven_main.cc b/src/trace-driven/gpgpusim_trace_driven_main.cc index 5e07ace..90dc769 100644 --- a/src/trace-driven/gpgpusim_trace_driven_main.cc +++ b/src/trace-driven/gpgpusim_trace_driven_main.cc @@ -16,7 +16,7 @@ #include "../gpgpu-sim/gpu-sim.h" #include "../../libcuda/gpgpu_context.h" #include "trace_driven.h" -#include "trace_opcode.h" +#include "ISA_Def/trace_opcode.h" #include "../gpgpusim_entrypoint.h" /* TO DO: diff --git a/src/trace-driven/trace_driven.cc b/src/trace-driven/trace_driven.cc index 22c527e..76eb7ca 100644 --- a/src/trace-driven/trace_driven.cc +++ b/src/trace-driven/trace_driven.cc @@ -19,11 +19,11 @@ #include "../gpgpu-sim/gpu-sim.h" #include "../../libcuda/gpgpu_context.h" #include "trace_driven.h" -#include "trace_opcode.h" -#include "volta_opcode.h" -#include "turing_opcode.h" -#include "pascal_opcode.h" -#include "kepler_opcode.h" +#include "ISA_Def/trace_opcode.h" +#include "ISA_Def/volta_opcode.h" +#include "ISA_Def/turing_opcode.h" +#include "ISA_Def/pascal_opcode.h" +#include "ISA_Def/kepler_opcode.h" #include "../gpgpusim_entrypoint.h" diff --git a/src/trace-driven/trace_driven.h b/src/trace-driven/trace_driven.h index 2888f86..e9fecd9 100644 --- a/src/trace-driven/trace_driven.h +++ b/src/trace-driven/trace_driven.h @@ -9,7 +9,7 @@ #include "../abstract_hardware_model.h" #include "../gpgpu-sim/shader.h" -#include "trace_opcode.h" +#include "ISA_Def/trace_opcode.h" class trace_function_info: public function_info { public: diff --git a/src/trace-driven/traces-generator/NVbit_tool/Makefile b/src/trace-driven/traces-generator/NVbit_tool/Makefile deleted file mode 100644 index fc0209a..0000000 --- a/src/trace-driven/traces-generator/NVbit_tool/Makefile +++ /dev/null @@ -1,23 +0,0 @@ -NVCC=nvcc -ccbin=`which gcc` -D_FORCE_INLINES -NVBIT_PATH=../../core -INCLUDES=-I$(NVBIT_PATH) -LIBS=-L$(NVBIT_PATH) -lnvbit -NVCC_PATH=-L $(subst bin/nvcc,lib64,$(shell which nvcc | tr -s /)) -SOURCES=$(wildcard *.cu) -OBJECTS=$(SOURCES:.cu=.o) -ARCH=35 - -mkfile_path := $(abspath $(lastword $(MAKEFILE_LIST))) -current_dir := $(notdir $(patsubst %/,%,$(dir $(mkfile_path)))) - -all: $(OBJECTS) $(NVBIT_PATH)/libnvbit.a - $(NVCC) -arch=sm_$(ARCH) -O3 *.o $(LIBS) $(NVCC_PATH) -lcuda -lcudart_static -shared -o ${current_dir}.so - -%.o: %.cu - $(NVCC) -dc -c -std=c++11 $(INCLUDES) -Xptxas -cloning=no -maxrregcount=16 -Xcompiler -Wall -arch=sm_$(ARCH) -O3 -Xcompiler -fPIC $< -o $@ - -$(NVBIT_PATH)/libnvbit.a: - make -C $(NVBIT_PATH) - -clean: - rm -f *.so *.o diff --git a/src/trace-driven/traces-generator/NVbit_tool/traceall.cu b/src/trace-driven/traces-generator/NVbit_tool/traceall.cu deleted file mode 100644 index 1426528..0000000 --- a/src/trace-driven/traces-generator/NVbit_tool/traceall.cu +++ /dev/null @@ -1,646 +0,0 @@ -/* Copyright (c) 2017, NVIDIA CORPORATION. All rights reserved. - * - * Redistribution and use in source and binary forms, with or without - * modification, are permitted provided that the following conditions - * are met: - * * Redistributions of source code must retain the above copyright - * notice, this list of conditions and the following disclaimer. - * * Redistributions in binary form must reproduce the above copyright - * notice, this list of conditions and the following disclaimer in the - * documentation and/or other materials provided with the distribution. - * * Neither the name of NVIDIA CORPORATION nor the names of its - * contributors may be used to endorse or promote products derived - * from this software without specific prior written permission. - * - * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY - * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE - * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR - * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR - * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, - * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, - * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR - * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY - * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT - * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE - * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. - */ - -/* Author: Oreste Villa, [email protected] - 2018 */ - -#include <assert.h> -#include <stdint.h> -#include <stdio.h> -#include <unistd.h> -#include <string> -#include <map> -#include <inttypes.h> -#include <unistd.h> -#include <vector> -#include <iostream> -#include <bitset> -#include <sys/stat.h> -#include <sstream> -#include <algorithm> -#include <iterator> -/* every tool needs to include this once */ -#include "nvbit_tool.h" - -/* nvbit interface file */ -#include "nvbit.h" - -/* for channel */ -#include "utils/channel.hpp" - -/* for _cuda_safe and GET_VAR* macros */ -#include "macros.h" - -/* Channel used to communicate from GPU to CPU receiving thread */ -#define CHANNEL_SIZE (1l << 20) -static __managed__ ChannelDev channel_dev; -static ChannelHost channel_host; - -/* receiving thread and its control variables */ -pthread_t recv_thread; -volatile bool recv_thread_started = false; -volatile bool recv_thread_receiving = false; - -/* skip flag used to avoid re-entry on the nvbit_callback when issuing - * flush_channel kernel call */ -bool skip_flag = false; - -/* global control variables for this tool */ -uint32_t instr_begin_interval = 0; -uint32_t instr_end_interval = UINT32_MAX; -int verbose = 0; - -/* opcode to id map and reverse map */ -std::map<std::string, int> opcode_to_id_map; -std::map<int, std::string> id_to_opcode_map; - -/* kernel instruction counter, updated by the GPU */ -static __managed__ uint64_t total_dynamic_instr_counter = 0; -static __managed__ uint64_t reported_dynamic_instr_counter = 0; -static __managed__ uint64_t dynamic_instr_limit = 0; -uint64_t dynamic_instr_limit_input = 0; //0 means no limit - -#define MAX_SRC 4 -/* information collected in the instrumentation function */ -typedef struct { - int cta_id_x; - int cta_id_y; - int cta_id_z; - int warpid_tb; - int warpid_sm; - int sm_id; - int opcode_id; - uint64_t addrs[32]; - uint32_t vpc; - bool is_mem; - int32_t GPRDst; - int32_t GPRSrcs[MAX_SRC]; - int32_t numSrcs; - int32_t width; - uint32_t active_mask; - -} mem_access_t; - -/* Instrumentation function that we want to inject, please note the use of - * 1. extern "C" __device__ __noinline__ - * To prevent "dead"-code elimination by the compiler. - * 2. NVBIT_EXPORT_FUNC(dev_func) - * To notify nvbit the name of the function we want to inject. - * This name must match exactly the function name. - */ -extern "C" __device__ __noinline__ void instrument_mem(int pred, int opcode_id, int32_t vpc, - uint32_t reg_high, - uint32_t reg_low, - int32_t imm, - int32_t srcReg1, int32_t srcReg2, int32_t desReg, int32_t width) { - if (!pred) { - return; - } - - uint32_t active_mask = __ballot(1); - const int laneid = get_laneid(); - const int first_laneid = __ffs(active_mask) - 1; - - if (dynamic_instr_limit && total_dynamic_instr_counter >= dynamic_instr_limit) - if (first_laneid == laneid) { - atomicAdd((unsigned long long*)&total_dynamic_instr_counter, 1); - return; - } - - mem_access_t ma; - - /* collect memory address information */ - int64_t base_addr = (((uint64_t)reg_high) << 32) | ((uint64_t)reg_low); - uint64_t addr = base_addr + imm; - for (int i = 0; i < 32; i++) { - ma.addrs[i] = __shfl(addr, i); - } - - int4 cta = get_ctaid(); - int uniqe_threadId = threadIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x + threadIdx.x; - ma.warpid_tb = uniqe_threadId/32; - - ma.cta_id_x = cta.x; - ma.cta_id_y = cta.y; - ma.cta_id_z = cta.z; - ma.warpid_sm = get_warpid(); - ma.opcode_id = opcode_id; - ma.is_mem = true; - ma.vpc = vpc; - ma.width = width; - ma.GPRDst = desReg; - ma.GPRSrcs[0] = srcReg1; - ma.GPRSrcs[1] = srcReg2; - ma.GPRSrcs[2] = -1; - ma.GPRSrcs[3] = -1; - ma.numSrcs = 2; - ma.active_mask = active_mask; - ma.sm_id = get_smid(); - - /* first active lane pushes information on the channel */ - if (first_laneid == laneid) { - channel_dev.push(&ma, sizeof(mem_access_t)); - atomicAdd((unsigned long long*)&total_dynamic_instr_counter, 1); - atomicAdd((unsigned long long*)&reported_dynamic_instr_counter, 1); - } -} -NVBIT_EXPORT_FUNC(instrument_mem); - - -extern "C" __device__ __noinline__ void instrument_inst(int pred, int opcode_id, - uint32_t vpc, int desReg, int srcReg1, int srcReg2, int srcReg3, int srcReg4, int srcNum) { - if (!pred) { - return; - } - - int active_mask = __ballot(1); - const int laneid = get_laneid(); - const int first_laneid = __ffs(active_mask) - 1; - - if (dynamic_instr_limit && total_dynamic_instr_counter >= dynamic_instr_limit) - if (first_laneid == laneid) { - atomicAdd((unsigned long long*)&total_dynamic_instr_counter, 1); - return; - } - - - mem_access_t ma; - - int4 cta = get_ctaid(); - int uniqe_threadId = threadIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x + threadIdx.x; - ma.warpid_tb = uniqe_threadId/32; - - ma.cta_id_x = cta.x; - ma.cta_id_y = cta.y; - ma.cta_id_z = cta.z; - ma.warpid_sm = get_warpid(); - ma.opcode_id = opcode_id; - ma.is_mem = false; - ma.vpc = vpc; - - ma.GPRDst = desReg; - ma.numSrcs = srcNum; //this is the total src number including the register and others - ma.GPRSrcs[0] = srcReg1; - ma.GPRSrcs[1] = srcReg2; - ma.GPRSrcs[2] = srcReg3; - ma.GPRSrcs[3] = srcReg4; - - ma.active_mask = active_mask; - ma.sm_id = get_smid(); - - /* first active lane pushes information on the channel */ - if (first_laneid == laneid) { - channel_dev.push(&ma, sizeof(mem_access_t)); - atomicAdd((unsigned long long*)&total_dynamic_instr_counter, 1); - atomicAdd((unsigned long long*)&reported_dynamic_instr_counter, 1); - } -} - -NVBIT_EXPORT_FUNC(instrument_inst); - -void nvbit_at_init() { - setenv("CUDA_MANAGED_FORCE_DEVICE_ALLOC", "1", 1); - GET_VAR_INT( - instr_begin_interval, "INSTR_BEGIN", 0, - "Beginning of the instruction interval where to apply instrumentation"); - GET_VAR_INT( - instr_end_interval, "INSTR_END", UINT32_MAX, - "End of the instruction interval where to apply instrumentation"); - GET_VAR_LONG( - dynamic_instr_limit_input, "DYNAMIC_INSTR_LIMIT", 0, - "Limit of the number instructions to be printed, 0 means no limit"); - GET_VAR_INT(verbose, "TOOL_VERBOSE", 0, "Enable verbosity inside the tool"); - std::string pad(100, '-'); - printf("%s\n", pad.c_str()); -} - -/* instrument each memory instruction adding a call to the above instrumentation - * function */ -void nvbit_at_function_first_load(CUcontext ctx, CUfunction f) { - - dynamic_instr_limit = dynamic_instr_limit_input; - - const std::vector<Instr *> &instrs = nvbit_get_instrs(ctx, f); - if (verbose) { - printf("Inspecting function %s at address 0x%lx\n", - nvbit_get_func_name(ctx, f), nvbit_get_func_addr(f)); - } - - uint32_t cnt = 0; - /* iterate on all the static instructions in the function */ - for (auto instr : instrs) { - if (cnt < instr_begin_interval || cnt >= instr_end_interval ) { - cnt++; - continue; - } - //if (verbose) { - instr->printDecoded(); - //} - - if (opcode_to_id_map.find(instr->getOpcode()) == - opcode_to_id_map.end()) { - int opcode_id = opcode_to_id_map.size(); - opcode_to_id_map[instr->getOpcode()] = opcode_id; - id_to_opcode_map[opcode_id] = instr->getOpcode(); - } - - int opcode_id = opcode_to_id_map[instr->getOpcode()]; - - //TO DO: handle generic and TEX memory space - if(instr->isLoad() && !instr->isStore() && instr->getMemOpType() != Instr::CONSTANT) { //Mem load inst //ignore constant for now - assert(instr->getNumOperands() == 2); - - /* get the operand */ - const Instr::operand_t *dst = instr->getOperand(0); - const Instr::operand_t *src = instr->getOperand(1); - - assert(dst->type == Instr::REG); - assert(src->type == Instr::MREF); - - /* insert call to the instrumentation function with its - * arguments */ - nvbit_insert_call(instr, "instrument_mem", IPOINT_BEFORE); - nvbit_add_call_arg_pred_val(instr); - nvbit_add_call_arg_const_val32(instr, opcode_id); - nvbit_add_call_arg_const_val32(instr, (int)instr->getOffset()); - if (instr->isExtended()) { - nvbit_add_call_arg_reg_val(instr, (int)src->value[0] + 1); - } else { - nvbit_add_call_arg_reg_val(instr, (int)Instr::RZ); - } - nvbit_add_call_arg_reg_val(instr, (int)src->value[0]); - nvbit_add_call_arg_const_val32(instr, (int)src->value[1]); - nvbit_add_call_arg_const_val32(instr, (int)src->value[0]); - nvbit_add_call_arg_const_val32(instr, -1); - nvbit_add_call_arg_const_val32(instr, (int)dst->value[0]); - nvbit_add_call_arg_const_val32(instr, (int)instr->getSize()); - } - else if(instr->isStore() && !instr->isLoad() && instr->getMemOpType() != Instr::CONSTANT) { //Mem store inst //ignore constant for now - assert(instr->getNumOperands() == 2); - - /* get the operand */ - const Instr::operand_t *dst = instr->getOperand(0); - const Instr::operand_t *src = instr->getOperand(1); - - assert(dst->type == Instr::MREF); - assert(src->type == Instr::REG); - - /* insert call to the instrumentation function with its - * arguments */ - nvbit_insert_call(instr, "instrument_mem", IPOINT_BEFORE); - nvbit_add_call_arg_pred_val(instr); - nvbit_add_call_arg_const_val32(instr, opcode_id); - nvbit_add_call_arg_const_val32(instr, (int)instr->getOffset()); - if (instr->isExtended()) { - nvbit_add_call_arg_reg_val(instr, (int)dst->value[0] + 1); - } else { - nvbit_add_call_arg_reg_val(instr, (int)Instr::RZ); - } - nvbit_add_call_arg_reg_val(instr, (int)dst->value[0]); - nvbit_add_call_arg_const_val32(instr, (int)dst->value[1]); - nvbit_add_call_arg_const_val32(instr, (int)dst->value[0]); - nvbit_add_call_arg_const_val32(instr, (int)src->value[0]); - nvbit_add_call_arg_const_val32(instr, -1); - nvbit_add_call_arg_const_val32(instr, (int)instr->getSize()); - } - else if(instr->isLoad() && instr->isStore() && instr->getMemOpType() != Instr::CONSTANT) { //if it is load and store i.e. atomic inst - assert(instr->getNumOperands() == 2); - - /* get the operand */ - const Instr::operand_t *dst = instr->getOperand(0); - const Instr::operand_t *src = instr->getOperand(1); - - assert(dst->type == Instr::MREF); - assert(src->type == Instr::REG); - - /* insert call to the instrumentation function with its - * arguments */ - nvbit_insert_call(instr, "instrument_mem", IPOINT_BEFORE); - nvbit_add_call_arg_pred_val(instr); - nvbit_add_call_arg_const_val32(instr, opcode_id); - nvbit_add_call_arg_const_val32(instr, (int)instr->getOffset()); - if (instr->isExtended()) { - nvbit_add_call_arg_reg_val(instr, (int)dst->value[0] + 1); - } else { - nvbit_add_call_arg_reg_val(instr, (int)Instr::RZ); - } - nvbit_add_call_arg_reg_val(instr, (int)dst->value[0]); - nvbit_add_call_arg_const_val32(instr, (int)dst->value[1]); - nvbit_add_call_arg_const_val32(instr, (int)dst->value[0]); - nvbit_add_call_arg_const_val32(instr, (int)src->value[0]); - nvbit_add_call_arg_const_val32(instr, -1); - nvbit_add_call_arg_const_val32(instr, (int)instr->getSize()); - } - else //Other ALU, FP, DP insts - { - - nvbit_insert_call(instr, "instrument_inst", IPOINT_BEFORE); - nvbit_add_call_arg_pred_val(instr); - nvbit_add_call_arg_const_val32(instr, opcode_id); - nvbit_add_call_arg_const_val32(instr, (int)instr->getOffset()); - int srcNum = 0; - for (int i = 0; i < MAX_SRC+1; i++) { - /* get the operand "i" */ - if(i < instr->getNumOperands()) { - const Instr::operand_t *op = instr->getOperand(i); - if (op->type == Instr::REG) - nvbit_add_call_arg_const_val32(instr, (int)op->value[0]); - else - nvbit_add_call_arg_const_val32(instr, -1); - - srcNum++; - } - else - nvbit_add_call_arg_const_val32(instr, -1); - } - nvbit_add_call_arg_const_val32(instr, srcNum); - } - cnt++; - } -} - -__global__ void flush_channel() { - /* push memory access with negative cta id to communicate the kernel is - * completed */ - mem_access_t ma; - ma.cta_id_x = -1; - channel_dev.push(&ma, sizeof(mem_access_t)); - - /* flush channel */ - channel_dev.flush(); -} - -static FILE *resultsFile = NULL; -static FILE *kernelsFile= NULL; -static FILE *statsFile= NULL; -static int kernelid = 1; - -unsigned old_total_insts = 0; -unsigned old_total_reported_insts = 0; - - -void nvbit_at_cuda_event(CUcontext ctx, int is_exit, nvbit_api_cuda_t cbid, - const char *name, void *params, CUresult *pStatus) { - if (skip_flag) return; - - if (cbid == API_CUDA_cuLaunchKernel_ptsz || - cbid == API_CUDA_cuLaunchKernel) { - cuLaunchKernel_params *p = (cuLaunchKernel_params *)params; - - if (!is_exit) { - - - if (mkdir("traces", S_IRWXU | S_IRWXG | S_IROTH | S_IXOTH) == -1) { - if( errno == EEXIST ) { - // alredy exists - } else { - // something else - std::cout << "cannot create folder error:" << strerror(errno) << std::endl; - return; - } - } - - int nregs; - _cuda_safe( - cuFuncGetAttribute(&nregs, CU_FUNC_ATTRIBUTE_NUM_REGS, p->f)); - - int shmem_static_nbytes; - _cuda_safe(cuFuncGetAttribute(&shmem_static_nbytes, - CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES, - p->f)); - - - - std::string func_name(nvbit_get_func_name(ctx, p->f)); - std::string::size_type end_pos = func_name.find('('); - if (end_pos != std::string::npos) - { - // std::string::size_type pos = func_name.find('<'); - //if (pos != std::string::npos) - // end_pos = pos; - - //std::string::size_type start_pos = func_name.find(' '); - //if (start_pos == std::string::npos) - // start_pos = 0; - //else - // start_pos++; - - func_name = func_name.substr(0, end_pos); - } - - char buffer[1024]; - sprintf (buffer, "./traces/%d-%s.trace", kernelid, func_name.c_str()); - - resultsFile = fopen(buffer, "w"); - - printf("Writing results to %s\n", buffer); - - fprintf(resultsFile, "-kernel name = %s", nvbit_get_func_name(ctx, p->f)); - fprintf(resultsFile, "\n"); - fprintf(resultsFile, "-kernel id = %d", kernelid); - fprintf(resultsFile, "\n"); - fprintf(resultsFile, "-grid dim = (%d,%d,%d)", p->gridDimX, p->gridDimY, p->gridDimZ); - fprintf(resultsFile, "\n"); - fprintf(resultsFile, "-block dim = (%d,%d,%d)", p->blockDimX, p->blockDimY, p->blockDimZ); - fprintf(resultsFile, "\n"); - fprintf(resultsFile, "-shmem = %d", shmem_static_nbytes + p->sharedMemBytes); - fprintf(resultsFile, "\n"); - fprintf(resultsFile, "-nregs = %d", nregs); - fprintf(resultsFile, "\n"); - fprintf(resultsFile, "-cuda stream id = %d", (uint64_t)p->hStream); - fprintf(resultsFile, "\n\n"); - - fprintf(resultsFile, "#traces format = threadblock_x threadblock_y threadblock_z warpid_tb sm_id warpid_sm PC mask dest_num reg_dests opcode src_num reg_srcs mem_width mem_addresses"); - fprintf(resultsFile, "\n"); - - if (kernelid == 1) { - kernelsFile = fopen("./traces/kernelslist", "w"); - statsFile = fopen("./traces/stats.csv", "w"); - fprintf(statsFile, "kernel name,total_insts,total_reported_insts\n"); - } - else { - kernelsFile = fopen("./traces/kernelslist", "a"); - statsFile = fopen("./traces/stats.csv", "a"); - } - - sprintf (buffer, "%d-%s.trace", kernelid, func_name.c_str()); - fprintf(kernelsFile, buffer); - fprintf(kernelsFile, "\n"); - fclose(kernelsFile); - - fprintf(statsFile, buffer); - fprintf(statsFile, ","); - - kernelid++; - recv_thread_receiving = true; - - } else { - /* make sure current kernel is completed */ - cudaDeviceSynchronize(); - assert(cudaGetLastError() == cudaSuccess); - - /* make sure we prevent re-entry on the nvbit_callback when issuing - * the flush_channel kernel */ - skip_flag = true; - - /* issue flush of channel so we are sure all the memory accesses - * have been pushed */ - flush_channel<<<1, 1>>>(); - cudaDeviceSynchronize(); - assert(cudaGetLastError() == cudaSuccess); - - /* unset the skip flag */ - skip_flag = false; - - /* wait here until the receiving thread has not finished with the - * current kernel */ - while (recv_thread_receiving) { - pthread_yield(); - } - - unsigned total_insts_per_kernel = total_dynamic_instr_counter - old_total_insts; - old_total_insts = total_dynamic_instr_counter; - - unsigned reported_insts_per_kernel = reported_dynamic_instr_counter - old_total_reported_insts; - old_total_reported_insts = reported_dynamic_instr_counter; - - fprintf(statsFile, ""); - fprintf(statsFile, "%d,%d",total_insts_per_kernel,reported_insts_per_kernel); - fprintf(statsFile, "\n"); - - - fclose(resultsFile); - fclose(statsFile); - } - } -} - -bool is_number(const std::string& s) -{ - std::string::const_iterator it = s.begin(); - while (it != s.end() && std::isdigit(*it)) ++it; - return !s.empty() && it == s.end(); -} - -void *recv_thread_fun(void *) { - char *recv_buffer = (char *)malloc(CHANNEL_SIZE); - - while (recv_thread_started) { - uint32_t num_recv_bytes = 0; - if (recv_thread_receiving && - (num_recv_bytes = channel_host.recv(recv_buffer, CHANNEL_SIZE)) > - 0) { - uint32_t num_processed_bytes = 0; - while (num_processed_bytes < num_recv_bytes) { - mem_access_t *ma = - (mem_access_t *)&recv_buffer[num_processed_bytes]; - - /* when we get this cta_id_x it means the kernel has completed - */ - if (ma->cta_id_x == -1) { - recv_thread_receiving = false; - break; - } - - fprintf(resultsFile, "%d ", ma->cta_id_x); - fprintf(resultsFile, "%d ", ma->cta_id_y); - fprintf(resultsFile, "%d ", ma->cta_id_z); - fprintf(resultsFile, "%d ", ma->warpid_tb); - fprintf(resultsFile, "%d ", ma->sm_id); - fprintf(resultsFile, "%d ", ma->warpid_sm); - fprintf(resultsFile, "0x%016lx ", ma->vpc); // Print the virtual PC. - fprintf(resultsFile, "%-8.8" PRIx32 " ", ma->active_mask); - if(ma->GPRDst >= 0) { - fprintf(resultsFile, "1 "); - fprintf(resultsFile, "R%d ", ma->GPRDst); - } - else - fprintf(resultsFile, "0 "); - - // Print the opcode. - fprintf(resultsFile, "%s ", id_to_opcode_map[ma->opcode_id].c_str()); - unsigned src_count=0; - for (int s = 0; s < MAX_SRC; s++) // GPR srcs count. - if(ma->GPRSrcs[s] >= 0) src_count++; - fprintf(resultsFile, "%d ", src_count); - - for (int s = 0; s < MAX_SRC; s++) // GPR srcs. - if(ma->GPRSrcs[s] >= 0) fprintf(resultsFile, "R%d ", ma->GPRSrcs[s]); - - //print addresses - std::bitset<32> mask(ma->active_mask); - if(ma->is_mem) { - //fprintf(resultsFile, "%d ", ma->width); - std::istringstream iss(id_to_opcode_map[ma->opcode_id]); - std::vector<std::string> tokens; - std::string token; - while (std::getline(iss, token, '.')) { - if (!token.empty()) - tokens.push_back(token); - } - if (tokens.size()>=3){ - if (is_number(tokens[2])){ - fprintf(resultsFile, "%d ", (std::stoi(tokens[2],nullptr)/8)); - } - else{ - fprintf(resultsFile, "%d ", 4); - } - } - else{ - fprintf(resultsFile, "%d ", 4); - } - - for (int s = 0; s < 32; s++) - if(mask.test(s)) - fprintf(resultsFile, "0x%016lx ", ma->addrs[s]); - } - else - { - fprintf(resultsFile, "0 "); - } - - fprintf(resultsFile, "\n"); - - num_processed_bytes += sizeof(mem_access_t); - } - } - } - free(recv_buffer); - return NULL; -} - -void nvbit_at_ctx_init(CUcontext ctx) { - recv_thread_started = true; - channel_host.init(0, CHANNEL_SIZE, &channel_dev, NULL); - pthread_create(&recv_thread, NULL, recv_thread_fun, NULL); -} - -void nvbit_at_ctx_term(CUcontext ctx) { - if (recv_thread_started) { - recv_thread_started = false; - pthread_join(recv_thread, NULL); - } -} diff --git a/src/trace-driven/traces-generator/README b/src/trace-driven/traces-generator/README deleted file mode 100644 index d79d075..0000000 --- a/src/trace-driven/traces-generator/README +++ /dev/null @@ -1 +0,0 @@ -TO DO diff --git a/src/trace-driven/traces-generator/traces-post-processing/Makefile b/src/trace-driven/traces-generator/traces-post-processing/Makefile deleted file mode 100755 index 63c3f38..0000000 --- a/src/trace-driven/traces-generator/traces-post-processing/Makefile +++ /dev/null @@ -1,11 +0,0 @@ -TARGET := post-traces-processing - -$(TARGET): post-traces-processing.cpp - g++ -o $@ $^ - -run: $(TARGET) - ./$(TARGET) - -clean: - rm -f $(TARGET) *.o - diff --git a/src/trace-driven/traces-generator/traces-post-processing/post-traces-processing.cpp b/src/trace-driven/traces-generator/traces-post-processing/post-traces-processing.cpp deleted file mode 100644 index 0dbb4e1..0000000 --- a/src/trace-driven/traces-generator/traces-post-processing/post-traces-processing.cpp +++ /dev/null @@ -1,195 +0,0 @@ -#include <vector>
-#include <iostream>
-#include <fstream>
-#include <string>
-#include <sstream>
-#include <stdio.h>
-#include <math.h>
-using namespace std;
-
-struct threadblock_info
-{
- bool initialized;
- unsigned tb_id_x, tb_id_y, tb_id_z;
- vector< vector< string > > warp_insts_array;
- threadblock_info() {
- initialized = false;
- tb_id_x = tb_id_y = tb_id_z = 0;
- }
-};
-
-void group_per_block(const char* filepath);
-void group_per_core(const char* filepath);
-
-//////////////////////////////////////////////////////////////////////////////////////////////////////////////////
-int main(int argc, char** argv)
-{
-
- string kernellist_filepath;
- bool is_per_core;
- if(argc == 1)
- {
- cout << "File path is missing\n";
- return 0;
- } else if(argc == 2)
- {
- kernellist_filepath = argv[1];
- is_per_core = true;
-
- } else if(argc == 3) {
- kernellist_filepath = argv[1];
- is_per_core = bool(argv[2]);
- }
- else {
- cout << "Too Many Arguemnts!\n";
- return 0;
- }
-
- ifstream ifs;
- ofstream ofs;
-
- ifs.open(kernellist_filepath.c_str());
- ofs.open((string(kernellist_filepath) + ".g").c_str());
-
- if (!ifs.is_open()) {
- cout << "Unable to open file: " <<kernellist_filepath<<endl;
- return 0;
- }
-
-
- string directory(kernellist_filepath);
- const size_t last_slash_idx = directory.rfind('/');
- if (std::string::npos != last_slash_idx)
- {
- directory = directory.substr(0, last_slash_idx);
- }
-
- string line;
- string filepath;
- while(!ifs.eof()) {
- getline(ifs, line);
- if(line.empty())
- continue;
- filepath = directory+"/"+line;
- group_per_block(filepath.c_str());
- ofs<<line + "g"<<endl;
- }
-
- ifs.close();
- ofs.close();
- return 0;
-}
-
-
-void group_per_block(const char* filepath) {
-
- ofstream ofs;
- ifstream ifs;
-
- ifs.open(filepath);
-
- if (!ifs.is_open()) {
- cout << "Unable to open file: " <<filepath<<endl;
- return;
- }
-
- cout << "Processing file " <<filepath<<endl;
- ofs.open((string(filepath) + "g").c_str());
-
- vector<threadblock_info> insts;
- unsigned grid_dim_x, grid_dim_y, grid_dim_z, tb_dim_x, tb_dim_y, tb_dim_z;
- unsigned tb_id_x, tb_id_y, tb_id_z, tb_id, warpid_tb;
- string line;
- stringstream ss;
- string string1, string2;
- bool found_grid_dim = false, found_block_dim = false;
-
- while(!ifs.eof()) {
- getline(ifs, line);
-
- if (line.length() == 0 || line[0] == '#') {
- ofs<<line<<endl;
- continue;
- }
-
- else if(line[0] == '-') {
- ss.str(line);
- ss.ignore();
- ss>>string1>>string2;
- if(string1 == "grid" && string2 == "dim") {
- sscanf(line.c_str(), "-grid dim = (%d,%d,%d)", &grid_dim_x, &grid_dim_y, &grid_dim_z);
- found_grid_dim = true;
- }
- else if (string1 == "block" && string2 == "dim") {
- sscanf(line.c_str(), "-block dim = (%d,%d,%d)", &tb_dim_x, &tb_dim_y, &tb_dim_z);
- found_block_dim = true;
- }
-
- if(found_grid_dim && found_block_dim) {
- insts.resize(grid_dim_x*grid_dim_y*grid_dim_z);
- for(unsigned i = 0; i<insts.size(); ++i) {
- insts[i].warp_insts_array.resize(ceil(float(tb_dim_x*tb_dim_y*tb_dim_z)/32));
- }
- }
- ofs<<line<<endl;
- continue;
- }
- else {
-
- ss.str(line);
- ss>>tb_id_x>>tb_id_y>>tb_id_z>>warpid_tb;
- tb_id = tb_id_z * grid_dim_y * grid_dim_x + tb_id_y * grid_dim_x + tb_id_x;
- if(!insts[tb_id].initialized) {
- insts[tb_id].tb_id_x = tb_id_x;
- insts[tb_id].tb_id_y = tb_id_y;
- insts[tb_id].tb_id_z = tb_id_z;
- insts[tb_id].initialized = true;
- }
- insts[tb_id].warp_insts_array[warpid_tb].push_back(line);
- }
-
- }
-
-
- for(unsigned i=0; i<insts.size(); ++i) {
- //ofs<<string<<endl;
- if(insts[i].initialized && insts[i].warp_insts_array.size() > 0) {
- ofs<<endl<<"#BEGIN_TB"<<endl;
- ofs<<endl<<"thread block = "<<insts[i].tb_id_x<<","<<insts[i].tb_id_y<<","<<insts[i].tb_id_z<<endl;
- }
- else {
- cout<<"Warning: Thread block "<<insts[i].tb_id_x<<","<<insts[i].tb_id_y<<","<<insts[i].tb_id_z<<" is empty"<<endl;
- continue;
- //ofs.close();
- //return;
- }
- for(unsigned j=0; j<insts[i].warp_insts_array.size(); ++j) {
- ofs<<endl<<"warp = "<<j<<endl;
- ofs<<"insts = "<<insts[i].warp_insts_array[j].size()<<endl;
- if(insts[i].warp_insts_array[j].size() == 0) {
- cout<<"Warning: Warp "<<j<<" in thread block"<<insts[i].tb_id_x<<","<<insts[i].tb_id_y<<","<<insts[i].tb_id_z<<" is empty"<<endl;
- // ofs.close();
- // return;
- }
- for(unsigned k=0; k<insts[i].warp_insts_array[j].size(); ++k) {
- ofs<<insts[i].warp_insts_array[j][k]<<endl;
- }
- }
- ofs<<endl<<"#END_TB"<<endl;
- }
-
-
-
- ofs.close();
- ifs.close();
-
-}
-
-
-void group_per_core(const char* filepath)
-{
-
- //TO DO
-
-
-}
|