[OPENMP][NVPTX]Fix barriers and parallel level counters, NFC.

[ABataev]

Summary

Parallel level counter should be volatile to prevent some dangerous
optimiations by the ptxas. Otherwise, ptxas optimizations lead to
undefined behaviour in some cases.
Also, use __threadfence() for #pragma omp flush and if the barrier
should not be used (we have only one thread in the team), still perform
flush operation since the standard requires implicit flush when
executing barriers.

[grokos]

LGTM

[jdoerfert]

I have a problem with this patch and this kind of patches/reviews:

1. How can a "Fix" commit be "NFC" at the same time?
2. The commit message basically mentions that multiple unrelated changes are combined in one patch.
3. I somehow doubt that it's the "ptx optimizations" that "lead to UB":

Parallel level counter should be volatile to prevent some dangerous optimiations by the ptxas. Otherwise, ptxas optimizations lead to undefined behaviour in some cases.

This does not sound like a "solution" but more like a hack. With volatile the problem goes away so let's make it volatile.

4. The more patches go through with this kind of "review" the more "suspicious" this looks to me.

[arsenm]

No tests

[ABataev]

In D62199#1512617, @jdoerfert wrote:

I have a problem with this patch and this kind of patches/reviews:

1. How can a "Fix" commit be "NFC" at the same time?
2. The commit message basically mentions that multiple unrelated changes are combined in one patch.
3. I somehow doubt that it's the "ptx optimizations" that "lead to UB":

Parallel level counter should be volatile to prevent some dangerous optimiations by the ptxas. Otherwise, ptxas optimizations lead to undefined behaviour in some cases.

This does not sound like a "solution" but more like a hack. With volatile the problem goes away so let's make it volatile.

4. The more patches go through with this kind of "review" the more "suspicious" this looks to me.

1. Because it is impossible to test them, that why it is NFC.
2. Yes, this is true. Bu the changes are rather small.
3. It is only your opinion.
4. Again, this is your problem, The patch was approved bythe code owner

In D62199#1512633, @arsenm wrote:

No tests

These changes cannot be tested. The problem with the original code appear only in very rare situations in highly optimized code.

1. parallelLevel is accessed by many threads simultaneously and without volatile can be optimized and lead to the undefined behavior.
2. I have no idea how to test the cache flushing.

That's why this patch is marked as NFC.

[arsenm]

In D62199#1512638, @ABataev wrote:

1. Because it is impossible to test them, that why it is NFC.

This doesn't make it NFC. There is intended change in behavior.

2. Yes, this is true. Bu the changes are rather small.

This isn't a reason to merge unrelated patches

3. It is only your opinion.

I agree with this issue. This is very hand wavy. If you cannot describe what the undefined behavior or problem is is, you can't come up with a sensible solution. As presented here, this does not justify why this fixes or works around the issue.

4. Again, this is your problem, The patch was approved bythe code owner

Patches can still be subject to post-commit reviews

[ABataev]

In D62199#1512697, @arsenm wrote:

In D62199#1512638, @ABataev wrote:

1. Because it is impossible to test them, that why it is NFC.

This doesn't make it NFC. There is intended change in behavior.

In beehavior of the compiler, but not the code itself.

2. Yes, this is true. Bu the changes are rather small.

This isn't a reason to merge unrelated patches

Yes, I agree , it would be better to split this patch into 2 parts.

3. It is only your opinion.

I agree with this issue. This is very hand wavy. If you cannot describe what the undefined behavior or problem is is, you can't come up with a sensible solution. As presented here, this does not justify why this fixes or works around the issue.

It is related somehow with the ptxas translator. I cannot explain why it causes the problem, but in some rare cases, when you mix simultaneos access and the atomic operations, the code is optimized in the unknown way and it leads to undefined behavior.

4. Again, this is your problem, The patch was approved bythe code owner

Patches can still be subject to post-commit reviews

Sure.

[Hahnfeld]

In D62199#1512697, @arsenm wrote:

In D62199#1512638, @ABataev wrote:

1. Because it is impossible to test them, that why it is NFC.

This doesn't make it NFC. There is intended change in behavior.

+1. NFC = no functional change, and apparently there are changes and they are noticeable.

2. Yes, this is true. Bu the changes are rather small.

This isn't a reason to merge unrelated patches

+1

3. It is only your opinion.

I agree with this issue. This is very hand wavy. If you cannot describe what the undefined behavior or problem is is, you can't come up with a sensible solution. As presented here, this does not justify why this fixes or works around the issue.

4. Again, this is your problem, The patch was approved bythe code owner

Patches can still be subject to post-commit reviews

Multiple things that bother me:

1. I'm not sure that @grokos is indeed the code owner. There was a proposal, but it was never formally accepted. Given that he's now working for a different vendor, does it make sense for him to be owner for code related to Nvidia code?
2. According to the Developer Policy:

Note that code ownership is completely different than reviewers: anyone can review a piece of code, and we welcome code review from anyone who is interested. Code owners are the “last line of defense” to guarantee that all patches that are committed are actually reviewed.

So saying "this is your problem" to a valid review is not appropriate IMO.

[hfinkel]

In D62199#1513126, @Hahnfeld wrote:

In D62199#1512697, @arsenm wrote:

In D62199#1512638, @ABataev wrote:

1. Because it is impossible to test them, that why it is NFC.

This doesn't make it NFC. There is intended change in behavior.

+1. NFC = no functional change, and apparently there are changes and they are noticeable.

2. Yes, this is true. Bu the changes are rather small.

This isn't a reason to merge unrelated patches

+1

3. It is only your opinion.

I agree with this issue. This is very hand wavy. If you cannot describe what the undefined behavior or problem is is, you can't come up with a sensible solution. As presented here, this does not justify why this fixes or works around the issue.

4. Again, this is your problem, The patch was approved bythe code owner

Patches can still be subject to post-commit reviews

Multiple things that bother me:

1. I'm not sure that @grokos is indeed the code owner. There was a proposal, but it was never formally accepted. Given that he's now working for a different vendor, does it make sense for him to be owner for code related to Nvidia code?

There need not be a relationship between one's employer and code ownership, even for backends.

2. According to the Developer Policy:

Note that code ownership is completely different than reviewers: anyone can review a piece of code, and we welcome code review from anyone who is interested. Code owners are the “last line of defense” to guarantee that all patches that are committed are actually reviewed.

So saying "this is your problem" to a valid review is not appropriate IMO.

That's correct: the response is inappropriate.

Code owners serve as reviewers of last resort. The entire project is run using a consensus-driven review process, and all interested parties are included. A code owner might take a leadership role in forming that consensus, but all code is subject to consensus-driven review both pre- and post-commit. Consensus does not necessarily mean unanimous agreement, but in this case there is clearly broad disagreement among the reviewers regarding the combining of small changes, the title of the commit and level of explanation regarding the lack of tests, and so on. These concerns should be addressed.

[arsenm]

I think this should be reverted. The well understood parts of the commit should absolutely be separated. There also needs to be a comment explaining the reason for volatile. For posterity, that should really be a separate commit in the history.

[grokos]

Let me address what has been written so far from my perspective and request some changes since there has been demand for them:

