NVlabs / NVBit

199 stars 18 forks source link

How to read out predicter value? #36

Open leiwen83 opened 3 years ago

leiwen83 commented 3 years ago

Hi,

If I want to read a predicter value in the inject code, How could I did the coding?

Thx, Lei

x-y-z commented 3 years ago

You can use nvbit_read_pred_reg() to read predicate registers. Also you need to include nvbit_reg_rw.h file in your injection function file.

liujian883 commented 3 years ago

how to index the pred reg with this api ?the general reg rw api can be set with reg index

leiwen83 commented 3 years ago
extern "C" __device__ __noinline__ int32_t nvbit_read_pred_reg() {
#pragma unroll
    for (int i = 0; i < 32; i++) __nvbit_var += i;
    return __nvbit_var;
}

The definition of this api is as above. how to parse the result? Like if I want to read what the current status for P5, whether it is true or not, how should I deal wtih the return value?

x-y-z commented 3 years ago

It reads out the entire predicate register and only the lower 8 bits ([P7 .. P0]) are used.

leiwen83 commented 3 years ago

I try to add the code, but read out as 0... I test it over 2080ti.

x-y-z commented 3 years ago

If the predicate register is not used in the code or is set to 0, the result will be 0.

WilliamWangPeng commented 2 years ago

I try to add the code, but read out as 0... I test it over 2080ti.

hi dear @leiwen83
I'm also interested in this nvbit_read_pred_reg() function. but how can you display the values of this function, I try to print the values, but there is an error here, image

image

thank you
best regards William

x-y-z commented 2 years ago

It is a device function and should be called only in your instrumentation functions in the inject_func.cu file.