Introduce KOKKOS_IMPL_ARCH_NVIDIA_GPU macro

[dalg24]

We already have KOKKOS_ARCH_INTEL_GPU. It is useful for backends like SYCL, OpenMPTarget, or OpenACC, that may target GPUs from different vendors.
I am not sure whether we want it to be a public or a private macro. We could consider defining it to be the compute capability as an integer but that means more complicate logic in the generated makefiles and in CMake.

Was proposed at last developer meeting when discussing #5867

[masterleinad]

Looks OK to me.

[masterleinad]

Not quite sure why GitHub CI is not running.

[dalg24]

Not quite sure why GitHub CI is not running.

Retriggering CI. Didn't change anything, just amended nothing to the last commit.

[masterleinad]

Only one of the HIP CI runs is failing with the usual errors.

[crtrott]

I think we really should consider making this thing have the CC as a value, but not blocking for now.

[dalg24]

I think we really should consider making this thing have the CC as a value, but not blocking for now.

I implemented your suggestion. Setting the macro in the code rather than with CMake as you wanted.

[PhilMiller]

Should we initially introduce this as KOKKOS_IMPL_ARCH_NVIDIA_GPU to not commit to exposing it?

[dalg24]

Should we initially introduce this as KOKKOS_IMPL_ARCH_NVIDIA_GPU to not commit to exposing it?

I brought this up at the developer meeting last week and it did not seem to be a concern. I do not feel strongly either way.

[PhilMiller]

Sorry, I missed that. It's fine by me as is. I just wanted to make sure it was considered

[dalg24]

Per dev meeting, make it an IMPL macro

[dalg24]

Retest this please

[dalg24]

Clang+CUDA build failed with

5: [ RUN      ] cuda_DeathTest.view_memory_access_violations_from_device
5: /var/jenkins/workspace/Kokkos/core/unit_test/TestViewMemoryAccessViolation.hpp:50: Failure
5: Death test: { Kokkos::parallel_for(Kokkos::RangePolicy<ExecutionSpace>(s, 0, 1), *this); Kokkos::fence(); }
5:     Result: died but not with expected error.
5:   Expected: contains regular expression "Kokkos::View ERROR: attempt to access inaccessible memory space.*UNAVAILABLE"
5: Actual msg:
5: [  DEATH   ] cudaDeviceSynchronize() error( cudaErrorAssert): device-side assert triggered /var/jenkins/workspace/Kokkos/core/src/Cuda/Kokkos_Cuda_Instance.cpp:135
5: [  DEATH   ] Backtrace:
5: [  DEATH   ]                        [0x1220b90]
5: [  DEATH   ]                        [0x1215178]
5: [  DEATH   ]                        [0x12151f8]
5: [  DEATH   ]                        [0x1228c8f]
5: [  DEATH   ]                        [0x1228593]
5: [  DEATH   ]                        [0x11fc97d]
5: [  DEATH   ]                         [0x982e5b]
5: [  DEATH   ]                         [0x982481]
5: [  DEATH   ]                         [0x973c1d]
5: [  DEATH   ]                        [0x11ea1a4]
5: [  DEATH   ]                        [0x11c2238]
5: [  DEATH   ]                        [0x11c35d0]
5: [  DEATH   ]                        [0x11c3e9b]
5: [  DEATH   ]                        [0x11d3cce]
5: [  DEATH   ]                        [0x11ead74]
5: [  DEATH   ]                        [0x11d36ec]
5: [  DEATH   ]                         [0x416980]
5: [  DEATH   ] __libc_start_main [0x7f588b3f3bf7]
5: [  DEATH   ]                         [0x41688a]
5: [  DEATH   ] 
5: [  FAILED  ] cuda_DeathTest.view_memory_access_violations_from_device (11253 ms)
5: [----------] 7 tests from cuda_DeathTest (123097 ms total)

Looks unrelated but rerunning to be safe.
Retest this please.

[masterleinad]

Looks OK to me (apart from Bruno's questions).