1. No doubt there are 2 changes here, unrelated to one another. Yes, they could be split into 2 patches to strictly comply with community guidelines. And I would prefer it if the author had submitted two distinct patches. On the other hand, I don't like missing the forest for the trees. This is a 4-line patch, I figured out it is OK to let both changes in at once. Sure, the small size of a patch is not a free pass to doing whatever we want, but in this case committing both changes in one go doesn't do any harm, at least none that I can think of so that I would make a big deal out of it. There's the letter of the law and the spirit of the law, in cases where these two collide I tend to go with the latter.
2. My bad that I didn't request a test for cache flushing, I was predisposed due to the NFC tag. I agree there is no way to test whether the cache has been flushed, but @ABataev you could provide a test case which fails without cache flushing so that you can demonstrate that this patch does indeed work. There must be such a test case, after all how did you figure out there is a problem there?
3. The NFC bit does not apply to the barrier fix. If we keep both changes in one patch, that should be reflected in the title. I agree that not being able to test a change does not make it NFC. You can call it NTC (non-testable change), but not NFC.
4. In case of volatile, it seems there's a gray line here between what constitutes a NFC patch. The way I see it, "volatile" prevents compiler optimizations, it doesn't change anything in the way this code works, the algorithms/data structures used etc., the only change is how the code is compiled. On the other hand, we could argue that since the compiled library behaves differently there is indeed a functional change. Once again, to me this is a gray line and there is no absolute answer. However, since here there are more objections than approvals, @ABataev can you change the patch's title and remove the NFC part altogether?
5. Regarding justification for the need for volatile, I agree that the explanation of the problem is very short, but the solution makes sense given the nature of that problem - at least to me. I could have blocked the patch until the author finds a way to describe the issue in detail - in the meantime there would be a correctness problem in the library which could cause inexplicable failures/wrong results to the users. Even if the "fix" is more of a hack - which I don't believe - I would rather let this hack in than leave the correctness problem lying around indefinitely. From the description I understand that a thread may cache values from the parallelLevel array but at the same time those values change in RAM without the thread knowing about it. This leads to corrupt states where some threads have a stale view of the parallel level. The only way to demonstrate the problem is to post ptx code, which I think is too much in the scope of a review, that's why I don't see the point in providing more details. The author could point to two places in the code, one where a thread caches a value from that array and another one where another thread modifies the global value, but it is not always possible to track where those changes occur. This is in case I understood the problem correctly. If not, then @ABataev I am also requesting a detailed explanation of the problem because the current description leads to wrong conclusions.
6. 1, 4 and 5 represent my point of view, which is of course subjective, anyone could have a different mentality. But calling my reviews "reviews" and my activity here "suspicious" doesn't contribute at all. In pretty much the same way the "this is your problem, I've got the owner's approval" reply is inappropriate (and in no way do I endorse this "argument" as a reply to reviewers' requests).
7. Since February 2018 I have been the owner of the libomptarget project. The issue of libomptarget having no owner came up in one of our biweekly multi-company telecons, I was the only one who volunteered and at a subsequent F2F there was consensus for my appointment since I had written most of the code for the base library and the CUDA and generic-elf64 plugins.
8. Regarding the nvptx device runtime, certainly it is a different project from the rest of libomptarget. I think the RTL is still considered "coupled" with libomptarget for historical reasons only (both the base library, the plugins and the RTL were all written by the same people at the same company in the scope of the same project). But in terms of further development and expertise needed, it is a completely different kettle of fish. From my current position I can only contribute to the RTL by reviewing patches, it is true that I no longer do any development in this RTL. If anyone feels that the nvptx RTL repository should be managed/owned separately from the rest of libomptarget, please speak up and I will raise the issue at the next community telecon.

[jdoerfert]

In D62199#1515159, @grokos wrote:

Let me address what has been written so far from my perspective and request some changes since there has been demand for them:

1. 1, 4 and 5 represent my point of view, which is of course subjective, anyone could have a different mentality. But calling my reviews "reviews" and my activity here "suspicious" doesn't contribute at all. In pretty much the same way the "this is your problem, I've got the owner's approval" reply is inappropriate (and in no way do I endorse this "argument" as a reply to reviewers' requests).

I was not only referring to you, but it was not the right choice of words to begin with, you're right about that.
I will stick to arguments from now on. What I should have said is:

I think the following "NFC" commits are not "NFC" or the commit message is not matching: D61801, D61785, D61526, D61523, D56274
This is one of the commits that explain the problem they fix but do not provide a test: D56332
This commit message isn't explanatory and doesn't match the patch: D56290

And I am concerned because really quick LGTM for non-trivial changes happened before (not including this one or any of the aforementioned), e.g.,
in D53141 @Hahnfeld pointed out problems.

[ABataev]

In D62199#1515182, @jdoerfert wrote:

In D62199#1515159, @grokos wrote:

Let me address what has been written so far from my perspective and request some changes since there has been demand for them:

1. 1, 4 and 5 represent my point of view, which is of course subjective, anyone could have a different mentality. But calling my reviews "reviews" and my activity here "suspicious" doesn't contribute at all. In pretty much the same way the "this is your problem, I've got the owner's approval" reply is inappropriate (and in no way do I endorse this "argument" as a reply to reviewers' requests).

I was not only referring to you, but it was not the right choice of words to begin with, you're right about that.
I will stick to arguments from now on. What I should have said is:

I think the following "NFC" commits are not "NFC" or the commit message is not matching: D61801, D61785, D61526, D61523, D56274
This is one of the commits that explain the problem they fix but do not provide a test: D56332
This commit message isn't explanatory and doesn't match the patch: D56290

And I am concerned because really quick LGTM for non-trivial changes happened before (not including this one or any of the aforementioned), e.g.,
in D53141 @Hahnfeld pointed out problems.

These patches are really NFC. They do not add a new functionality, they simplify the existing one. And we already have the tests for these changes in NVPTX part of the libomptarget. These tests are not affected by these patches. That's why they are NFC.

[grokos]

In D62199#1515182, @jdoerfert wrote:

I agree with some of your complaints, I disagree with others:

I think the following "NFC" commits are not "NFC" or the commit message is not matching: D61801, D61785, D61526, D61523, D56274

