16e712e7c3
Allow directly storing an immediate instead of requiring that it first be moved into a register. This makes for more compact and readable PTX. An approach similar to this (using a ComplexPattern) this could be used for most PTX instructions to avoid the need for `_[ri]+` variants and boiler-plate.
21 lines
634 B
LLVM
21 lines
634 B
LLVM
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
|
|
; RUN: llc -mtriple=nvptx64 < %s | FileCheck %s
|
|
|
|
define i64 @test() nounwind readnone {
|
|
; CHECK-LABEL: test(
|
|
; CHECK: {
|
|
; CHECK-NEXT: .reg .b64 %rd<3>;
|
|
; CHECK-EMPTY:
|
|
; CHECK-NEXT: // %bb.0:
|
|
; CHECK-NEXT: mov.b64 %rd1, 1;
|
|
; CHECK-NEXT: st.b64 [%rd1], 42;
|
|
; CHECK-NEXT: ld.global.b64 %rd2, [%rd1];
|
|
; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
|
|
; CHECK-NEXT: ret;
|
|
%addr0 = inttoptr i64 1 to ptr
|
|
%addr1 = inttoptr i64 1 to ptr addrspace(1)
|
|
store i64 42, ptr %addr0
|
|
%res = load i64, ptr addrspace(1) %addr1
|
|
ret i64 %res
|
|
}
|