[AMDGPU] `-munsafe-fp-atomics` does nothing on gfx908

_This issue is closely related to https://github.com/intel/llvm/issues/7179 but I'm opening another ticket as that is related to gfx90a, which looks like to not have the same problem_

**Describe the bug**
***Contex***
On MI100, FP atomic support is not complete, and by default FP operations are implemented as a CAS loop. However, on gfx908 ROCm exposes an alternative function `atomiAddNoRet` which returns `void`, that is controlled via the `-munsafe-fp-atomics` flag. The hip compiler should detect if the return value of the atomic_add is used and, if not, issue the optimized function. When used, it generates the `global_atomic_add_f32` builtin function, avoiding the CAS loop.
***Bug***
However, when compiling with DPC++ the flag is ignored and the compiler always produces a very inefficient CAS loop, even it would be possible to use the optimized function.
The same code on AdaptiveCpp with the aforementioned flag correctly produce the native builtin and runs 20x faster.

**To Reproduce**
Code is taken from https://github.com/intel/llvm/issues/7179
```c++
#include

#include
#include

#define atomicAdd(x, y) (sycl::atomic_ref(*(x)) +=(y))

const int factor = 4;
const int num_points = 10240 * 1 * 256;
const int num_wgs = 10240 * factor;
const int wg_size = 256 / factor;

const int tm_size = 32;
const int num_iters = 1000;

// add 1.0f to array element if it is in my 'team'
void kernel(float* data, sycl::nd_item<1> item)
{
size_t tid = item.get_global_linear_id();
size_t sid = tid / tm_size;

for (size_t i = 0; i < tm_size; i++) {
atomicAdd(data + (sid * tm_size + i), 1.0f);
}
}

int main()
{
try {
sycl::device dev = sycl::device(sycl::default_selector_v);
sycl::queue q(dev);

float* data_d = sycl::malloc_device(num_points, q);
float* data_h = sycl::malloc_host( num_points, q);

// warm-up run:
for (int iter = 0; iter < 1; iter++) {
q.memset(data_d, 0.0, num_points * sizeof(float));
q.wait();

q.parallel_for(
sycl::nd_range<1>(num_points, wg_size),
[=](sycl::nd_item<1> item) {
kernel(data_d, item);
}
);
q.wait_and_throw();
}

// timed run:
std::chrono::time_point time11;
std::chrono::time_point time12;
double duration = 0.0;

for (int iter = 0; iter < num_iters; iter++) {
q.memset(data_d, 0.0, num_points * sizeof(float));
q.wait();

time11 = std::chrono::steady_clock::now();
q.parallel_for(
sycl::nd_range<1>(num_points, wg_size),
[=](sycl::nd_item<1> item) {
kernel(data_d, item);
}
);
q.wait_and_throw();
time12 = std::chrono::steady_clock::now();
duration += std::chrono::duration(time12 - time11).count();
}

q.memcpy(data_h, data_d, num_points * sizeof(float));
q.wait();

for (size_t i = 0; i < num_points; i++) {
if (data_h[i] != (float)tm_size) {
std::cerr << "mismatch found at index " << i << std::endl;
return 3;
}
}
std::cout << "Success!\n";
std::cout << "num_points: " << num_points << std::endl;
std::cout << "num_wgs: " << num_wgs << std::endl;
std::cout << "wg_size: " << wg_size << std::endl;
std::cout << "tm_size: " << tm_size << std::endl;
std::cout << "Total time for " << num_iters << " iterations: " << duration << " us\n";

} catch (sycl::exception const &exc) {
std::cerr << exc.what() << "Exception caught at file:" << __FILE__
<< ", line:" << __LINE__ << std::endl;
return 1;
} catch (std::exception& e) {
std::cerr << "Uncaught exception: " << e.what() << std::endl;
return 2;
}

return 0;
}
```

Generated code for DPC++
```
v_add_f32_e32 v2, 1.0, v3
global_atomic_cmpswap v2, v[0:1], v[2:3], off glc
s_waitcnt vmcnt(0)
v_cmp_eq_u32_e32 vcc, v2, v3
s_or_b64 s[4:5], vcc, s[4:5]
v_mov_b32_e32 v3, v2
s_andn2_b64 exec, exec, s[4:5]
s_cbranch_execnz .LBB1_63
```
AdaptiveCpp:
```
global_atomic_add_f32 v[2:3], v4, off
global_atomic_add_f32 v[2:3], v4, off offset:4
global_atomic_add_f32 v[2:3], v4, off offset:8
global_atomic_add_f32 v[2:3], v4, off offset:12
global_atomic_add_f32 v[2:3], v4, off offset:16
global_atomic_add_f32 v[2:3], v4, off offset:20
global_atomic_add_f32 v[2:3], v4, off offset:24
global_atomic_add_f32 v[2:3], v4, off offset:28
global_atomic_add_f32 v[2:3], v4, off offset:32
```
DPC++ also warns that the flag is not used:
```bash
clang++: warning: argument unused during compilation: '-munsafe-fp-atomics' [-Wunused-command-line-argument]
```

[Here](https://github.com/intel/llvm/blob/0e8ce8c7fa53c5c55cbf587bb985ff66434c6d6e/libclc/amdgcn-amdhsa/libspirv/atomic/atomic_add.cl) looks like the fetch add is lowered to `__hip_atomic_fetch_add`. I tried a simple hip kernel using that builtin with the unsafe flag, and hipcc also generated the `global_atomic_add_f32` function.

**Environment**

- OS: Ubuntu 20.04
- DPC++ version: 0e8ce8c7fa53c5c55cbf587bb985ff66434c6d6e
- ROCm version: 5.5.0

**Additional Context**
https://github.com/AdaptiveCpp/AdaptiveCpp/issues/729
https://github.com/ROCm-Developer-Tools/hipamd/issues/19