I think all of them are NFC. D56274 is preventing LLVM from applying an optimization which can break syncthreads(). It replaces syncthreads() with the equivalent asm instruction. Apart from that bit (which doesn't change functionality, syncthreads() and the asm intrinsic are equivalent), the resulting code is identical and that's why the patch is marked as NFC. D61523 was superseded by D61801 (if you noticed, it hasn't been committed and I guess the author should abandon it). D61526 is an optimization of the thread-limit counter. It used to be a field in TaskDesc_items (and reside in global memory) and now it has been turned into a shared variable (residing in shared memory). The variable is accessed faster, but its functionality is the very same. From this perspective, I am OK with the NFC tag for that patch. Yes, it is a gray area once again as with volatile and no absolute answer could be given. D61785 does the same with the threads-in-team counter and D61801 does the same with the threads-in-parallel-region counter. For these 2 last patches you are right that the patch description and the commit messages are unclear or even wrong; D61785 should state explicitly that we are taking about the threads-in-team counter (instead of a vague "thread counter"), whereas in D61801 "handling of thread limit" is plain wrong and should be replaced with "handling of threads-in-parallel-region counter". I can edit the commit messages for the sake of clarity and consistency.

This is one of the commits that explain the problem they fix but do not provide a test: D56332

True, a test should be added. I was more concerned with fixing the dynamic scheduling bug, that was the important bit, but nonetheless that change should be accompanied by a corresponding dynamic scheduling test (especially since we now have a testing infrastructure in libomptarget).

This commit message isn't explanatory and doesn't match the patch: D56290

That's just a code polishing patch, I don't know what more the author could describe about it. He did some refactoring and cleanup, marked functions as inline and/or const and/or static etc. OK, just saying "Formatting" in the description may look poor, but then again there's the spirit of the law I mentioned earlier. I wouldn't make a big deal out of a poor message for such a patch.

And I am concerned because really quick LGTM for non-trivial changes happened before (not including this one or any of the aforementioned), e.g.,
in D53141 @Hahnfeld pointed out problems.

You are right, there was an issue I didn't catch. I never claimed to be infallible and no one is, that's why we like to have more than one pair of eyes review a patch. There are issues you may catch and I may miss and vice versa. Usually, if I don't have any comments about a patch I wait a day or two before accepting it, so that I give time to other reviewers to speak up.

[ABataev]

In D62199#1515159, @grokos wrote:

Let me address what has been written so far from my perspective and request some changes since there has been demand for them:

1. No doubt there are 2 changes here, unrelated to one another. Yes, they could be split into 2 patches to strictly comply with community guidelines. And I would prefer it if the author had submitted two distinct patches. On the other hand, I don't like missing the forest for the trees. This is a 4-line patch, I figured out it is OK to let both changes in at once. Sure, the small size of a patch is not a free pass to doing whatever we want, but in this case committing both changes in one go doesn't do any harm, at least none that I can think of so that I would make a big deal out of it. There's the letter of the law and the spirit of the law, in cases where these two collide I tend to go with the latter.
2. My bad that I didn't request a test for cache flushing, I was predisposed due to the NFC tag. I agree there is no way to test whether the cache has been flushed, but @ABataev you could provide a test case which fails without cache flushing so that you can demonstrate that this patch does indeed work. There must be such a test case, after all how did you figure out there is a problem there?

This change was not caused by the test failure or something similar, I just read the standard carefully and noticed that implicit flushing should happen each time the barrier construct is executed. I can try to add the test, anyway. Or provide a reference to a test, if it exist already in the codenase.

1. The NFC bit does not apply to the barrier fix. If we keep both changes in one patch, that should be reflected in the title. I agree that not being able to test a change does not make it NFC. You can call it NTC (non-testable change), but not NFC.
2. In case of volatile, it seems there's a gray line here between what constitutes a NFC patch. The way I see it, "volatile" prevents compiler optimizations, it doesn't change anything in the way this code works, the algorithms/data structures used etc., the only change is how the code is compiled. On the other hand, we could argue that since the compiled library behaves differently there is indeed a functional change. Once again, to me this is a gray line and there is no absolute answer. However, since here there are more objections than approvals, @ABataev can you change the patch's title and remove the NFC part altogether?

I will revert the commit and split the patch into 3 parts.

1. Regarding justification for the need for volatile, I agree that the explanation of the problem is very short, but the solution makes sense given the nature of that problem - at least to me. I could have blocked the patch until the author finds a way to describe the issue in detail - in the meantime there would be a correctness problem in the library which could cause inexplicable failures/wrong results to the users. Even if the "fix" is more of a hack - which I don't believe - I would rather let this hack in than leave the correctness problem lying around indefinitely. From the description I understand that a thread may cache values from the parallelLevel array but at the same time those values change in RAM without the thread knowing about it. This leads to corrupt states where some threads have a stale view of the parallel level. The only way to demonstrate the problem is to post ptx code, which I think is too much in the scope of a review, that's why I don't see the point in providing more details. The author could point to two places in the code, one where a thread caches a value from that array and another one where another thread modifies the global value, but it is not always possible to track where those changes occur. This is in case I understood the problem correctly. If not, then @ABataev I am also requesting a detailed explanation of the problem because the current description leads to wrong conclusions.
2. 1, 4 and 5 represent my point of view, which is of course subjective, anyone could have a different mentality. But calling my reviews "reviews" and my activity here "suspicious" doesn't contribute at all. In pretty much the same way the "this is your problem, I've got the owner's approval" reply is inappropriate (and in no way do I endorse this "argument" as a reply to reviewers' requests).
3. Since February 2018 I have been the owner of the libomptarget project. The issue of libomptarget having no owner came up in one of our biweekly multi-company telecons, I was the only one who volunteered and at a subsequent F2F there was consensus for my appointment since I had written most of the code for the base library and the CUDA and generic-elf64 plugins.
4. Regarding the nvptx device runtime, certainly it is a different project from the rest of libomptarget. I think the RTL is still considered "coupled" with libomptarget for historical reasons only (both the base library, the plugins and the RTL were all written by the same people at the same company in the scope of the same project). But in terms of further development and expertise needed, it is a completely different kettle of fish. From my current position I can only contribute to the RTL by reviewing patches, it is true that I no longer do any development in this RTL. If anyone feels that the nvptx RTL repository should be managed/owned separately from the rest of libomptarget, please speak up and I will raise the issue at the next community telecon.

[arsenm]

In D62199#1515385, @ABataev wrote:

These patches are really NFC. They do not add a new functionality, they simplify the existing one. And we already have the tests for these changes in NVPTX part of the libomptarget. These tests are not affected by these patches. That's why they are NFC.

This seems like a misunderstanding of the intent of NFC to me. NFC is for commits that are not supposed to have any observable effect, like refactoring. Anything that fixes a problem, while not providing what might be considered a feature or new functionality, is not NFC

[jdoerfert]

In D62199#1515511, @grokos wrote:

In D62199#1515182, @jdoerfert wrote:

I agree with some of your complaints, I disagree with others:

I think the following "NFC" commits are not "NFC" or the commit message is not matching: D61801, D61785, D61526, D61523, D56274

I think all of them are NFC.

I do not. Maybe we disagree what NFC is supposed to mean.
As I mentioned for this commit, I don't think a "FIX" can be "NFC". I also do not consider optimizations as NFC.
If something improves "performance" and "shared memory", it has to change functionality.

This is one of the commits that explain the problem they fix but do not provide a test: D56332

True, a test should be added. I was more concerned with fixing the dynamic scheduling bug, that was the important bit, but nonetheless that change should be accompanied by a corresponding dynamic scheduling test (especially since we now have a testing infrastructure in libomptarget).

I don't like the argument "we need to fix a bug first", neither for this patch nor for the one in the comment above.

This commit message isn't explanatory and doesn't match the patch: D56290

That's just a code polishing patch, I don't know what more the author could describe about it. He did some refactoring and cleanup, marked functions as inline and/or const and/or static etc. OK, just saying "Formatting" in the description may look poor, but then again there's the spirit of the law I mentioned earlier. I wouldn't make a big deal out of a poor message for such a patch.

I see a difference between "formating" and this patch. Functions are removed, code rewritten, a comment about a deadlock was removed because it seemed the problem might have been something else and it works now. That is not only more than formating but another example of a change that should not be rubber-stamped.

And I am concerned because really quick LGTM for non-trivial changes happened before (not including this one or any of the aforementioned), e.g.,
in D53141 @Hahnfeld pointed out problems.

You are right, there was an issue I didn't catch. I never claimed to be infallible and no one is, that's why we like to have more than one pair of eyes review a patch. There are issues you may catch and I may miss and vice versa.

I don't want anyone to be infallible, nor do I want to point out every mistake made by a person. This is not personal to me and I don't want you to get that impression.
What I want is to discuss a systematic problem in the review process. I don't claim this is intentional, but I claim there is *at least* one.

Usually, if I don't have any comments about a patch I wait a day or two before accepting it, so that I give time to other reviewers to speak up.

One problem is that it is not a day or two before accepting (see below). Another problem is that comments by others are not always taken into account (https://reviews.llvm.org/D51875#1229582), e.g., by checking in if they have been resolved (D60918).

Time (hours:minutes) between upload to LGTM for the patches I mentioned earlier:

D62199 3:30
D61801 6:30
D61785 4:30
D61526 0:50
D61523 2:00
D56274 0:15
D56332 6:30
D56290 19:06 (ping after 17:30)
D53141 0:21 (objection 6m later)

[hfinkel]

In D62199#1517071, @arsenm wrote:

In D62199#1515385, @ABataev wrote:

These patches are really NFC. They do not add a new functionality, they simplify the existing one. And we already have the tests for these changes in NVPTX part of the libomptarget. These tests are not affected by these patches. That's why they are NFC.

This seems like a misunderstanding of the intent of NFC to me. NFC is for commits that are not supposed to have any observable effect, like refactoring. Anything that fixes a problem, while not providing what might be considered a feature or new functionality, is not NFC

@arsenm is certainly correct. The developer policy speaks to this as well, stating, "Avoid committing formatting- or whitespace-only changes outside of code you plan to make subsequent changes to. Also, try to separate formatting or whitespace changes from functional changes, either by correcting the format first (ideally) or afterward. Such changes should be highly localized and the commit message should clearly state that the commit is not intended to change functionality, usually by stating it is NFC." The policy does not explicitly state that this is the only valid use of NFC, and it is sometimes used for refactoring as well, but if a bug is fixed, we don't consider that to be NFC (because the effective functionality of the product has changed).

[ABataev]

In D62199#1517078, @hfinkel wrote:

In D62199#1517071, @arsenm wrote:

In D62199#1515385, @ABataev wrote:

These patches are really NFC. They do not add a new functionality, they simplify the existing one. And we already have the tests for these changes in NVPTX part of the libomptarget. These tests are not affected by these patches. That's why they are NFC.

This seems like a misunderstanding of the intent of NFC to me. NFC is for commits that are not supposed to have any observable effect, like refactoring. Anything that fixes a problem, while not providing what might be considered a feature or new functionality, is not NFC

@arsenm is certainly correct. The developer policy speaks to this as well, stating, "Avoid committing formatting- or whitespace-only changes outside of code you plan to make subsequent changes to. Also, try to separate formatting or whitespace changes from functional changes, either by correcting the format first (ideally) or afterward. Such changes should be highly localized and the commit message should clearly state that the commit is not intended to change functionality, usually by stating it is NFC." The policy does not explicitly state that this is the only valid use of NFC, and it is sometimes used for refactoring as well, but if a bug is fixed, we don't consider that to be NFC (because the effective functionality of the product has changed).

In some cases a "fix" can be NFC, if it fixes the bug in the implementation of the non-functional requirement, for example.
Functional changes - the changes that change functionality, i.e. change the expected output for the subset of inputs. Generally speaking, "volatile" qualifier does not change the logical dependency between inputs and expected outputs, it changes how the compiler handles such variables and optimizes them. But the logic remains the same.
I'm not sure that some internal compiler/programming language requirements, can be considered as functional.
Functional requirements are supported by non-functional requirements (also known as "quality requirements"), which impose constraints on the design or implementation. In this case, "volatile" modifier can be considered as the implementation constraint.

[hfinkel]

In D62199#1517095, @ABataev wrote:

In D62199#1517078, @hfinkel wrote:

In D62199#1517071, @arsenm wrote:

In D62199#1515385, @ABataev wrote:

These patches are really NFC. They do not add a new functionality, they simplify the existing one. And we already have the tests for these changes in NVPTX part of the libomptarget. These tests are not affected by these patches. That's why they are NFC.

This seems like a misunderstanding of the intent of NFC to me. NFC is for commits that are not supposed to have any observable effect, like refactoring. Anything that fixes a problem, while not providing what might be considered a feature or new functionality, is not NFC

@arsenm is certainly correct. The developer policy speaks to this as well, stating, "Avoid committing formatting- or whitespace-only changes outside of code you plan to make subsequent changes to. Also, try to separate formatting or whitespace changes from functional changes, either by correcting the format first (ideally) or afterward. Such changes should be highly localized and the commit message should clearly state that the commit is not intended to change functionality, usually by stating it is NFC." The policy does not explicitly state that this is the only valid use of NFC, and it is sometimes used for refactoring as well, but if a bug is fixed, we don't consider that to be NFC (because the effective functionality of the product has changed).

In some cases a "fix" can be NFC, if it fixes the bug in the implementation of the non-functional requirement, for example.
Functional changes - the changes that change functionality, i.e. change the expected output for the subset of inputs. Generally speaking, "volatile" qualifier does not change the logical dependency between inputs and expected outputs, it changes how the compiler handles such variables and optimizes them. But the logic remains the same.
I'm not sure that some internal compiler/programming language requirements, can be considered as functional.
Functional requirements are supported by non-functional requirements (also known as "quality requirements"), which impose constraints on the design or implementation. In this case, "volatile" modifier can be considered as the implementation constraint.

I understand what you're saying, but that's not how the LLVM community uses the term NFC. One way to think about it is the following: If any change in the behavior of the system: functionality, correctness, performance (at compile time or run time), could be bisected to this commit, then it's not NFC. NFC is for things like reformatting code - unless a mistake was made, no observable behavior of the system should bisect to a commit that only reformats code. In some cases, *small* refactorings are labeled as NFC. If it were to turn out that a refactoring happened to introduce a compile-time-performance change, for example, then labeling the commit as NFC would have been in error (likely accidentally). Thus, as the LLVM community uses the term NFC, a "fix" can never be NFC.

[ABataev]

In D62199#1517099, @hfinkel wrote:

In D62199#1517095, @ABataev wrote:

In D62199#1517078, @hfinkel wrote:

In D62199#1517071, @arsenm wrote:

In D62199#1515385, @ABataev wrote:

These patches are really NFC. They do not add a new functionality, they simplify the existing one. And we already have the tests for these changes in NVPTX part of the libomptarget. These tests are not affected by these patches. That's why they are NFC.

This seems like a misunderstanding of the intent of NFC to me. NFC is for commits that are not supposed to have any observable effect, like refactoring. Anything that fixes a problem, while not providing what might be considered a feature or new functionality, is not NFC

@arsenm is certainly correct. The developer policy speaks to this as well, stating, "Avoid committing formatting- or whitespace-only changes outside of code you plan to make subsequent changes to. Also, try to separate formatting or whitespace changes from functional changes, either by correcting the format first (ideally) or afterward. Such changes should be highly localized and the commit message should clearly state that the commit is not intended to change functionality, usually by stating it is NFC." The policy does not explicitly state that this is the only valid use of NFC, and it is sometimes used for refactoring as well, but if a bug is fixed, we don't consider that to be NFC (because the effective functionality of the product has changed).

In some cases a "fix" can be NFC, if it fixes the bug in the implementation of the non-functional requirement, for example.
Functional changes - the changes that change functionality, i.e. change the expected output for the subset of inputs. Generally speaking, "volatile" qualifier does not change the logical dependency between inputs and expected outputs, it changes how the compiler handles such variables and optimizes them. But the logic remains the same.
I'm not sure that some internal compiler/programming language requirements, can be considered as functional.
Functional requirements are supported by non-functional requirements (also known as "quality requirements"), which impose constraints on the design or implementation. In this case, "volatile" modifier can be considered as the implementation constraint.

I understand what you're saying, but that's not how the LLVM community uses the term NFC. One way to think about it is the following: If any change in the behavior of the system: functionality, correctness, performance (at compile time or run time), could be bisected to this commit, then it's not NFC. NFC is for things like reformatting code - unless a mistake was made, no observable behavior of the system should bisect to a commit that only reformats code. In some cases, *small* refactorings are labeled as NFC. If it were to turn out that a refactoring happened to introduce a compile-time-performance change, for example, then labeling the commit as NFC would have been in error (likely accidentally). Thus, as the LLVM community uses the term NFC, a "fix" can never be NFC.

Let me disagree with you. At first, you said "The policy does not explicitly state that this is the only valid use of NFC, and it is sometimes used for refactoring as well." Now you're saying that this only for formatting.
Functionality is well determined term. I don't think there is a different meaning of it in the community.

[hfinkel]

In D62199#1517144, @ABataev wrote:

In D62199#1517099, @hfinkel wrote:

In D62199#1517095, @ABataev wrote:

In D62199#1517078, @hfinkel wrote:

In D62199#1517071, @arsenm wrote:

In D62199#1515385, @ABataev wrote:

These patches are really NFC. They do not add a new functionality, they simplify the existing one. And we already have the tests for these changes in NVPTX part of the libomptarget. These tests are not affected by these patches. That's why they are NFC.

This seems like a misunderstanding of the intent of NFC to me. NFC is for commits that are not supposed to have any observable effect, like refactoring. Anything that fixes a problem, while not providing what might be considered a feature or new functionality, is not NFC

@arsenm is certainly correct. The developer policy speaks to this as well, stating, "Avoid committing formatting- or whitespace-only changes outside of code you plan to make subsequent changes to. Also, try to separate formatting or whitespace changes from functional changes, either by correcting the format first (ideally) or afterward. Such changes should be highly localized and the commit message should clearly state that the commit is not intended to change functionality, usually by stating it is NFC." The policy does not explicitly state that this is the only valid use of NFC, and it is sometimes used for refactoring as well, but if a bug is fixed, we don't consider that to be NFC (because the effective functionality of the product has changed).

In some cases a "fix" can be NFC, if it fixes the bug in the implementation of the non-functional requirement, for example.
Functional changes - the changes that change functionality, i.e. change the expected output for the subset of inputs. Generally speaking, "volatile" qualifier does not change the logical dependency between inputs and expected outputs, it changes how the compiler handles such variables and optimizes them. But the logic remains the same.
I'm not sure that some internal compiler/programming language requirements, can be considered as functional.
Functional requirements are supported by non-functional requirements (also known as "quality requirements"), which impose constraints on the design or implementation. In this case, "volatile" modifier can be considered as the implementation constraint.

I understand what you're saying, but that's not how the LLVM community uses the term NFC. One way to think about it is the following: If any change in the behavior of the system: functionality, correctness, performance (at compile time or run time), could be bisected to this commit, then it's not NFC. NFC is for things like reformatting code - unless a mistake was made, no observable behavior of the system should bisect to a commit that only reformats code. In some cases, *small* refactorings are labeled as NFC. If it were to turn out that a refactoring happened to introduce a compile-time-performance change, for example, then labeling the commit as NFC would have been in error (likely accidentally). Thus, as the LLVM community uses the term NFC, a "fix" can never be NFC.

Let me disagree with you. At first, you said "The policy does not explicitly state that this is the only valid use of NFC, and it is sometimes used for refactoring as well." Now you're saying that this only for formatting.
Functionality is well determined term. I don't think there is a different meaning of it in the community.

"Now you're saying that this only for formatting." - I specifically did not say that. I did specifically say, "it is sometimes used for refactoring as well." The functionality of the compiler and its runtime libraries is changed if a bug is fixed, Alexey. I don't believe there's disagreement in the community on that point.

[ABataev]

In D62199#1517150, @hfinkel wrote:

In D62199#1517144, @ABataev wrote:

In D62199#1517099, @hfinkel wrote:

In D62199#1517095, @ABataev wrote:

In D62199#1517078, @hfinkel wrote:

In D62199#1517071, @arsenm wrote:

In D62199#1515385, @ABataev wrote:

These patches are really NFC. They do not add a new functionality, they simplify the existing one. And we already have the tests for these changes in NVPTX part of the libomptarget. These tests are not affected by these patches. That's why they are NFC.

This seems like a misunderstanding of the intent of NFC to me. NFC is for commits that are not supposed to have any observable effect, like refactoring. Anything that fixes a problem, while not providing what might be considered a feature or new functionality, is not NFC

@arsenm is certainly correct. The developer policy speaks to this as well, stating, "Avoid committing formatting- or whitespace-only changes outside of code you plan to make subsequent changes to. Also, try to separate formatting or whitespace changes from functional changes, either by correcting the format first (ideally) or afterward. Such changes should be highly localized and the commit message should clearly state that the commit is not intended to change functionality, usually by stating it is NFC." The policy does not explicitly state that this is the only valid use of NFC, and it is sometimes used for refactoring as well, but if a bug is fixed, we don't consider that to be NFC (because the effective functionality of the product has changed).

In some cases a "fix" can be NFC, if it fixes the bug in the implementation of the non-functional requirement, for example.
Functional changes - the changes that change functionality, i.e. change the expected output for the subset of inputs. Generally speaking, "volatile" qualifier does not change the logical dependency between inputs and expected outputs, it changes how the compiler handles such variables and optimizes them. But the logic remains the same.
I'm not sure that some internal compiler/programming language requirements, can be considered as functional.
Functional requirements are supported by non-functional requirements (also known as "quality requirements"), which impose constraints on the design or implementation. In this case, "volatile" modifier can be considered as the implementation constraint.

I understand what you're saying, but that's not how the LLVM community uses the term NFC. One way to think about it is the following: If any change in the behavior of the system: functionality, correctness, performance (at compile time or run time), could be bisected to this commit, then it's not NFC. NFC is for things like reformatting code - unless a mistake was made, no observable behavior of the system should bisect to a commit that only reformats code. In some cases, *small* refactorings are labeled as NFC. If it were to turn out that a refactoring happened to introduce a compile-time-performance change, for example, then labeling the commit as NFC would have been in error (likely accidentally). Thus, as the LLVM community uses the term NFC, a "fix" can never be NFC.

Let me disagree with you. At first, you said "The policy does not explicitly state that this is the only valid use of NFC, and it is sometimes used for refactoring as well." Now you're saying that this only for formatting.
Functionality is well determined term. I don't think there is a different meaning of it in the community.

"Now you're saying that this only for formatting." - I specifically did not say that. I did specifically say, "it is sometimes used for refactoring as well." The functionality of the compiler and its runtime libraries is changed if a bug is fixed, Alexey. I don't believe there's disagreement in the community on that point.

Hal, again, you arguing against fotmal definition of functional and non-functional, functional requirement and non-functional requirement, in the computer science. If the "fix" fixes the implementation of the non-functional requirement, it can be a non-functional fix.
I'm not saying that you cannot redefine those formal definitions and use your own. You absolutely can do this! But you need to define the new meaning of those old terms strictly and introduce the new terms for the redfined ones. Only after that we can use these terms with their new meanings.

[hfinkel]

In D62199#1517258, @ABataev wrote:

In D62199#1517150, @hfinkel wrote:

In D62199#1517144, @ABataev wrote:

In D62199#1517099, @hfinkel wrote:

In D62199#1517095, @ABataev wrote:

In D62199#1517078, @hfinkel wrote:

In D62199#1517071, @arsenm wrote:

In D62199#1515385, @ABataev wrote:

These patches are really NFC. They do not add a new functionality, they simplify the existing one. And we already have the tests for these changes in NVPTX part of the libomptarget. These tests are not affected by these patches. That's why they are NFC.

This seems like a misunderstanding of the intent of NFC to me. NFC is for commits that are not supposed to have any observable effect, like refactoring. Anything that fixes a problem, while not providing what might be considered a feature or new functionality, is not NFC

@arsenm is certainly correct. The developer policy speaks to this as well, stating, "Avoid committing formatting- or whitespace-only changes outside of code you plan to make subsequent changes to. Also, try to separate formatting or whitespace changes from functional changes, either by correcting the format first (ideally) or afterward. Such changes should be highly localized and the commit message should clearly state that the commit is not intended to change functionality, usually by stating it is NFC." The policy does not explicitly state that this is the only valid use of NFC, and it is sometimes used for refactoring as well, but if a bug is fixed, we don't consider that to be NFC (because the effective functionality of the product has changed).

In some cases a "fix" can be NFC, if it fixes the bug in the implementation of the non-functional requirement, for example.
Functional changes - the changes that change functionality, i.e. change the expected output for the subset of inputs. Generally speaking, "volatile" qualifier does not change the logical dependency between inputs and expected outputs, it changes how the compiler handles such variables and optimizes them. But the logic remains the same.
I'm not sure that some internal compiler/programming language requirements, can be considered as functional.
Functional requirements are supported by non-functional requirements (also known as "quality requirements"), which impose constraints on the design or implementation. In this case, "volatile" modifier can be considered as the implementation constraint.

I understand what you're saying, but that's not how the LLVM community uses the term NFC. One way to think about it is the following: If any change in the behavior of the system: functionality, correctness, performance (at compile time or run time), could be bisected to this commit, then it's not NFC. NFC is for things like reformatting code - unless a mistake was made, no observable behavior of the system should bisect to a commit that only reformats code. In some cases, *small* refactorings are labeled as NFC. If it were to turn out that a refactoring happened to introduce a compile-time-performance change, for example, then labeling the commit as NFC would have been in error (likely accidentally). Thus, as the LLVM community uses the term NFC, a "fix" can never be NFC.

Let me disagree with you. At first, you said "The policy does not explicitly state that this is the only valid use of NFC, and it is sometimes used for refactoring as well." Now you're saying that this only for formatting.
Functionality is well determined term. I don't think there is a different meaning of it in the community.

"Now you're saying that this only for formatting." - I specifically did not say that. I did specifically say, "it is sometimes used for refactoring as well." The functionality of the compiler and its runtime libraries is changed if a bug is fixed, Alexey. I don't believe there's disagreement in the community on that point.

Hal, again, you arguing against fotmal definition of functional and non-functional, functional requirement and non-functional requirement, in the computer science. If the "fix" fixes the implementation of the non-functional requirement, it can be a non-functional fix.
I'm not saying that you cannot redefine those formal definitions and use your own. You absolutely can do this! But you need to define the new meaning of those old terms strictly and introduce the new terms for the redfined ones. Only after that we can use these terms with their new meanings.

Alexey, I feel like we're talking past each other. Tagging commits as NFC is not a theoretical exercise, it is intended to be a practically-usable note to those reviewing the commit history that, "this commit shouldn't have changed any observable behavior of the system.". When you say that a commit "they simplify the existing [functionality]", that sounds like refactoring, and that could very well be NFC. I don't think we're in disagreement on that. The disagreement seems to be over things like this:

D56274 is preventing LLVM from applying an optimization which can break syncthreads(). It replaces syncthreads() with the equivalent asm instruction. Apart from that bit (which doesn't change functionality, syncthreads() and the asm intrinsic are equivalent), the resulting code is identical

And we have several underlying issues. It sounds like you're saying that you've never observed a difference between using the intrinsic and the inline asm, but you're changing to the inline asm because the intrinsic is "known to be buggy." Thus, while you've never observed the bug yourself in this context, the new form simply seems safer than the old form. Is that correct? I can see an argument for marking such a change as NFC, but should someone compile the code with a compiler and in some configuration where the bug does manifest, it would not be NFC *for them*. Because this risk exists, we should not mark the change as NFC. If you have observed the bug, then the change is certainly not NFC (because the behavior of the system changed from before to after the commit, at least for some system configurations). We also shouldn't tag commits adding volatile qualifiers as NFC for the same reason (you're doing this because the commit might change the behavior of the system for someone, even if the exact set of such configurations is unknown).

[ABataev]

In D62199#1517419, @hfinkel wrote:

In D62199#1517258, @ABataev wrote:

In D62199#1517150, @hfinkel wrote:

In D62199#1517144, @ABataev wrote:

In D62199#1517099, @hfinkel wrote:

In D62199#1517095, @ABataev wrote:

In D62199#1517078, @hfinkel wrote:

In D62199#1517071, @arsenm wrote:

In D62199#1515385, @ABataev wrote:

These patches are really NFC. They do not add a new functionality, they simplify the existing one. And we already have the tests for these changes in NVPTX part of the libomptarget. These tests are not affected by these patches. That's why they are NFC.

This seems like a misunderstanding of the intent of NFC to me. NFC is for commits that are not supposed to have any observable effect, like refactoring. Anything that fixes a problem, while not providing what might be considered a feature or new functionality, is not NFC

@arsenm is certainly correct. The developer policy speaks to this as well, stating, "Avoid committing formatting- or whitespace-only changes outside of code you plan to make subsequent changes to. Also, try to separate formatting or whitespace changes from functional changes, either by correcting the format first (ideally) or afterward. Such changes should be highly localized and the commit message should clearly state that the commit is not intended to change functionality, usually by stating it is NFC." The policy does not explicitly state that this is the only valid use of NFC, and it is sometimes used for refactoring as well, but if a bug is fixed, we don't consider that to be NFC (because the effective functionality of the product has changed).

In some cases a "fix" can be NFC, if it fixes the bug in the implementation of the non-functional requirement, for example.
Functional changes - the changes that change functionality, i.e. change the expected output for the subset of inputs. Generally speaking, "volatile" qualifier does not change the logical dependency between inputs and expected outputs, it changes how the compiler handles such variables and optimizes them. But the logic remains the same.
I'm not sure that some internal compiler/programming language requirements, can be considered as functional.
Functional requirements are supported by non-functional requirements (also known as "quality requirements"), which impose constraints on the design or implementation. In this case, "volatile" modifier can be considered as the implementation constraint.

I understand what you're saying, but that's not how the LLVM community uses the term NFC. One way to think about it is the following: If any change in the behavior of the system: functionality, correctness, performance (at compile time or run time), could be bisected to this commit, then it's not NFC. NFC is for things like reformatting code - unless a mistake was made, no observable behavior of the system should bisect to a commit that only reformats code. In some cases, *small* refactorings are labeled as NFC. If it were to turn out that a refactoring happened to introduce a compile-time-performance change, for example, then labeling the commit as NFC would have been in error (likely accidentally). Thus, as the LLVM community uses the term NFC, a "fix" can never be NFC.

Let me disagree with you. At first, you said "The policy does not explicitly state that this is the only valid use of NFC, and it is sometimes used for refactoring as well." Now you're saying that this only for formatting.
Functionality is well determined term. I don't think there is a different meaning of it in the community.

"Now you're saying that this only for formatting." - I specifically did not say that. I did specifically say, "it is sometimes used for refactoring as well." The functionality of the compiler and its runtime libraries is changed if a bug is fixed, Alexey. I don't believe there's disagreement in the community on that point.

Hal, again, you arguing against fotmal definition of functional and non-functional, functional requirement and non-functional requirement, in the computer science. If the "fix" fixes the implementation of the non-functional requirement, it can be a non-functional fix.
I'm not saying that you cannot redefine those formal definitions and use your own. You absolutely can do this! But you need to define the new meaning of those old terms strictly and introduce the new terms for the redfined ones. Only after that we can use these terms with their new meanings.

Alexey, I feel like we're talking past each other. Tagging commits as NFC is not a theoretical exercise, it is intended to be a practically-usable note to those reviewing the commit history that, "this commit shouldn't have changed any observable behavior of the system.". When you say that a commit "they simplify the existing [functionality]", that sounds like refactoring, and that could very well be NFC. I don't think we're in disagreement on that. The disagreement seems to be over things like this:

D56274 is preventing LLVM from applying an optimization which can break syncthreads(). It replaces syncthreads() with the equivalent asm instruction. Apart from that bit (which doesn't change functionality, syncthreads() and the asm intrinsic are equivalent), the resulting code is identical

And we have several underlying issues. It sounds like you're saying that you've never observed a difference between using the intrinsic and the inline asm, but you're changing to the inline asm because the intrinsic is "known to be buggy." Thus, while you've never observed the bug yourself in this context, the new form simply seems safer than the old form. Is that correct? I can see an argument for marking such a change as NFC, but should someone compile the code with a compiler and in some configuration where the bug does manifest, it would not be NFC *for them*. Because this risk exists, we should not mark the change as NFC. If you have observed the bug, then the change is certainly not NFC (because the behavior of the system changed from before to after the commit, at least for some system configurations). We also shouldn't tag commits adding volatile qualifiers as NFC for the same reason (you're doing this because the commit might change the behavior of the system for someone, even if the exact set of such configurations is unknown).

Hal, you don't see a difference between functional and non-functional requirements. The fact, that behavior ID changed because of the optimization, is the implementation constraint. It is non-functional requirement at the moment to use the asm form instead of the intrinsic. Because, there is no logical changes in the dependency between inputs and outputs. These 2 forms has the same functionality. Thus, it is non-functional requirement to use the second form instead of the original one at the moment. And because of that this is a non-functional change, because the logical dependency between inputs and outputs remains the same it was before.
The same is with volatile. This is just a compiler constraint to use this modifier, this not a functional requirement. And thus the change can be marked as non-functional.

[hfinkel]

In D62199#1517434, @ABataev wrote:

In D62199#1517419, @hfinkel wrote:

In D62199#1517258, @ABataev wrote:

In D62199#1517150, @hfinkel wrote:

In D62199#1517144, @ABataev wrote:

In D62199#1517099, @hfinkel wrote:

In D62199#1517095, @ABataev wrote:

In D62199#1517078, @hfinkel wrote:

In D62199#1517071, @arsenm wrote:

In D62199#1515385, @ABataev wrote:

These patches are really NFC. They do not add a new functionality, they simplify the existing one. And we already have the tests for these changes in NVPTX part of the libomptarget. These tests are not affected by these patches. That's why they are NFC.

This seems like a misunderstanding of the intent of NFC to me. NFC is for commits that are not supposed to have any observable effect, like refactoring. Anything that fixes a problem, while not providing what might be considered a feature or new functionality, is not NFC

@arsenm is certainly correct. The developer policy speaks to this as well, stating, "Avoid committing formatting- or whitespace-only changes outside of code you plan to make subsequent changes to. Also, try to separate formatting or whitespace changes from functional changes, either by correcting the format first (ideally) or afterward. Such changes should be highly localized and the commit message should clearly state that the commit is not intended to change functionality, usually by stating it is NFC." The policy does not explicitly state that this is the only valid use of NFC, and it is sometimes used for refactoring as well, but if a bug is fixed, we don't consider that to be NFC (because the effective functionality of the product has changed).

In some cases a "fix" can be NFC, if it fixes the bug in the implementation of the non-functional requirement, for example.
Functional changes - the changes that change functionality, i.e. change the expected output for the subset of inputs. Generally speaking, "volatile" qualifier does not change the logical dependency between inputs and expected outputs, it changes how the compiler handles such variables and optimizes them. But the logic remains the same.
I'm not sure that some internal compiler/programming language requirements, can be considered as functional.
Functional requirements are supported by non-functional requirements (also known as "quality requirements"), which impose constraints on the design or implementation. In this case, "volatile" modifier can be considered as the implementation constraint.

I understand what you're saying, but that's not how the LLVM community uses the term NFC. One way to think about it is the following: If any change in the behavior of the system: functionality, correctness, performance (at compile time or run time), could be bisected to this commit, then it's not NFC. NFC is for things like reformatting code - unless a mistake was made, no observable behavior of the system should bisect to a commit that only reformats code. In some cases, *small* refactorings are labeled as NFC. If it were to turn out that a refactoring happened to introduce a compile-time-performance change, for example, then labeling the commit as NFC would have been in error (likely accidentally). Thus, as the LLVM community uses the term NFC, a "fix" can never be NFC.

Let me disagree with you. At first, you said "The policy does not explicitly state that this is the only valid use of NFC, and it is sometimes used for refactoring as well." Now you're saying that this only for formatting.
Functionality is well determined term. I don't think there is a different meaning of it in the community.

"Now you're saying that this only for formatting." - I specifically did not say that. I did specifically say, "it is sometimes used for refactoring as well." The functionality of the compiler and its runtime libraries is changed if a bug is fixed, Alexey. I don't believe there's disagreement in the community on that point.

Hal, again, you arguing against fotmal definition of functional and non-functional, functional requirement and non-functional requirement, in the computer science. If the "fix" fixes the implementation of the non-functional requirement, it can be a non-functional fix.
I'm not saying that you cannot redefine those formal definitions and use your own. You absolutely can do this! But you need to define the new meaning of those old terms strictly and introduce the new terms for the redfined ones. Only after that we can use these terms with their new meanings.

Alexey, I feel like we're talking past each other. Tagging commits as NFC is not a theoretical exercise, it is intended to be a practically-usable note to those reviewing the commit history that, "this commit shouldn't have changed any observable behavior of the system.". When you say that a commit "they simplify the existing [functionality]", that sounds like refactoring, and that could very well be NFC. I don't think we're in disagreement on that. The disagreement seems to be over things like this:

D56274 is preventing LLVM from applying an optimization which can break syncthreads(). It replaces syncthreads() with the equivalent asm instruction. Apart from that bit (which doesn't change functionality, syncthreads() and the asm intrinsic are equivalent), the resulting code is identical

And we have several underlying issues. It sounds like you're saying that you've never observed a difference between using the intrinsic and the inline asm, but you're changing to the inline asm because the intrinsic is "known to be buggy." Thus, while you've never observed the bug yourself in this context, the new form simply seems safer than the old form. Is that correct? I can see an argument for marking such a change as NFC, but should someone compile the code with a compiler and in some configuration where the bug does manifest, it would not be NFC *for them*. Because this risk exists, we should not mark the change as NFC. If you have observed the bug, then the change is certainly not NFC (because the behavior of the system changed from before to after the commit, at least for some system configurations). We also shouldn't tag commits adding volatile qualifiers as NFC for the same reason (you're doing this because the commit might change the behavior of the system for someone, even if the exact set of such configurations is unknown).

Hal, you don't see a difference between functional and non-functional requirements. The fact, that behavior ID changed because of the optimization, is the implementation constraint. It is non-functional requirement at the moment to use the asm form instead of the intrinsic. Because, there is no logical changes in the dependency between inputs and outputs. These 2 forms has the same functionality. Thus, it is non-functional requirement to use the second form instead of the original one at the moment. And because of that this is a non-functional change, because the logical dependency between inputs and outputs remains the same it was before.
The same is with volatile. This is just a compiler constraint to use this modifier, this not a functional requirement. And thus the change can be marked as non-functional.

Alexey, I believe that I understand the distinction that you are making. That's not the way the term is used within the LLVM community. The observable behavior of the compiler and its runtime libraries is what matters. The fact that some implementation technique should not be required to make that observable behavior correct is irrelevant in our usage. The introduction of that technique in order to make the observable behavior correct, where it was not previously correct, is a functionality-changing change to the compiler. Thus, it is not NFC.

[ABataev]

In D62199#1517437, @hfinkel wrote:

In D62199#1517434, @ABataev wrote:

In D62199#1517419, @hfinkel wrote:

In D62199#1517258, @ABataev wrote:

In D62199#1517150, @hfinkel wrote:

In D62199#1517144, @ABataev wrote:

In D62199#1517099, @hfinkel wrote:

In D62199#1517095, @ABataev wrote:

In D62199#1517078, @hfinkel wrote:

In D62199#1517071, @arsenm wrote:

In D62199#1515385, @ABataev wrote:

These patches are really NFC. They do not add a new functionality, they simplify the existing one. And we already have the tests for these changes in NVPTX part of the libomptarget. These tests are not affected by these patches. That's why they are NFC.

This seems like a misunderstanding of the intent of NFC to me. NFC is for commits that are not supposed to have any observable effect, like refactoring. Anything that fixes a problem, while not providing what might be considered a feature or new functionality, is not NFC

@arsenm is certainly correct. The developer policy speaks to this as well, stating, "Avoid committing formatting- or whitespace-only changes outside of code you plan to make subsequent changes to. Also, try to separate formatting or whitespace changes from functional changes, either by correcting the format first (ideally) or afterward. Such changes should be highly localized and the commit message should clearly state that the commit is not intended to change functionality, usually by stating it is NFC." The policy does not explicitly state that this is the only valid use of NFC, and it is sometimes used for refactoring as well, but if a bug is fixed, we don't consider that to be NFC (because the effective functionality of the product has changed).

In some cases a "fix" can be NFC, if it fixes the bug in the implementation of the non-functional requirement, for example.
Functional changes - the changes that change functionality, i.e. change the expected output for the subset of inputs. Generally speaking, "volatile" qualifier does not change the logical dependency between inputs and expected outputs, it changes how the compiler handles such variables and optimizes them. But the logic remains the same.
I'm not sure that some internal compiler/programming language requirements, can be considered as functional.
Functional requirements are supported by non-functional requirements (also known as "quality requirements"), which impose constraints on the design or implementation. In this case, "volatile" modifier can be considered as the implementation constraint.

I understand what you're saying, but that's not how the LLVM community uses the term NFC. One way to think about it is the following: If any change in the behavior of the system: functionality, correctness, performance (at compile time or run time), could be bisected to this commit, then it's not NFC. NFC is for things like reformatting code - unless a mistake was made, no observable behavior of the system should bisect to a commit that only reformats code. In some cases, *small* refactorings are labeled as NFC. If it were to turn out that a refactoring happened to introduce a compile-time-performance change, for example, then labeling the commit as NFC would have been in error (likely accidentally). Thus, as the LLVM community uses the term NFC, a "fix" can never be NFC.

Let me disagree with you. At first, you said "The policy does not explicitly state that this is the only valid use of NFC, and it is sometimes used for refactoring as well." Now you're saying that this only for formatting.
Functionality is well determined term. I don't think there is a different meaning of it in the community.

"Now you're saying that this only for formatting." - I specifically did not say that. I did specifically say, "it is sometimes used for refactoring as well." The functionality of the compiler and its runtime libraries is changed if a bug is fixed, Alexey. I don't believe there's disagreement in the community on that point.

Hal, again, you arguing against fotmal definition of functional and non-functional, functional requirement and non-functional requirement, in the computer science. If the "fix" fixes the implementation of the non-functional requirement, it can be a non-functional fix.
I'm not saying that you cannot redefine those formal definitions and use your own. You absolutely can do this! But you need to define the new meaning of those old terms strictly and introduce the new terms for the redfined ones. Only after that we can use these terms with their new meanings.

Alexey, I feel like we're talking past each other. Tagging commits as NFC is not a theoretical exercise, it is intended to be a practically-usable note to those reviewing the commit history that, "this commit shouldn't have changed any observable behavior of the system.". When you say that a commit "they simplify the existing [functionality]", that sounds like refactoring, and that could very well be NFC. I don't think we're in disagreement on that. The disagreement seems to be over things like this:

D56274 is preventing LLVM from applying an optimization which can break syncthreads(). It replaces syncthreads() with the equivalent asm instruction. Apart from that bit (which doesn't change functionality, syncthreads() and the asm intrinsic are equivalent), the resulting code is identical

And we have several underlying issues. It sounds like you're saying that you've never observed a difference between using the intrinsic and the inline asm, but you're changing to the inline asm because the intrinsic is "known to be buggy." Thus, while you've never observed the bug yourself in this context, the new form simply seems safer than the old form. Is that correct? I can see an argument for marking such a change as NFC, but should someone compile the code with a compiler and in some configuration where the bug does manifest, it would not be NFC *for them*. Because this risk exists, we should not mark the change as NFC. If you have observed the bug, then the change is certainly not NFC (because the behavior of the system changed from before to after the commit, at least for some system configurations). We also shouldn't tag commits adding volatile qualifiers as NFC for the same reason (you're doing this because the commit might change the behavior of the system for someone, even if the exact set of such configurations is unknown).

Hal, you don't see a difference between functional and non-functional requirements. The fact, that behavior ID changed because of the optimization, is the implementation constraint. It is non-functional requirement at the moment to use the asm form instead of the intrinsic. Because, there is no logical changes in the dependency between inputs and outputs. These 2 forms has the same functionality. Thus, it is non-functional requirement to use the second form instead of the original one at the moment. And because of that this is a non-functional change, because the logical dependency between inputs and outputs remains the same it was before.
The same is with volatile. This is just a compiler constraint to use this modifier, this not a functional requirement. And thus the change can be marked as non-functional.

Alexey, I believe that I understand the distinction that you are making. That's not the way the term is used within the LLVM community. The observable behavior of the compiler and its runtime libraries is what matters. The fact that some implementation technique should not be required to make that observable behavior correct is irrelevant in our usage. The introduction of that technique in order to make the observable behavior correct, where it was not previously correct, is a functionality-changing change to the compiler. Thus, it is not NFC.

It is better to express this explicitly in the policy, because it differs from the classical definition.

[Hahnfeld]

In D62199#1517460, @ABataev wrote:

In D62199#1517437, @hfinkel wrote:

Alexey, I believe that I understand the distinction that you are making. That's not the way the term is used within the LLVM community. The observable behavior of the compiler and its runtime libraries is what matters. The fact that some implementation technique should not be required to make that observable behavior correct is irrelevant in our usage. The introduction of that technique in order to make the observable behavior correct, where it was not previously correct, is a functionality-changing change to the compiler. Thus, it is not NFC.

It is better to express this explicitly in the policy, because it differs from the classical definition.

See https://llvm.org/docs/Lexicon.html#nfc:

“No functional change”. Used in a commit message to indicate that a patch is a pure refactoring/cleanup. Usually used in the first line, so it is visible without opening the actual commit email.

AFAICS that has nothing to do with "non-functional" requirements. [ We could continue arguing about "implementation constraints" that fix bugs in certain situations being related to functional or non-functional requirements, but this is really not the question. ]

[ABataev]

In D62199#1517522, @Hahnfeld wrote:

In D62199#1517460, @ABataev wrote:

In D62199#1517437, @hfinkel wrote:

Alexey, I believe that I understand the distinction that you are making. That's not the way the term is used within the LLVM community. The observable behavior of the compiler and its runtime libraries is what matters. The fact that some implementation technique should not be required to make that observable behavior correct is irrelevant in our usage. The introduction of that technique in order to make the observable behavior correct, where it was not previously correct, is a functionality-changing change to the compiler. Thus, it is not NFC.

It is better to express this explicitly in the policy, because it differs from the classical definition.

See https://llvm.org/docs/Lexicon.html#nfc:

“No functional change”. Used in a commit message to indicate that a patch is a pure refactoring/cleanup. Usually used in the first line, so it is visible without opening the actual commit email.

This is my last comment. I'm not going yo continue this discussion anymore.
Almost all my patches marked as NFC are refactoring patches. So, the patches were correctly marked as NFC per the definition in the lexicon. Thanks for the reference, Jonas.
If you don't think that "volatile" is a refactoring, I'm sorry. Anyway, the patch were reverted and split into 3 parts.

AFAICS that has nothing to do with "non-functional" requirements. [ We could continue arguing about "implementation constraints" that fix bugs in certain situations being related to functional or non-functional requirements, but this is really not the question. ]

Sorry, you're wrong. The programming is the work with requirements expressed explicitly or implicitly. and all changes in code are to implement those requirements or to fix bugs/refactor the code in accordance with the requirements. So, the changes are also functional and non-functional. If the patch explicitly changes the logical dependency between inputs and outputs, this is a functional change. Otherwise, the change is non-functional. It may be a refactoring or a fix for non-functional requirement/constraint.

openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omp_data.cu

__device__
     omptarget_nvptx_Queue<omptarget_nvptx_ThreadPrivateContext, OMP_STATE_COUNT>
         omptarget_nvptx_device_State[MAX_SM];
 
 __device__ omptarget_nvptx_SimpleMemoryManager
     omptarget_nvptx_simpleMemoryManager;
 __device__ __shared__ uint32_t usedMemIdx;
 __device__ __shared__ uint32_t usedSlotIdx;
 
-__device__ __shared__ uint8_t parallelLevel[MAX_THREADS_PER_TEAM / WARPSIZE];
+__device__ __shared__ volatile uint8_t
+    parallelLevel[MAX_THREADS_PER_TEAM / WARPSIZE];
 __device__ __shared__ uint16_t threadLimit;
 __device__ __shared__ uint16_t threadsInTeam;
 __device__ __shared__ uint16_t nThreads;
 // Pointer to this team's OpenMP state object
 __device__ __shared__
     omptarget_nvptx_ThreadPrivateContext *omptarget_nvptx_threadPrivateContext;
 
 ////////////////////////////////////////////////////////////////////////////////

openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h

 ////////////////////////////////////////////////////////////////////////////////
 // global data tables
 ////////////////////////////////////////////////////////////////////////////////
 
 extern __device__ omptarget_nvptx_SimpleMemoryManager
     omptarget_nvptx_simpleMemoryManager;
 extern __device__ __shared__ uint32_t usedMemIdx;
 extern __device__ __shared__ uint32_t usedSlotIdx;
-extern __device__ __shared__ uint8_t
+extern __device__ __shared__ volatile uint8_t
     parallelLevel[MAX_THREADS_PER_TEAM / WARPSIZE];
 extern __device__ __shared__ uint16_t threadLimit;
 extern __device__ __shared__ uint16_t threadsInTeam;
 extern __device__ __shared__ uint16_t nThreads;
 extern __device__ __shared__
     omptarget_nvptx_ThreadPrivateContext *omptarget_nvptx_threadPrivateContext;
 
 extern __device__ __shared__ uint32_t execution_param;

openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu

if (numberOfActiveOMPThreads > 1) {
             WARPSIZE * ((numberOfActiveOMPThreads + WARPSIZE - 1) / WARPSIZE);
 
         PRINT(LD_SYNC,
               "call kmpc_barrier with %d omp threads, sync parameter %d\n",
               (int)numberOfActiveOMPThreads, (int)threads);
         // Barrier #1 is for synchronization among active threads.
         named_sync(L1_BARRIER, threads);
       }
+    } else {
+      __kmpc_flush(loc_ref);
     } // numberOfActiveOMPThreads > 1
     PRINT0(LD_SYNC, "completed kmpc_barrier\n");
   }
 }
 
 // Emit a simple barrier call in SPMD mode.  Assumes the caller is in an L0
 // parallel region and that all worker threads participate.
 EXTERN void __kmpc_barrier_simple_spmd(kmp_Ident *loc_ref, int32_t tid) {

 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // Flush
 ////////////////////////////////////////////////////////////////////////////////
 
 EXTERN void __kmpc_flush(kmp_Ident *loc) {
   PRINT0(LD_IO, "call kmpc_flush\n");
-  __threadfence_system();
+  __threadfence();
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // Vote
 ////////////////////////////////////////////////////////////////////////////////
 
 EXTERN int32_t __kmpc_warp_active_thread_mask() {
   PRINT0(LD_IO, "call __kmpc_warp_active_thread_mask\n");
   return __ACTIVEMASK();
 }