[llvm] [NVPTX] Support address offsets added with disjoint or (PR #122042)
Alex MacLean via llvm-commits
llvm-commits at lists.llvm.org
Wed Jan 8 12:22:35 PST 2025
================
@@ -0,0 +1,25 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_35 -verify-machineinstrs | FileCheck %s
+; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_35 | %ptxas-verify %}
+target triple = "nvptx64-nvidia-cuda"
+
+ at a = external global ptr align 16
+
+define i32 @test_v2i8(i16 %a) {
+; CHECK-LABEL: test_v2i8(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: mov.u64 %rd1, a;
+; CHECK-NEXT: cvta.global.u64 %rd2, %rd1;
+; CHECK-NEXT: ld.volatile.u32 %r1, [%rd2+8];
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %a1 = ptrtoint ptr @a to i64
+ %a2 = or disjoint i64 %a1, 8
+ %a3 = inttoptr i64 %a2 to ptr
+ %v = load volatile i32, ptr %a3
----------------
AlexMaclean wrote:
No. Removed.
https://github.com/llvm/llvm-project/pull/122042
More information about the llvm-commits
mailing list