Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
58 commits
Select commit Hold shift + click to select a range
7aeb264
[ASan] Make most tests run under internal shell on Darwin
boomanaiden154 Nov 18, 2025
38c1a58
[flang][NFC] Strip trailing whitespace from tests (6 of N)
tarunprabhu Nov 18, 2025
67d5c14
[llvm][AddressSanitizer] option for applying AddressSanitizer to spec…
etsal Nov 18, 2025
1fb8e3d
[lldb] Support integer registers with more than 64 bits. (#166363)
sedymrak Nov 18, 2025
93a8ca8
[AArch64][GISel] Don't crash in known-bits when copying from vectors …
cofibrant Nov 18, 2025
2675dcd
[lldb] update lldb-server platform help parsing (attempt 3) (#164904)
cs01 Nov 18, 2025
c7d2ed4
Reland [Support][Jobserver][Tests] Simplify default executor init (#1…
yxsamliu Nov 18, 2025
727ee7e
[APInt] Introduce carry-less multiply primitives (#168527)
artagnon Nov 18, 2025
cb58129
[AMDGPU][GlobalISel] Add RegBankLegalize support for G_IS_FPCLASS (#1…
vangthao95 Nov 18, 2025
6d3971d
[AsmParser] Use a range-based for loop (NFC) (#168488)
kazutakahirata Nov 18, 2025
58cffea
[InstCombine] Canonicalize signed saturated additions (#153053)
AZero13 Nov 18, 2025
b533712
[CI] Only run normal check targets if requested (#168412)
boomanaiden154 Nov 18, 2025
94e9bfb
[AArch64] Reorder Comparison Trees to Facilitate CSE (#168064)
mskamp Nov 18, 2025
3cf1f0c
[ARM] Pattern match Low Overhead Loops pseudos (NFC) (#168209)
s-barannikov Nov 18, 2025
0b82415
[AMDGPU] Consider FLAT instructions for VMEM hazard detection (#137170)
ro-i Nov 18, 2025
c88ae6e
[flang][OpenMP] Move two utilities from Semantics to Parser, NFC (#16…
kparzysz Nov 18, 2025
bd8c941
Reapply "[Github] Update PR labeller to v6.0.1 (#167246)"
boomanaiden154 Nov 18, 2025
d772663
Revert "[Github] Update PR labeller to v6.0.1 (#167246)"
boomanaiden154 Nov 18, 2025
5af0398
[X86] Add test examples of build vectors of reversed scalar loads tha…
RKSimon Nov 18, 2025
5407e62
Revert "[MLIR][NVVM] Add tcgen05.mma MLIR Ops" (#168583)
joker-eph Nov 18, 2025
8bdd82c
[CI] Skip Running Premerge Advisor on AArch64 (#168404)
boomanaiden154 Nov 18, 2025
40ed57c
[CI] Prefer Bash Tests over Empty String Comparisons (#168575)
boomanaiden154 Nov 18, 2025
0ae2bcc
[ARM] TableGen-erate node descriptions (#168212)
s-barannikov Nov 18, 2025
523bd2d
[GISel][RISCV] Compute CTPOP of small odd-sized integer correctly (#1…
XChy Nov 18, 2025
46565f3
[LLDB] Add log channel for InstrumentationRuntime plugins (#168508)
delcypher Nov 18, 2025
3f61402
[clang][DependencyScanning] Add Test Coverage of `StabeDirs` during B…
qiongsiwu Nov 18, 2025
8f67759
[NFC][TableGen] Remove `close` member from various CodeGenHelpers (#1…
jurahul Nov 18, 2025
4ab2423
[ConstantFolding] Generalize constant folding for vector_interleave2 …
topperc Nov 18, 2025
96e58b8
[RISCV] Legalize misaligned unmasked vp.load/vp.store to vle8/vse8. (…
topperc Nov 18, 2025
0dd3cb5
Reland instr-ref-target-hooks-sp-clobber.mir (#168136)
rastogishubham Nov 18, 2025
b630721
[bazel] Fix #164904 (#168593)
pranavk Nov 18, 2025
e93763e
[dsymutil] Specify that -flat is for testing in the help output (#168…
JDevlieghere Nov 18, 2025
2ad93b4
[X86] getRoundingModeX86 - add missing "clang-format on" toggle comme…
RKSimon Nov 18, 2025
ac6e48d
Modify llvm-dwp to be able to emit string tables over 4GB without los…
clayborg Nov 18, 2025
58b8e6e
[DebugInfo][IR] Verifier checks for the extraData (#167971)
laxmansole Nov 18, 2025
04a1fd5
[RISCV] Make XFAIL test UNSUPPORTED. (#168525)
mgudim Nov 18, 2025
576e1af
[NFC][AMDGPU] IGLP: Fixes for unsigned int handling (#135090)
ro-i Nov 18, 2025
124fa5c
[AArch64] - Improve costing for Identity shuffles for SVE targets. (#…
pawan-nirpal-031 Nov 18, 2025
4155cdc
Mips: Remove manual libcall name search and table (#168595)
arsenm Nov 18, 2025
8aca6c3
[AllocToken] Test compatibility with -fsanitize=kcfi,memtag (#168600)
melver Nov 18, 2025
e1bb50b
[bazel] fix #168212 (#168598)
pranavk Nov 18, 2025
56b1d42
[CIR] Mark globals as constants (#168463)
andykaylor Nov 18, 2025
1157a22
[GISel] Use getScalarSizeInBits in LegalizerHelper::lowerBitCount (#1…
topperc Nov 18, 2025
3e8dc4d
[clang][deps] NFC: Use qualified names for function definitions (#168…
jansvoboda11 Nov 18, 2025
d3c2973
[lldb/aarch64] Add STR/LDR instructions for FP registers to Emulator …
igorkudrin Nov 18, 2025
507f236
[VPlan] Fix OpType-mismatch in getFlagsFromIndDesc (#168560)
artagnon Nov 18, 2025
8fce476
Implement a more seamless way to provide missing functions on z/OS (#…
perry-ca Nov 18, 2025
31ec633
[clang-tidy] Fix bugs in misc-coroutine-hostile-raii check (#167947)
higher-performance Nov 18, 2025
c4898f3
[HLSL][DirectX] Use a padding type for HLSL buffers. (#167404)
bogner Nov 18, 2025
5cde345
[runtimes] Remove pstl from the list of supported runtimes (#168414)
ldionne Nov 18, 2025
1e3ea03
[VPlan] VPIRFlags kind for FCmp with predicate + fast-math flags (NFCI).
fhahn Nov 18, 2025
6665642
[AMDGPU] Don't fold an i64 immediate value if it can't be replicated …
shiltian Nov 18, 2025
db71cc5
[libc] Implement pkey_alloc/free/get/set/mprotect for x86_64 linux (#…
jtstogel Nov 18, 2025
e47e9f3
[NVPTX] TableGen-erate SDNode descriptions (#168367)
s-barannikov Nov 18, 2025
ed78ab7
[orc-rt] Introduce Task and TaskDispatcher APIs and implementations. …
lhames Nov 18, 2025
3e499e9
[CIR] Add support for common linkage (#168613)
andykaylor Nov 18, 2025
5e80358
[llvm][ARM] Allow MOVT and MOVW on the offset between two labels (#16…
hwti Nov 18, 2025
ba0964a
merge main into amd-staging
ronlieb Nov 19, 2025
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
12 changes: 7 additions & 5 deletions .ci/monolithic-linux.sh
Original file line number Diff line number Diff line change
Expand Up @@ -64,11 +64,13 @@ cmake -S "${MONOREPO_ROOT}"/llvm -B "${BUILD_DIR}" \

start-group "ninja"

# Targets are not escaped as they are passed as separate arguments.
ninja -C "${BUILD_DIR}" -k 0 ${targets} |& tee ninja.log
cp ${BUILD_DIR}/.ninja_log ninja.ninja_log
if [[ -n "${targets}" ]]; then
# Targets are not escaped as they are passed as separate arguments.
ninja -C "${BUILD_DIR}" -k 0 ${targets} |& tee ninja.log
cp ${BUILD_DIR}/.ninja_log ninja.ninja_log
fi

if [[ "${runtime_targets}" != "" ]]; then
if [[ -n "${runtime_targets}" ]]; then
start-group "ninja Runtimes"

ninja -C "${BUILD_DIR}" ${runtime_targets} |& tee ninja_runtimes.log
Expand All @@ -77,7 +79,7 @@ fi

# Compiling runtimes with just-built Clang and running their tests
# as an additional testing for Clang.
if [[ "${runtime_targets_needs_reconfig}" != "" ]]; then
if [[ -n "${runtime_targets_needs_reconfig}" ]]; then
start-group "CMake Runtimes C++26"

cmake \
Expand Down
10 changes: 6 additions & 4 deletions .ci/monolithic-windows.sh
Original file line number Diff line number Diff line change
Expand Up @@ -51,11 +51,13 @@ cmake -S "${MONOREPO_ROOT}"/llvm -B "${BUILD_DIR}" \

start-group "ninja"

# Targets are not escaped as they are passed as separate arguments.
ninja -C "${BUILD_DIR}" -k 0 ${targets} |& tee ninja.log
cp ${BUILD_DIR}/.ninja_log ninja.ninja_log
if [[ -n "${targets}" ]]; then
# Targets are not escaped as they are passed as separate arguments.
ninja -C "${BUILD_DIR}" -k 0 ${targets} |& tee ninja.log
cp ${BUILD_DIR}/.ninja_log ninja.ninja_log
fi

if [[ "${runtimes_targets}" != "" ]]; then
if [[ -n "${runtimes_targets}" ]]; then
start-group "ninja runtimes"

ninja -C "${BUILD_DIR}" -k 0 ${runtimes_targets} |& tee ninja_runtimes.log
Expand Down
2 changes: 1 addition & 1 deletion .ci/premerge_advisor_explain.py
Original file line number Diff line number Diff line change
Expand Up @@ -148,7 +148,7 @@ def main(

# Skip looking for results on AArch64 for now because the premerge advisor
# service is not available on AWS currently.
if platform.machine() == "arm64":
if platform.machine() == "arm64" or platform.machine() == "aarch64":
sys.exit(0)

main(
Expand Down
2 changes: 1 addition & 1 deletion .ci/premerge_advisor_upload.py
Original file line number Diff line number Diff line change
Expand Up @@ -59,7 +59,7 @@ def main(commit_sha, workflow_run_number, build_log_files):

# Skip uploading results on AArch64 for now because the premerge advisor
# service is not available on AWS currently.
if platform.machine() == "arm64":
if platform.machine() == "arm64" or platform.machine() == "aarch64":
sys.exit(0)

main(args.commit_sha, args.workflow_run_number, args.build_log_files)
10 changes: 5 additions & 5 deletions .ci/utils.sh
Original file line number Diff line number Diff line change
Expand Up @@ -33,7 +33,7 @@ function at-exit {
# If building fails there will be no results files.
shopt -s nullglob

if [[ "$GITHUB_ACTIONS" != "" ]]; then
if [[ -n "$GITHUB_ACTIONS" ]]; then
python "${MONOREPO_ROOT}"/.ci/generate_test_report_github.py \
$retcode "${BUILD_DIR}"/test-results.*.xml "${MONOREPO_ROOT}"/ninja*.log \
>> $GITHUB_STEP_SUMMARY
Expand All @@ -44,7 +44,7 @@ function at-exit {
fi

if [[ "$retcode" != "0" ]]; then
if [[ "$GITHUB_ACTIONS" != "" ]]; then
if [[ -n "$GITHUB_ACTIONS" ]]; then
python "${MONOREPO_ROOT}"/.ci/premerge_advisor_upload.py \
$(git rev-parse HEAD~1) $GITHUB_RUN_NUMBER \
"${BUILD_DIR}"/test-results.*.xml "${MONOREPO_ROOT}"/ninja*.log
Expand All @@ -59,10 +59,10 @@ trap at-exit EXIT

function start-group {
groupname=$1
if [[ "$GITHUB_ACTIONS" != "" ]]; then
if [[ -n "$GITHUB_ACTIONS" ]]; then
echo "::endgroup"
echo "::group::$groupname"
elif [[ "$POSTCOMMIT_CI" != "" ]]; then
elif [[ -n "$POSTCOMMIT_CI" ]]; then
echo "@@@$STEP@@@"
else
echo "Starting $groupname"
Expand All @@ -73,6 +73,6 @@ export PIP_BREAK_SYSTEM_PACKAGES=1
pip install -q -r "${MONOREPO_ROOT}"/.ci/all_requirements.txt

# The ARM64 builders run on AWS and don't have access to the GCS cache.
if [[ "$GITHUB_ACTIONS" != "" ]] && [[ "$RUNNER_ARCH" != "ARM64" ]]; then
if [[ -n "$GITHUB_ACTIONS" ]] && [[ "$RUNNER_ARCH" != "ARM64" ]]; then
python .ci/cache_lit_timing_files.py download
fi
14 changes: 9 additions & 5 deletions clang-tools-extra/clang-tidy/misc/CoroutineHostileRAIICheck.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -67,6 +67,11 @@ static auto typeWithNameIn(const std::vector<StringRef> &Names) {
hasCanonicalType(hasDeclaration(namedDecl(hasAnyName(Names)))));
}

static auto functionWithNameIn(const std::vector<StringRef> &Names) {
auto Call = callExpr(callee(functionDecl(hasAnyName(Names))));
return anyOf(expr(cxxBindTemporaryExpr(has(Call))), expr(Call));
}

CoroutineHostileRAIICheck::CoroutineHostileRAIICheck(StringRef Name,
ClangTidyContext *Context)
: ClangTidyCheck(Name, Context),
Expand All @@ -83,9 +88,8 @@ void CoroutineHostileRAIICheck::registerMatchers(MatchFinder *Finder) {
hasAttr(attr::Kind::ScopedLockable)))))
.bind("scoped-lockable");
auto OtherRAII = varDecl(typeWithNameIn(RAIITypesList)).bind("raii");
auto AllowedSuspend = awaitable(
anyOf(typeWithNameIn(AllowedAwaitablesList),
callExpr(callee(functionDecl(hasAnyName(AllowedCallees))))));
auto AllowedSuspend = awaitable(anyOf(typeWithNameIn(AllowedAwaitablesList),
functionWithNameIn(AllowedCallees)));
Finder->addMatcher(
expr(anyOf(coawaitExpr(unless(AllowedSuspend)), coyieldExpr()),
forEachPrevStmt(
Expand Down Expand Up @@ -113,9 +117,9 @@ void CoroutineHostileRAIICheck::storeOptions(
ClangTidyOptions::OptionMap &Opts) {
Options.store(Opts, "RAIITypesList",
utils::options::serializeStringList(RAIITypesList));
Options.store(Opts, "SafeAwaitableList",
Options.store(Opts, "AllowedAwaitablesList",
utils::options::serializeStringList(AllowedAwaitablesList));
Options.store(Opts, "SafeCallees",
Options.store(Opts, "AllowedCallees",
utils::options::serializeStringList(AllowedCallees));
}
} // namespace clang::tidy::misc
Original file line number Diff line number Diff line change
Expand Up @@ -2,7 +2,7 @@
// RUN: -config="{CheckOptions: {\
// RUN: misc-coroutine-hostile-raii.RAIITypesList: 'my::Mutex; ::my::other::Mutex', \
// RUN: misc-coroutine-hostile-raii.AllowedAwaitablesList: 'safe::awaitable; ::transformable::awaitable', \
// RUN: misc-coroutine-hostile-raii.AllowedCallees: 'safe::AwaitFunc; ::safe::Obj::AwaitMethod' \
// RUN: misc-coroutine-hostile-raii.AllowedCallees: 'safe::AwaitFunc; ::safe::Obj::AwaitMethod; retExemptedAwaitable' \
// RUN: }}"

namespace std {
Expand Down Expand Up @@ -163,7 +163,10 @@ ReturnObject RAIISafeSuspendTest() {
// ================================================================================
// Safe transformable awaitable
// ================================================================================
struct transformable { struct awaitable{}; };
struct transformable {
struct awaitable{};
struct unsafe_awaitable{};
};
using alias_transformable_awaitable = transformable::awaitable;
struct UseTransformAwaitable {
struct promise_type {
Expand All @@ -172,13 +175,18 @@ struct UseTransformAwaitable {
std::suspend_always final_suspend() noexcept { return {}; }
void unhandled_exception() {}
std::suspend_always await_transform(transformable::awaitable) { return {}; }
std::suspend_always await_transform(transformable::unsafe_awaitable) {
return {};
}
};
};

auto retAwaitable() { return transformable::awaitable{}; }
auto retExemptedAwaitable() { return transformable::unsafe_awaitable{}; }
UseTransformAwaitable RAIISafeSuspendTest2() {
absl::Mutex a;
co_await retAwaitable();
co_await retExemptedAwaitable();
co_await transformable::awaitable{};
co_await alias_transformable_awaitable{};
}
Expand Down
1 change: 0 additions & 1 deletion clang/include/clang/CIR/MissingFeatures.h
Original file line number Diff line number Diff line change
Expand Up @@ -29,7 +29,6 @@ struct MissingFeatures {

// Unhandled global/linkage information.
static bool opGlobalThreadLocal() { return false; }
static bool opGlobalConstant() { return false; }
static bool opGlobalWeakRef() { return false; }
static bool opGlobalUnnamedAddr() { return false; }
static bool opGlobalSection() { return false; }
Expand Down
1 change: 0 additions & 1 deletion clang/lib/AST/ByteCode/Context.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -21,7 +21,6 @@
#include "clang/AST/ASTLambda.h"
#include "clang/AST/Expr.h"
#include "clang/Basic/TargetInfo.h"
#include "llvm/Support/SystemZ/zOSSupport.h"

using namespace clang;
using namespace clang::interp;
Expand Down
3 changes: 2 additions & 1 deletion clang/lib/CIR/CodeGen/CIRGenDecl.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -526,7 +526,8 @@ cir::GlobalOp CIRGenFunction::addInitializerToStaticVarDecl(
bool needsDtor =
d.needsDestruction(getContext()) == QualType::DK_cxx_destructor;

assert(!cir::MissingFeatures::opGlobalConstant());
gv.setConstant(d.getType().isConstantStorage(
getContext(), /*ExcludeCtor=*/true, !needsDtor));
gv.setInitialValueAttr(init);

emitter.finalize(gv);
Expand Down
28 changes: 21 additions & 7 deletions clang/lib/CIR/CodeGen/CIRGenModule.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -675,7 +675,10 @@ CIRGenModule::getOrCreateCIRGlobal(StringRef mangledName, mlir::Type ty,
errorNYI(d->getSourceRange(), "OpenMP target global variable");

gv.setAlignmentAttr(getSize(astContext.getDeclAlign(d)));
assert(!cir::MissingFeatures::opGlobalConstant());
// FIXME: This code is overly simple and should be merged with other global
// handling.
gv.setConstant(d->getType().isConstantStorage(
astContext, /*ExcludeCtor=*/false, /*ExcludeDtor=*/false));

setLinkageForGV(gv, d);

Expand Down Expand Up @@ -864,7 +867,11 @@ void CIRGenModule::emitGlobalVarDefinition(const clang::VarDecl *vd,
if (emitter)
emitter->finalize(gv);

assert(!cir::MissingFeatures::opGlobalConstant());
// If it is safe to mark the global 'constant', do so now.
gv.setConstant((vd->hasAttr<CUDAConstantAttr>() && langOpts.CUDAIsDevice) ||
(!needsGlobalCtor && !needsGlobalDtor &&
vd->getType().isConstantStorage(
astContext, /*ExcludeCtor=*/true, /*ExcludeDtor=*/true)));
assert(!cir::MissingFeatures::opGlobalSection());

// Set CIR's linkage type as appropriate.
Expand All @@ -876,8 +883,17 @@ void CIRGenModule::emitGlobalVarDefinition(const clang::VarDecl *vd,
// FIXME(cir): setLinkage should likely set MLIR's visibility automatically.
gv.setVisibility(getMLIRVisibilityFromCIRLinkage(linkage));
assert(!cir::MissingFeatures::opGlobalDLLImportExport());
if (linkage == cir::GlobalLinkageKind::CommonLinkage)
errorNYI(initExpr->getSourceRange(), "common linkage");
if (linkage == cir::GlobalLinkageKind::CommonLinkage) {
// common vars aren't constant even if declared const.
gv.setConstant(false);
// Tentative definition of global variables may be initialized with
// non-zero null pointers. In this case they should have weak linkage
// since common linkage must have zero initializer and must not have
// explicit section therefore cannot have non-zero initial value.
std::optional<mlir::Attribute> initializer = gv.getInitialValue();
if (initializer && !getBuilder().isNullValue(*initializer))
gv.setLinkage(cir::GlobalLinkageKind::WeakAnyLinkage);
}

setNonAliasAttributes(vd, gv);

Expand Down Expand Up @@ -1231,10 +1247,8 @@ cir::GlobalLinkageKind CIRGenModule::getCIRLinkageForDeclarator(
// linkage.
if (!getLangOpts().CPlusPlus && isa<VarDecl>(dd) &&
!isVarDeclStrongDefinition(astContext, *this, cast<VarDecl>(dd),
getCodeGenOpts().NoCommon)) {
errorNYI(dd->getBeginLoc(), "common linkage", dd->getDeclKindName());
getCodeGenOpts().NoCommon))
return cir::GlobalLinkageKind::CommonLinkage;
}

// selectany symbols are externally visible, so use weak instead of
// linkonce. MSVC optimizes away references to const selectany globals, so
Expand Down
4 changes: 1 addition & 3 deletions clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -1995,7 +1995,6 @@ void CIRToLLVMGlobalOpLowering::setupRegionInitializedLLVMGlobalOp(
// attributes are available on cir.global ops. This duplicates code
// in CIRToLLVMGlobalOpLowering::matchAndRewrite() but that will go
// away when the placeholders are no longer needed.
assert(!cir::MissingFeatures::opGlobalConstant());
const bool isConst = op.getConstant();
assert(!cir::MissingFeatures::addressSpace());
const unsigned addrSpace = 0;
Expand Down Expand Up @@ -2055,8 +2054,7 @@ mlir::LogicalResult CIRToLLVMGlobalOpLowering::matchAndRewrite(
convertTypeForMemory(*getTypeConverter(), dataLayout, cirSymType);
// FIXME: These default values are placeholders until the the equivalent
// attributes are available on cir.global ops.
assert(!cir::MissingFeatures::opGlobalConstant());
const bool isConst = false;
const bool isConst = op.getConstant();
assert(!cir::MissingFeatures::addressSpace());
const unsigned addrSpace = 0;
const bool isDsoLocal = op.getDsoLocal();
Expand Down
21 changes: 16 additions & 5 deletions clang/lib/CodeGen/CGExpr.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -4641,11 +4641,17 @@ LValue CodeGenFunction::EmitArraySubscriptExpr(const ArraySubscriptExpr *E,
LHS.getBaseInfo(), TBAAAccessInfo());
}

// The HLSL runtime handle the subscript expression on global resource arrays.
if (getLangOpts().HLSL && (E->getType()->isHLSLResourceRecord() ||
E->getType()->isHLSLResourceRecordArray())) {
std::optional<LValue> LV =
CGM.getHLSLRuntime().emitResourceArraySubscriptExpr(E, *this);
// The HLSL runtime handles subscript expressions on global resource arrays
// and objects with HLSL buffer layouts.
if (getLangOpts().HLSL) {
std::optional<LValue> LV;
if (E->getType()->isHLSLResourceRecord() ||
E->getType()->isHLSLResourceRecordArray()) {
LV = CGM.getHLSLRuntime().emitResourceArraySubscriptExpr(E, *this);
} else if (E->getType().getAddressSpace() == LangAS::hlsl_constant) {
LV = CGM.getHLSLRuntime().emitBufferArraySubscriptExpr(E, *this,
EmitIdxAfterBase);
}
if (LV.has_value())
return *LV;
}
Expand Down Expand Up @@ -5110,6 +5116,11 @@ LValue CodeGenFunction::EmitMemberExpr(const MemberExpr *E) {
EmitIgnoredExpr(E->getBase());
return EmitDeclRefLValue(DRE);
}
if (getLangOpts().HLSL &&
E->getType().getAddressSpace() == LangAS::hlsl_constant) {
// We have an HLSL buffer - emit using HLSL's layout rules.
return CGM.getHLSLRuntime().emitBufferMemberExpr(*this, E);
}

Expr *BaseExpr = E->getBase();
// Check whether the underlying base pointer is a constant null.
Expand Down
4 changes: 4 additions & 0 deletions clang/lib/CodeGen/CGExprAgg.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -2278,6 +2278,10 @@ void CodeGenFunction::EmitAggregateCopy(LValue Dest, LValue Src, QualType Ty,
}
}

if (getLangOpts().HLSL && Ty.getAddressSpace() == LangAS::hlsl_constant)
if (CGM.getHLSLRuntime().emitBufferCopy(*this, DestPtr, SrcPtr, Ty))
return;

// Aggregate assignment turns into llvm.memcpy. This is almost valid per
// C99 6.5.16.1p3, which states "If the value being stored in an object is
// read from another object that overlaps in anyway the storage of the first
Expand Down
Loading