jhuber6 updated this revision to Diff 432146.
jhuber6 added a comment.
Add test for #line.
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D125904/new/
https://reviews.llvm.org/D125904
Files:
clang/lib/CodeGen/CGCUDANV.cpp
clang/lib/CodeGen/CodeGenModule.cpp
clang/lib/CodeGen/CodeGenModule.h
clang/test/CodeGenCUDA/device-fun-linkage.cu
clang/test/CodeGenCUDA/static-device-var-rdc.cu
Index: clang/test/CodeGenCUDA/static-device-var-rdc.cu
===================================================================
--- clang/test/CodeGenCUDA/static-device-var-rdc.cu
+++ clang/test/CodeGenCUDA/static-device-var-rdc.cu
@@ -2,12 +2,12 @@
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa -fcuda-is-device \
-// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s | FileCheck \
-// RUN: -check-prefixes=DEV,INT-DEV %s
+// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o %t.nocuid.dev -x hip %s
+// RUN: cat %t.nocuid.dev | FileCheck -check-prefixes=DEV,INT-DEV %s
// RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-gnu-linux \
-// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s | FileCheck \
-// RUN: -check-prefixes=HOST,INT-HOST %s
+// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o %t.nocuid.host -x hip %s
+// RUN: cat %t.nocuid.host | FileCheck -check-prefixes=HOST,INT-HOST %s
// RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa -fcuda-is-device -cuid=abc \
// RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s > %t.dev
@@ -21,6 +21,7 @@
// variable names.
// RUN: cat %t.dev %t.host | FileCheck -check-prefix=POSTFIX %s
+// RUN: cat %t.nocuid.dev %t.nocuid.host | FileCheck -check-prefix=POSTFIX-ID %s
// Negative tests.
@@ -48,6 +49,9 @@
#include "Inputs/cuda.h"
+// Make sure we can still mangle with a line directive.
+#line 0 "-"
+
// Test function scope static device variable, which should not be externalized.
// DEV-DAG: @_ZZ6kernelPiPPKiE1w = internal addrspace(4) constant i32 1
@@ -56,8 +60,8 @@
// HOST-DAG: @_ZL1y = internal global i32 undef
// Test normal static device variables
-// INT-DEV-DAG: @_ZL1x = addrspace(1) externally_initialized global i32 0
-// INT-HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x\00"
+// INT-DEV-DAG: @_ZL1x[[FILEID:.*]] = addrspace(1) externally_initialized global i32 0
+// INT-HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x[[FILEID:.*]]\00"
// Test externalized static device variables
// EXT-DEV-DAG: @_ZL1x.static.[[HASH:.*]] = addrspace(1) externally_initialized global i32 0
@@ -66,6 +70,8 @@
// POSTFIX: @_ZL1x.static.[[HASH:.*]] = addrspace(1) externally_initialized global i32 0
// POSTFIX: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x.static.[[HASH]]\00"
+// POSTFIX-ID: @_ZL1x.static.[[FILEID:.*]] = addrspace(1) externally_initialized global i32 0
+// POSTFIX-ID: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x.static.[[FILEID]]\00"
static __device__ int x;
@@ -75,8 +81,8 @@
static __device__ int x2;
// Test normal static device variables
-// INT-DEV-DAG: @_ZL1y = addrspace(4) externally_initialized global i32 0
-// INT-HOST-DAG: @[[DEVNAMEY:[0-9]+]] = {{.*}}c"_ZL1y\00"
+// INT-DEV-DAG: @_ZL1y[[FILEID:.*]] = addrspace(4) externally_initialized global i32 0
+// INT-HOST-DAG: @[[DEVNAMEY:[0-9]+]] = {{.*}}c"_ZL1y[[FILEID:.*]]\00"
// Test externalized static device variables
// EXT-DEV-DAG: @_ZL1y.static.[[HASH]] = addrspace(4) externally_initialized global i32 0
Index: clang/test/CodeGenCUDA/device-fun-linkage.cu
===================================================================
--- clang/test/CodeGenCUDA/device-fun-linkage.cu
+++ clang/test/CodeGenCUDA/device-fun-linkage.cu
@@ -23,10 +23,10 @@
// Ensure that unused static device function is eliminated
static __device__ void static_func() {}
// NORDC-NEG-NOT: define{{.*}} void @_ZL13static_funcv()
-// RDC-NEG-NOT: define{{.*}} void @_ZL13static_funcv()
+// RDC-NEG-NOT: define{{.*}} void @_ZL13static_funcv[[FILEID:.*]]()
// Ensure that kernel function has external or weak_odr
// linkage regardless static specifier
static __global__ void static_kernel() {}
// NORDC: define void @_ZL13static_kernelv()
-// RDC: define weak_odr void @_ZL13static_kernelv()
+// RDC: define weak_odr void @_ZL13static_kernelv[[FILEID:.*]]()
Index: clang/lib/CodeGen/CodeGenModule.h
===================================================================
--- clang/lib/CodeGen/CodeGenModule.h
+++ clang/lib/CodeGen/CodeGenModule.h
@@ -1467,7 +1467,10 @@
bool stopAutoInit();
/// Print the postfix for externalized static variable or kernels for single
- /// source offloading languages CUDA and HIP.
+ /// source offloading languages CUDA and HIP. The unique postfix is created
+ /// using either the CUID argument, or the file's UniqueID and active macros.
+ /// The fallback method without a CUID requires that the offloading toolchain
+ /// does not define separate macros via the -cc1 options.
void printPostfixForExternalizedDecl(llvm::raw_ostream &OS,
const Decl *D) const;
Index: clang/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenModule.cpp
+++ clang/lib/CodeGen/CodeGenModule.cpp
@@ -1416,8 +1416,9 @@
// Make unique name for device side static file-scope variable for HIP.
if (CGM.getContext().shouldExternalize(ND) &&
CGM.getLangOpts().GPURelocatableDeviceCode &&
- CGM.getLangOpts().CUDAIsDevice && !CGM.getLangOpts().CUID.empty())
+ CGM.getLangOpts().CUDAIsDevice)
CGM.printPostfixForExternalizedDecl(Out, ND);
+
return std::string(Out.str());
}
@@ -6825,12 +6826,38 @@
void CodeGenModule::printPostfixForExternalizedDecl(llvm::raw_ostream &OS,
const Decl *D) const {
- StringRef Tag;
// ptxas does not allow '.' in symbol names. On the other hand, HIP prefers
// postfix beginning with '.' since the symbol name can be demangled.
if (LangOpts.HIP)
- Tag = (isa<VarDecl>(D) ? ".static." : ".intern.");
+ OS << (isa<VarDecl>(D) ? ".static." : ".intern.");
else
- Tag = (isa<VarDecl>(D) ? "__static__" : "__intern__");
- OS << Tag << getContext().getCUIDHash();
+ OS << (isa<VarDecl>(D) ? "__static__" : "__intern__");
+
+ // If the CUID is not specified we try to generate a unique postfix.
+ if (getLangOpts().CUID.empty()) {
+ SourceManager &SM = getContext().getSourceManager();
+ PresumedLoc PLoc = SM.getPresumedLoc(D->getLocation());
+ assert(PLoc.isValid() && "Source location is expected to be valid.");
+
+ // Get the hash of the user defined macros.
+ llvm::MD5 Hash;
+ llvm::MD5::MD5Result Result;
+ for (const auto &Arg : PreprocessorOpts.Macros)
+ Hash.update(Arg.first);
+ Hash.final(Result);
+
+ // Get the UniqueID for the file containing the decl.
+ llvm::sys::fs::UniqueID ID;
+ if (auto EC = llvm::sys::fs::getUniqueID(PLoc.getFilename(), ID)) {
+ PLoc = SM.getPresumedLoc(D->getLocation(), /*UseLineDirectives=*/false);
+ assert(PLoc.isValid() && "Source location is expected to be valid.");
+ if (auto EC = llvm::sys::fs::getUniqueID(PLoc.getFilename(), ID))
+ SM.getDiagnostics().Report(diag::err_cannot_open_file)
+ << PLoc.getFilename() << EC.message();
+ }
+ OS << llvm::format("%x", ID.getFile()) << llvm::format("%x", ID.getDevice())
+ << "_" << llvm::utohexstr(Result.low(), /*LowerCase=*/true, /*Width=*/8);
+ } else {
+ OS << getContext().getCUIDHash();
+ }
}
Index: clang/lib/CodeGen/CGCUDANV.cpp
===================================================================
--- clang/lib/CodeGen/CGCUDANV.cpp
+++ clang/lib/CodeGen/CGCUDANV.cpp
@@ -285,8 +285,7 @@
// Make unique name for device side static file-scope variable for HIP.
if (CGM.getContext().shouldExternalize(ND) &&
- CGM.getLangOpts().GPURelocatableDeviceCode &&
- !CGM.getLangOpts().CUID.empty()) {
+ CGM.getLangOpts().GPURelocatableDeviceCode) {
SmallString<256> Buffer;
llvm::raw_svector_ostream Out(Buffer);
Out << DeviceSideName;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
