For AMDGPU the visibility of these symbols (OpenCL kernels, CUDA __global__ functions, and CUDA __device__ variables) must not be hidden, as we rely on them being available in the dynamic symbol table in the final DSO.
This patch implements this by considering language attributes as a source of explicit visibility, but rather than attributing any one visibility to them they are simply coerced to be a non-hidden visibility. This allows for the optimization of using protected visibility when these symbols are known to be dso_local.
This patch also adds diagnostics for explicitly setting a hidden visibility on these symbols.
I imagine there are a number of issues with the patch in its current state, but I wanted to get something implemented before reaching out to OpenCL/CUDA maintainers to see if this is a reasonable change. @Anastasia and @tra I wasn't certain if you would be good candidates to discuss this change, so please let me know if I need to keep looking.