Skip to content

[NVPTX] Switch to imm offset variants for LDG and LDU #128270

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 3 commits into from
Feb 25, 2025
Merged

[NVPTX] Switch to imm offset variants for LDG and LDU #128270

Changes from 1 commit
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
391e451
[NVPTX] Switch to imm offset variants for LDG and LDU
AlexMaclean Feb 21, 2025
c3ff962
address comments -- add tests
AlexMaclean Feb 25, 2025
6a08c58
address comments -- auto-gen tests
AlexMaclean Feb 25, 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
16 changes: 16 additions & 0 deletions llvm/test/CodeGen/NVPTX/ldu-ldg.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -148,3 +148,19 @@ define <2 x half> @test_ldg_v2f16(ptr addrspace(1) %ptr) {
%val = tail call <2 x half> @llvm.nvvm.ldg.global.f.v2f16.p1(ptr addrspace(1) %ptr, i32 4)
ret <2 x half> %val
}

@g = addrspace(1) global i32 0

; CHECK-LABEL: test_ldg_asi
Copy link
Member

@Artem-B Artem-B Feb 25, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Let's automate the check generation while we're here, too.

define i32 @test_ldg_asi() {
; CHECK: ld.global.nc.u32 %r{{.*}}, [g+4]
%val = tail call i32 @llvm.nvvm.ldg.global.i.i32.p1(ptr addrspace(1) getelementptr (i8, ptr addrspace(1) @g, i32 4), i32 4)
ret i32 %val
}

; CHECK-LABEL: test_lug_asi
define i32 @test_lug_asi() {
; CHECK: ldu.global.u32 %r{{.*}}, [g+4]
%val = tail call i32 @llvm.nvvm.ldu.global.i.i32.p1(ptr addrspace(1) getelementptr (i8, ptr addrspace(1) @g, i32 4), i32 4)
ret i32 %val
}