Comments (5)
Ohh, I just found this limitation in the paper.
Injected functions may not use shared and constant memory because that memory can be used by the application itself.
But in my inject_kernel, I just want to read the shared memory of the origin kernel instead of modifying it. That is also not possible?
from nvbit.
We currently don't have an interface for reading shared memory within an injected function.
We will keep it in mind for next version, but unfortunately we will not be able to explore this in short time as it requires non trivial amount of plumbing.
from nvbit.
Can I pass the c[0x0][0x18]
and c[0x0][0x1c]
as the base address of the shared memory? (BTW, I'm using GV100).
It works for me in this simple test:
extern "C" __device__ __noinline__ void inject_kernel(unsigned int r2, unsigned int r3) {
unsigned long x = ((unsigned long)r3 << 32) | r2;
long *shared_mem_ptr = (long *)x;
printf("%ld\n", *shared_mem_ptr);
}
from nvbit.
We never tried that explicitly, but it seems promising.
From the earlier post I thought you were looking for an API to dump the entire content of shmem (without knowing a particular address), but if you know the exact address I guess you can do the way you are doing above... it should be safe.
from nvbit.
Thanks for your explanation.
from nvbit.
Related Issues (20)
- TensorFlow stuck in infinite loop when executed with NVBit.
- How to close the WARNING in NVBit
- control flow graph can not be extended
- F2FP or Uf2FP?
- RNNT error, operation not permitted when stream is capturing HOT 2
- Texture Memory Reference HOT 1
- How to flush instruction cache? HOT 1
- What is the maximum number of bits/bytes allowed after nvbit_insert_call?
- Loading NVBit tool with dlopen() explicitly HOT 2
- Finding source line number (in the host code) for cudaMalloc and similar API calls
- Getting PC from NVBit HOT 2
- NVBit hangs when creating Cuda Contexts in parallel (multi-gpu)
- Nvbit misses kernels compared to Nsight products
- Using Nvbit for a specific region of code
- How to use APIs in our own application code
- CUDA 12.0 / Driver > 510 - Unsupported? HOT 1
- NVbit can not work with Lammps? HOT 1
- mem_trace address type
- Non-deterministic behavior
- CUDA 12 and more recent nvidia driver support HOT 2
Recommend Projects
-
React
A declarative, efficient, and flexible JavaScript library for building user interfaces.
-
Vue.js
🖖 Vue.js is a progressive, incrementally-adoptable JavaScript framework for building UI on the web.
-
Typescript
TypeScript is a superset of JavaScript that compiles to clean JavaScript output.
-
TensorFlow
An Open Source Machine Learning Framework for Everyone
-
Django
The Web framework for perfectionists with deadlines.
-
Laravel
A PHP framework for web artisans
-
D3
Bring data to life with SVG, Canvas and HTML. 📊📈🎉
-
Recommend Topics
-
javascript
JavaScript (JS) is a lightweight interpreted programming language with first-class functions.
-
web
Some thing interesting about web. New door for the world.
-
server
A server is a program made to process requests and deliver data to clients.
-
Machine learning
Machine learning is a way of modeling and interpreting data that allows a piece of software to respond intelligently.
-
Visualization
Some thing interesting about visualization, use data art
-
Game
Some thing interesting about game, make everyone happy.
Recommend Org
-
Facebook
We are working to build community through open source technology. NB: members must have two-factor auth.
-
Microsoft
Open source projects and samples from Microsoft.
-
Google
Google ❤️ Open Source for everyone.
-
Alibaba
Alibaba Open Source for everyone
-
D3
Data-Driven Documents codes.
-
Tencent
China tencent open source team.
from nvbit.