llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Levon Ter-Grigoryan (PatriosTheGreat) <details> <summary>Changes</summary> Promoting __constant__ to external variables includes them to PTX which then leads to nvlinker failure. See changes at device-use-host-var test. Befor this change those variables was included to PTX without definition. --- Full diff: https://github.com/llvm/llvm-project/pull/73549.diff 2 Files Affected: - (modified) clang/lib/Sema/SemaCUDA.cpp (+1) - (modified) clang/test/CodeGenCUDA/device-use-host-var.cu (+16) ``````````diff diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index 318174f7be8fa95..f9d72e571e7b98b 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -783,6 +783,7 @@ void Sema::MaybeAddCUDAConstantAttr(VarDecl *VD) { (VD->isFileVarDecl() || VD->isStaticDataMember()) && !IsDependentVar(VD) && ((VD->isConstexpr() || VD->getType().isConstQualified()) && + VD->getStorageClass() != SC_Extern && HasAllowedCUDADeviceStaticInitializer(*this, VD, CICK_DeviceOrConstant))) { VD->addAttr(CUDAConstantAttr::CreateImplicit(getASTContext())); diff --git a/clang/test/CodeGenCUDA/device-use-host-var.cu b/clang/test/CodeGenCUDA/device-use-host-var.cu index 64de57e41b4b9f5..807a485f4c14972 100644 --- a/clang/test/CodeGenCUDA/device-use-host-var.cu +++ b/clang/test/CodeGenCUDA/device-use-host-var.cu @@ -2,6 +2,8 @@ // RUN: -fcuda-is-device -emit-llvm -o - -x hip %s | FileCheck %s // RUN: %clang_cc1 -std=c++14 -triple amdgcn-amd-amdhsa \ // RUN: -fcuda-is-device -emit-llvm -o - -x hip %s | FileCheck -check-prefix=NEG %s +// RUN: %clang_cc1 -std=c++14 -triple nvptx64-nvidia-cuda \ +// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s | FileCheck -check-prefix=NEG %s #include "Inputs/cuda.h" @@ -104,3 +106,17 @@ void fun() { (void) b<double>; (void) var_host_only; } + +// NEG-NOT: external_func +extern __global__ void external_func(); +// NEG-NOT: @external_dep +extern void* const external_dep[] = { + (void*)(external_func) +}; +// NEG-NOT: @external_arr +extern void* const external_arr[] = {}; + +void* host_fun() { + (void) external_dep; + (void) external_arr; +} `````````` </details> https://github.com/llvm/llvm-project/pull/73549 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
