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

[via cfe-commits]

Author: Artem Belevich
Date: 2024-08-12T10:21:23-07:00
New Revision: 5629249575f56f6135fb63e2f0d4ca9a7375167c

URL: https://github.com/llvm/llvm-project/commit/5629249575f56f6135fb63e2f0d4ca9a7375167c
DIFF: https://github.com/llvm/llvm-project/commit/5629249575f56f6135fb63e2f0d4ca9a7375167c.diff

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

Fixes https://github.com/llvm/llvm-project/issues/101560

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

Modified: 
    clang/include/clang/AST/ASTContext.h

Removed: 
    


################################################################################
diff  --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h
index 94cf9113cd43a4..58a820508da42b 100644
--- a/clang/include/clang/AST/ASTContext.h
+++ 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.

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..1e52887b894b17
--- /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_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;
+    }
+}