Skip to content

Commit 539a6b1

Browse files
authored
1 parent 9dce22d commit 539a6b1

File tree

3 files changed

+9
-6
lines changed

3 files changed

+9
-6
lines changed

cmake/llvm-hash.txt

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1 +1 @@
1-
0ea4fb92648b2aa7cbab486bb493e122b4dcc062
1+
71a977d0d611f3e9f6137a6b8a26b730b2886ce9

python/test/unit/language/test_compile_only.py

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -54,7 +54,7 @@ def simple_dot(a_base, b_base, out):
5454
assert re.search(pattern, str(ttgir)), "The TTGIR does not match the expected pattern."
5555

5656
ptx = k.asm["ptx"]
57-
pattern = (r"mov\.u32 %r(?P<G>\d+), global_smem;"
57+
pattern = (r"mov\.b32 %r(?P<G>\d+), global_smem;"
5858
r"(.|\n)*"
5959
r"tcgen05\.alloc\.cta_group::1\.sync\.aligned\.shared::cta\.b32 \[%r(?P=G)], 64"
6060
r"(.|\n)*"

third_party/amd/lib/TritonAMDGPUToLLVM/LoadStoreOpToLLVM.cpp

Lines changed: 7 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -588,12 +588,15 @@ struct AsyncCopyGlobalToLocalOpConversion
588588
// size
589589
for (int i = 0; i < shmemAddrs.size(); i++) {
590590
auto srcIdx = i * maxVec;
591-
auto srcPtr = srcElems[srcIdx];
591+
Value srcPtr = srcElems[srcIdx];
592592

593593
if (maskElems.empty()) {
594594
rewriter.create<ROCDL::GlobalLoadLDSOp>(
595-
loc, srcPtr, shmemAddrs[i], vecBytesVal, /*offset=*/b.i32_val(0),
596-
cacheModifiers);
595+
loc,
596+
/*globalPtr=*/srcPtr, /*ldsPtr=*/shmemAddrs[i],
597+
/*size=*/vecBytesVal, /*offset=*/b.i32_val(0),
598+
/*aux=*/cacheModifiers, /*alias_scopes=*/nullptr,
599+
/*noalias_scopes=*/nullptr, /*tbaa=*/nullptr);
597600
continue;
598601
}
599602

@@ -607,7 +610,7 @@ struct AsyncCopyGlobalToLocalOpConversion
607610
rewriter.setInsertionPointToStart(loadBlock);
608611
rewriter.create<ROCDL::GlobalLoadLDSOp>(
609612
loc, srcPtr, shmemAddrs[i], vecBytesVal, /*offset=*/b.i32_val(0),
610-
cacheModifiers);
613+
cacheModifiers, nullptr, nullptr, nullptr);
611614

612615
rewriter.create<LLVM::BrOp>(loc, afterLoad);
613616
rewriter.setInsertionPointToStart(afterLoad);

0 commit comments

Comments
 (0)