[flang][OpenMP] critical does not serialize lanes in target offload, giving wrong results

A `critical` region inside a target region does not serialize the lanes of a wavefront in flang, so a conforming program silently gets wrong results. The C equivalent passes on the same GPU.

## Reproducer

```fortran
program p
integer :: s, i, tl
character(len=16) :: a
call get_command_argument(1, a); read (a, *) tl
s = 0
!$omp target parallel do num_threads(tl) map(tofrom: s)
do i = 1, tl
!$omp critical
s = s + 1
!$omp end critical
end do
print '(A,I5,A,I6)', "threads=", tl, " critical=", s
end program p
```

```console
$ flang -fopenmp --offload-arch=gfx90a -O2 crit.f90 -o crit && ./crit 256
threads= 256 critical= 4 # expected 256
```

The same shape in C is correct:

```c
#pragma omp target parallel for num_threads(tl) map(tofrom: s)
for (int i = 0; i < tl; ++i) {
#pragma omp critical
{ s += 1; }
}
```

```console
$ clang -fopenmp --offload-arch=gfx90a -O2 crit.c -o critc && ./critc 256
threads= 256 critical= 256
```

## What the number means

The surviving count is the number of wavefronts, not the number of threads:

| threads | wavefronts | flang `critical` | clang `critical` |
|---|---|---|---|
| 64 | 1 | 1 | 64 |
| 128 | 2 | 2 | 128 |
| 256 | 4 | 4 | 256 |

Lanes serialize between wavefronts and never within one.

Three controls, all in the same construct and on the same mapped variable:

* `atomic update` instead of `critical` gives the full count, so the variable is genuinely shared and written back correctly.
* Running the same binary with `OMP_TARGET_OFFLOAD=DISABLED` gives the full count, so the source semantics are as expected.
* Wrapping the `critical` in a hand-written lane-serialization loop (`me = omp_get_thread_num()`, then `do t = 0, tl-1; if (me == t) then ... end if; end do`) gives the full count with the same `critical` construct.

## Cause

`setCriticalLock` in `openmp/device/src/Synchronization.cpp` acquires the lock only on the lowest active lane of a wavefront:

```cpp
void setCriticalLock(omp_lock_t *Lock) {
uint64_t LowestActiveThread = utils::ctz(mapping::activemask());
if (mapping::getThreadIdInWarp() == LowestActiveThread) {
...
}
}
```

That is sound only if exactly one lane per wavefront reaches it. clang guarantees this: `CGOpenMPRuntimeGPU::emitCriticalRegion` wraps the region in a loop over `__kmpc_get_hardware_num_threads_in_block()`, lets only the matching lane enter, and calls `__kmpc_syncwarp` between turns, before delegating to the generic emission for the lock itself.

`OpenMPIRBuilder::createCritical`, which flang reaches through `convertOmpCritical`, has no device path. It emits `__kmpc_critical` / body / `__kmpc_end_critical` directly, so every active lane enters the body at once and all but one update is lost.

Device IR from the Fortran reproducer, with no serialization around the region:

```llvm
tail call void @__kmpc_critical(ptr addrspacecast (ptr addrspace(1) @1 to ptr), i32 %omp_global_thread_num7, ptr nonnull @.gomp_critical_user_.var)
%3 = load i32, ptr %loadgep_2, align 4
%4 = add i32 %3, 1
store i32 %4, ptr %loadgep_2, align 4
tail call void @__kmpc_end_critical(ptr addrspacecast (ptr addrspace(1) @1 to ptr), i32 %omp_global_thread_num7, ptr nonnull @.gomp_critical_user_.var)
```

## Environment

flang and clang 24.0.0git, `--offload-arch=gfx90a`, MI250X. Also reproduces with AMD's downstream amdflang on AFAR 22.2.0, 23.1.0, 23.2.0 and 23.2.1, and the C case is correct on all of them.

The root cause analysis and the reduced test cases in this report were produced with Claude; I reviewed and verified them.

Start by reproducing the Fortran case with flang and then read openmp/device/src/Synchronization.cpp, especially setCriticalLock. Trace flang's convertOmpCritical into OpenMPIRBuilder::createCritical and compare it with clang's GPU critical-region handling described in the issue. Done means the offloaded reproducer returns the full count for multiple thread and wavefront sizes without regressing the existing C behavior.

Written by the indexing model from the issue text.