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

[Artem Belevich via cfe-commits]

https://github.com/Artem-B updated https://github.com/llvm/llvm-project/pull/102661

>From 0f3944e1c12baa958f52c3c015a0cf5f9aeff1ed Mon Sep 17 00:00:00 2001
From: Artem Belevich <tra at google.com>
Date: Fri, 9 Aug 2024 11:51:23 -0700
Subject: [PATCH] [CUDA] Emit used function list in deterministic order.

Fixes https://github.com/llvm/llvm-project/issues/101560
---
 clang/lib/CodeGen/CodeGenModule.cpp           |  3 +++
 .../host-used-extern-determinism.cu           | 21 +++++++++++++++++++
 2 files changed, 24 insertions(+)
 create mode 100644 clang/test/CodeGenCUDA/host-used-extern-determinism.cu

diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index 9aaf90ccfe04ff..aefedeffab614a 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -950,6 +950,9 @@ void CodeGenModule::Release() {
       UsedArray.push_back(llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast(
           GetAddrOfGlobal(GD), Int8PtrTy));
     }
+    // Sort decls by name to always emit them in deterministic order.
+    llvm::sort(UsedArray,
+               [](auto A, auto B) { return A->getName() < B->getName(); });
 
     llvm::ArrayType *ATy = llvm::ArrayType::get(Int8PtrTy, UsedArray.size());
 
diff --git a/clang/test/CodeGenCUDA/host-used-extern-determinism.cu b/clang/test/CodeGenCUDA/host-used-extern-determinism.cu
new file mode 100644
index 00000000000000..0c4e966f5de6f8
--- /dev/null
+++ b/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_externalx = internal {{.*}}global
+// References to the kernels must be in sorted order.
+// CHECK-SAME: [ptr @_Z6kernelILi0EEvPi, ptr @_Z6kernelILi1EEvPi, ptr @_Z6kernelILi2EEvPi, ptr @_Z6kernelILi3EEvPi]
+
+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;
+    }
+}