================ @@ -0,0 +1,16 @@ +// Tests that during incremental HIP device compilation an inline __device__ +// helper referenced by a kernel has its definition emitted into the device +// module (rather than left as an external declaration). In incremental mode +// device code objects are not linked together, so a referenced inline/internal +// __device__ function must be re-emitted into the referencing module. + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -fincremental-extensions -emit-llvm -o - %s | FileCheck %s ---------------- yxsamliu wrote:
The current test compiles `test_device` and `test_kernel` together with one `%clang_cc1` command, so both are generated in the same module. It does not create two incremental inputs and therefore does not test the problem described by this PR. The existing incremental code remembers emitted internal, linkonce, and weak functions. `Release()` puts them back into `DeferredDecls`, and `moveLazyEmissionStates()` moves them into the next module. A later reference normally finds the function there and emits its definition. The new fallback only runs when the function is missing from `DeferredDecls`, but its definition is still available through `FD->getDefinition()`. The simple helper and kernel inputs in the current test do not seem to trigger this fallback, so a more specific input is needed. This could be tested in `clang/unittests/Interpreter/IncrementalProcessingTest.cpp`. `EmitCXXGlobalInitFunc` shows how to call `Interpreter::Parse()` twice and inspect `PartialTranslationUnit::TheModule`. A no-op executor can follow `InterpreterExtensionsTest.CustomIncrementalExecutor`, since AMDGPU cannot be JIT-executed. The test should demonstrate a case that fails before the change and passes after it. https://github.com/llvm/llvm-project/pull/218784 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
