llvm / llvm/Polygeist

Question: Support for differentiating between volatile and non-volatile memory operations

Open
#418 1 comment 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

Dominant language
C++
Stars
624
Forks
170
PR merge metrics
No merged PRs in 30d

Description

Hi,

After lowering a CUDA program to mlir format using cgeist, I wanted to run some analysis on memory operations. I want to differentiate memory operations which are 'volatile.'

```
__global__ void test(volatile int* a, int* b) {
a[some_idx] = data();
b[another_idx] = data();
}
```

For example, I want to differentiate stores to array ```a``` and ```b``` as volatile has special semantics in CUDA. In the LLVM IR, it is possible (```isVolatile``` method). Can something similar be achieved here?

Contributor guide

No contributing guide indexed for this repository

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

Research direction

The issue does not name any source file, test, or entry point. First clarify whether support for exposing volatile memory operations is requested, then identify the relevant lowering path and define a test that distinguishes the two stores.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
compilers
Issue type
Feature
Difficulty
4/5
Estimated time
3-5 days
Activity status
Stale
Clarity
Needs clarification
Newbie friendliness
25/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.