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