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; + } +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
