diff options
Diffstat (limited to 'src/trace-driven/traces-generator/NVbit_tool')
| -rw-r--r-- | src/trace-driven/traces-generator/NVbit_tool/Makefile | 23 | ||||
| -rw-r--r-- | src/trace-driven/traces-generator/NVbit_tool/traceall.cu | 646 |
2 files changed, 669 insertions, 0 deletions
diff --git a/src/trace-driven/traces-generator/NVbit_tool/Makefile b/src/trace-driven/traces-generator/NVbit_tool/Makefile new file mode 100644 index 0000000..fc0209a --- /dev/null +++ b/src/trace-driven/traces-generator/NVbit_tool/Makefile @@ -0,0 +1,23 @@ +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 new file mode 100644 index 0000000..1426528 --- /dev/null +++ b/src/trace-driven/traces-generator/NVbit_tool/traceall.cu @@ -0,0 +1,646 @@ +/* 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); + } +} |