Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

How to read out predicter value? #36

Open
leiwen83 opened this issue Jan 20, 2021 · 8 comments
Open

How to read out predicter value? #36

leiwen83 opened this issue Jan 20, 2021 · 8 comments

Comments

@leiwen83
Copy link

Hi,

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

Thx,
Lei

@x-y-z
Copy link
Collaborator

x-y-z commented Jan 20, 2021

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
Copy link

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

@leiwen83
Copy link
Author

leiwen83 commented Jan 20, 2021

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
Copy link
Collaborator

x-y-z commented Jan 20, 2021

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

@leiwen83
Copy link
Author

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

@x-y-z
Copy link
Collaborator

x-y-z commented Jan 20, 2021

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

@WilliamWangPeng
Copy link

WilliamWangPeng commented Jul 28, 2021

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
Copy link
Collaborator

x-y-z commented Jul 28, 2021

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

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

No branches or pull requests

4 participants