[CUDA] Emit used function list in deterministic order. (#102661)

Fixes #101560

Author: Artem-B

clang/include/clang/AST/ASTContext.h

@@ -34,6 +34,7 @@
 #include "llvm/ADT/MapVector.h"
 #include "llvm/ADT/PointerIntPair.h"
 #include "llvm/ADT/PointerUnion.h"
+#include "llvm/ADT/SetVector.h"
 #include "llvm/ADT/SmallVector.h"
 #include "llvm/ADT/StringMap.h"
 #include "llvm/ADT/StringRef.h"
@@ -1194,8 +1195,8 @@ class ASTContext : public RefCountedBase<ASTContext> {
   llvm::DenseSet<const VarDecl *> CUDADeviceVarODRUsedByHost;
 
   /// Keep track of CUDA/HIP external kernels or device variables ODR-used by
-  /// host code.
-  llvm::DenseSet<const ValueDecl *> CUDAExternalDeviceDeclODRUsedByHost;
+  /// host code. SetVector is used to maintain the order.
+  llvm::SetVector<const ValueDecl *> CUDAExternalDeviceDeclODRUsedByHost;
 
   /// Keep track of CUDA/HIP implicit host device functions used on device side
   /// in device compilation.

clang/test/CodeGenCUDA/host-used-extern-determinism.cu

@@ -0,0 +1,21 @@
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
+// RUN:   -fgpu-rdc -std=c++11 -emit-llvm -o - -target-cpu gfx906 | FileCheck %s
+
+#include "Inputs/cuda.h"
+
+// CHECK-LABEL: @__clang_gpu_used_external = internal {{.*}}global
+// References to the kernels must be in order of appearance.
+// CHECK-SAME: [ptr @_Z6kernelILi3EEvPi, ptr @_Z6kernelILi1EEvPi, ptr @_Z6kernelILi2EEvPi, ptr @_Z6kernelILi0EEvPi]
+
+template <int N>
+__global__ void kernel(int* out) { *out = N; }
+
+void host(int n) {
+    void * k;
+    switch (n) {
+        case 3: k = (void*)&kernel<3>; break;
+        case 1: k = (void*)&kernel<1>; break;
+        case 2: k = (void*)&kernel<2>; break;
+        case 0: k = (void*)&kernel<0>; break;
+    }
+}