Skip to content

Commit

Permalink
format code
Browse files Browse the repository at this point in the history
  • Loading branch information
@barnes88
barnes88 committed Jul 11, 2024
1 parent b45bb82 commit 219a76d
Show file tree
Hide file tree
Showing 2 changed files with 31 additions and 23 deletions.
16 changes: 8 additions & 8 deletions util/tracer_nvbit/tracer_tool/inject_funcs.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -18,14 +18,14 @@
* extern "C" __device__ __noinline__
* To prevent "dead"-code elimination by the compiler.
*/
extern "C" __device__ __noinline__ void instrument_inst(
int pred, int opcode_id, int32_t vpc, bool is_mem, uint64_t addr,
int32_t width, int32_t desReg, int32_t srcReg1, int32_t srcReg2,
int32_t srcReg3, int32_t srcReg4, int32_t srcReg5, int32_t srcNum,
uint64_t immediate, uint64_t pchannel_dev,
uint64_t ptotal_dynamic_instr_counter,
uint64_t preported_dynamic_instr_counter, uint64_t pstop_report,
uint32_t line_num) {
extern "C" __device__ __noinline__ void
instrument_inst(int pred, int opcode_id, int32_t vpc, bool is_mem,
uint64_t addr, int32_t width, int32_t desReg, int32_t srcReg1,
int32_t srcReg2, int32_t srcReg3, int32_t srcReg4,
int32_t srcReg5, int32_t srcNum, uint64_t immediate,
uint64_t pchannel_dev, uint64_t ptotal_dynamic_instr_counter,
uint64_t preported_dynamic_instr_counter, uint64_t pstop_report,
uint32_t line_num) {
const int active_mask = __ballot_sync(__activemask(), 1);
const int predicate_mask = __ballot_sync(__activemask(), pred);
const int laneid = get_laneid();
Expand Down
38 changes: 23 additions & 15 deletions util/tracer_nvbit/tracer_tool/tracer_tool.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -75,8 +75,8 @@ std::string stats_location = cwd + "/traces/stats.csv";

/* kernel instruction counter, updated by the GPU */
uint64_t dynamic_kernel_limit_start =
0; // 0 means start from the begging kernel
uint64_t dynamic_kernel_limit_end = 0; // 0 means no limit
0; // 0 means start from the begging kernel
uint64_t dynamic_kernel_limit_end = 0; // 0 means no limit

enum address_format { list_all = 0, base_stride = 1, base_delta = 2 };

Expand Down Expand Up @@ -231,7 +231,7 @@ void instrument_function_if_needed(CUcontext ctx, CUfunction func) {
if (mem_oper_idx >= 0) {
nvbit_add_call_arg_const_val32(instr, 1);
assert(num_mref <= 2);
if (num_mref == 2) { // LDGSTS
if (num_mref == 2) { // LDGSTS
nvbit_add_call_arg_mref_addr64(instr, 1 - mem_oper_idx);
} else {
nvbit_add_call_arg_mref_addr64(instr, mem_oper_idx);
Expand Down Expand Up @@ -296,7 +296,8 @@ 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 (skip_flag)
return;

if (first_call == true) {
first_call = false;
Expand All @@ -316,7 +317,8 @@ void nvbit_at_cuda_event(CUcontext ctx, int is_exit, nvbit_api_cuda_t cbid,
dynamic_kernel_limit_start == 1)
active_region = true;
else {
if (active_from_start) active_region = false;
if (active_from_start)
active_region = false;
}

if (user_defined_folders == 1) {
Expand Down Expand Up @@ -527,7 +529,8 @@ void nvbit_at_cuda_event(CUcontext ctx, int is_exit, nvbit_api_cuda_t cbid,

bool is_number(const std::string &s) {
std::string::const_iterator it = s.begin();
while (it != s.end() && std::isdigit(*it)) ++it;
while (it != s.end() && std::isdigit(*it))
++it;
return !s.empty() && it == s.end();
}

Expand All @@ -543,13 +546,14 @@ unsigned get_datawidth_from_opcode(const std::vector<std::string> &opcode) {
}
}

return 4; // default is 4 bytes
return 4; // default is 4 bytes
}

bool check_opcode_contain(const std::vector<std::string> &opcode,
std::string param) {
for (unsigned i = 0; i < opcode.size(); ++i)
if (opcode[i] == param) return true;
if (opcode[i] == param)
return true;

return false;
}
Expand Down Expand Up @@ -637,7 +641,7 @@ void *recv_thread_fun(void *) {
if (lineinfo) {
fprintf(resultsFile, "%d ", ma->line_num);
}
fprintf(resultsFile, "%04x ", ma->vpc); // Print the virtual PC
fprintf(resultsFile, "%04x ", ma->vpc); // Print the virtual PC
fprintf(resultsFile, "%08x ", ma->active_mask & ma->predicate_mask);
if (ma->GPRDst >= 0) {
fprintf(resultsFile, "1 ");
Expand All @@ -648,12 +652,14 @@ void *recv_thread_fun(void *) {
// 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++;
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]);
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 & ma->predicate_mask);
Expand All @@ -662,7 +668,8 @@ void *recv_thread_fun(void *) {
std::vector<std::string> tokens;
std::string token;
while (std::getline(iss, token, '.')) {
if (!token.empty()) tokens.push_back(token);
if (!token.empty())
tokens.push_back(token);
}
fprintf(resultsFile, "%d ", get_datawidth_from_opcode(tokens));

Expand Down Expand Up @@ -696,7 +703,8 @@ void *recv_thread_fun(void *) {
// list all the addresses
fprintf(resultsFile, "%u ", address_format::list_all);
for (int s = 0; s < 32; s++) {
if (mask.test(s)) fprintf(resultsFile, "0x%016lx ", ma->addrs[s]);
if (mask.test(s))
fprintf(resultsFile, "0x%016lx ", ma->addrs[s]);
}
}
} else {
Expand Down

0 comments on commit 219a76d

Please sign in to comment.