[NVPTX] Drop memory references of LDG/LDU

[asavonic]

Summary

This patch fixes machine verifier errors:

    *** Bad machine code: Missing mayLoad flag ***
    - function:    foo1
    - basic block: %bb.0  (0x5560fc64ef08)
    - instruction: %4:float32regs =
	INT_PTX_LDG_GLOBAL_f32areg64 killed %3:int64regs
	:: (load (s32) from %ir.from1, addrspace 1)

LDG/LDU should not be treated as loads because they operates on
read-only memory, and therefore these instructions should not have mayLoad flag.
Machine verifier checks for memoperands to determine whether an instruction is
a load, so dropping them during lowering fixes the problem.

[asavonic]

It appears that mayLoad was omitted for a reason, but I cannot figure out why. If a load is done from a read-only memory, it should still be treated as a load, right? Is there anything that requires special handling for these instructions?

Git blame shows that these comments were introduced by @jlebar in https://reviews.llvm.org/D17471, but there are no details either, so maybe I'm missing something.

[jlebar]

It appears that mayLoad was omitted for a reason, but I cannot figure out why. If a load is done from a read-only memory, it should still be treated as a load, right?

I think the notion is, if the load is from read-only memory that always has the same value for the entirety of the program, in what way is it like a "load", as opposed to e.g. "move immediate into register"?

You can safely reorder these instructions with any other memory operation.

Definitely a hack, and maybe there is today a better way of representing this in LLVM.

[jdoerfert]

Loads from constant memory that is known dereferenceable are not modeled as loads/reads in the IR, e.g., we derive readnone in their presence.
At least if you have "constant mem" TBAA annotations or some other encodings that inform the analyses about the constness of the memory.
That said, if the verifier complains we need to teach the verifier or avoid this. I am inclined to suggest the former if there is no reason not to.

[asavonic]

That makes sense, thanks a lot for the explanation!
Machine verifier checks for memoperands to determine whether an instruction is a load or a store. Perhaps we can just drop them for LDG/LDU if we don't want them to be counted as loads?

[jdoerfert]

I cannot really comment on the verifier change.

[tra]

While ldu does indeed specify that it loads from read-only memory, I do not think we can treat ld.global.nc the same way.
PTX spec says Load register variable d from the location specified by the source address operand a in the global state space, and optionally cache in non-coherent texture cache. Since the cache is non-coherent, the data should be read-only within the kernel's process.

The way I read it -- it's a regular load that bypasses cache. Unlike ldu, it does not specify that it's a read-only data. While read-only data will make it work correctly, it's not the only valid use case.
E.g. data may be changed from the host and two subsequent ld.global.nc will return different values.

[asavonic]

In D112466#3103803, @tra wrote:

While ldu does indeed specify that it loads from read-only memory, I do not think we can treat ld.global.nc the same way.
PTX spec says Load register variable d from the location specified by the source address operand a in the global state space, and optionally cache in non-coherent texture cache. Since the cache is non-coherent, the data should be read-only within the kernel's process.

The way I read it -- it's a regular load that bypasses cache.

I'm not sure, but I think the spec says that the load is cached:
Load [...] from the location [...] in the global state space, and optionally cache in non-coherent texture cache.
The problem is that the cache is non-coherent and optional.

Unlike ldu, it does not specify that it's a read-only data. While read-only data will make it work correctly, it's not the only valid use case.
E.g. data may be changed from the host and two subsequent ld.global.nc will return different values.

You mean something like this, right?

ld.global.nc r1, [a]
st.global [a], r2           ; may be updated from host?
ld.global.nc r3, [a]

I think the result is undefined: r1 and r3 can be equal if the first instruction cached the result. Otherwise, since the cache is optional, r2 and r3 can be equal.
However, what happens for this code?

st.global [a], r4
ld.global.nc r5, [a]

Provided that a is not cached before the store, ld.global.nc should always load the stored value, right?
If this is correct, then ld.global.nc can behave like a "load" in some cases, and therefore should have mayLoad flag.

[hliao]

These machine instruction should be marked with mayLoad as they are really load instructions. Check the attached patch fixing the machine verifier.

0001-nvptx-Mark-LDG-LDU-mayLoad.patch7 KBDownload

[asavonic]

In D112466#3106139, @hliao wrote:

These machine instruction should be marked with mayLoad as they are really load instructions. Check the attached patch fixing the machine verifier.

0001-nvptx-Mark-LDG-LDU-mayLoad.patch7 KBDownload

This is where we started from in this review request (see the revision 382033).
I think the consensus is that LDU is not a load, because it operates on read-only memory (and therefore can be safely excluded from memory-related analysis).
I'm not sure about LDG though.

[tra]

In D112466#3105940, @asavonic wrote:

In D112466#3103803, @tra wrote:

While ldu does indeed specify that it loads from read-only memory, I do not think we can treat ld.global.nc the same way.
PTX spec says Load register variable d from the location specified by the source address operand a in the global state space, and optionally cache in non-coherent texture cache. Since the cache is non-coherent, the data should be read-only within the kernel's process.

The way I read it -- it's a regular load that bypasses cache.

I'm not sure, but I think the spec says that the load is cached:
Load [...] from the location [...] in the global state space, and optionally cache in non-coherent texture cache.
The problem is that the cache is non-coherent and optional.

ld.global.nc is an odd instruction.

The optionally cache in texture cache applies only to the new variant of the instruction which specifies caching hints.

ld.global.nc{.level::eviction_priority}{.level::cache_hint}.type      d, [a]{, cache-policy};
ld.global.nc{.level::eviction_priority}{.level::cache_hint}.vec.type  d, [a]{, cache-policy};
...
Support for .level::eviction_priority and .level::cache_hint qualifiers introduced in PTX ISA version 7.4.
Support for .level::eviction_priority qualifier requires sm_70 or higher.
Support for .level::cache_hint qualifier requires sm_80 or higher.

I think the optionally-caching variants may have to be treated as separate instructions and those would be closer to ldu functionality-wise.

Unlike ldu, it does not specify that it's a read-only data. While read-only data will make it work correctly, it's not the only valid use case.
E.g. data may be changed from the host and two subsequent ld.global.nc will return different values.

You mean something like this, right?

ld.global.nc r1, [a]
st.global [a], r2           ; may be updated from host?

I mean that host could do cudaMemcpy(host->device, a, host_ptr, N).

ld.global.nc r3, [a]

I think the result is undefined: r1 and r3 can be equal if the first instruction cached the result. Otherwise, since the cache is optional, r2 and r3 can be equal.

ld.global.nc does no caching. If host copies something into 'a' between two loads, r1 and r3 will be different.

However, what happens for this code?

st.global [a], r4
ld.global.nc r5, [a]

Provided that a is not cached before the store, ld.global.nc should always load the stored value, right?
If this is correct, then ld.global.nc can behave like a "load" in some cases, and therefore should have mayLoad flag.

Result of this code is undefined. In general, writes are posted and the data may not make it to the memory by the time ld.global.nc attempts to read it.
According to PTX docs memory consistency does not apply to ld.global.nc: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#scope-and-applicability

The memory consistency model does not apply to texture (including ld.global.nc) and surface accesses.

My example of host/device accesses may be a bad one. It would equally affect regular loads/stores, too, and right now we generally assume that only GPU can change memory.
If that's the case, then ld.global.nc should not have mayLoad -- we can't guarantee consistency with any store ops anyways, so there's no point affecting their order.

[hliao]

In D112466#3106496, @asavonic wrote:

In D112466#3106139, @hliao wrote:

These machine instruction should be marked with mayLoad as they are really load instructions. Check the attached patch fixing the machine verifier.

0001-nvptx-Mark-LDG-LDU-mayLoad.patch7 KBDownload

This is where we started from in this review request (see the revision 382033).
I think the consensus is that LDU is not a load, because it operates on read-only memory (and therefore can be safely excluded from memory-related analysis).
I'm not sure about LDG though.

Memory analysis already knows any loads from constant memory won't inference with other loads or stores (pointsToConstantMemory). Loads from constant memory are still a load. They needs to be marked with mayLoad.

[asavonic]

In D112466#3106921, @hliao wrote:

In D112466#3106496, @asavonic wrote:

In D112466#3106139, @hliao wrote:

These machine instruction should be marked with mayLoad as they are really load instructions. Check the attached patch fixing the machine verifier.

0001-nvptx-Mark-LDG-LDU-mayLoad.patch7 KBDownload

This is where we started from in this review request (see the revision 382033).
I think the consensus is that LDU is not a load, because it operates on read-only memory (and therefore can be safely excluded from memory-related analysis).
I'm not sure about LDG though.

Memory analysis already knows any loads from constant memory won't inference with other loads or stores (pointsToConstantMemory). Loads from constant memory are still a load. They needs to be marked with mayLoad.

It seems that pointsToConstantMemory is not implemented for NVPTX, so we'll have fix that.
@jlebar, does this sound like a better solution?

[asavonic]

It seems that pointsToConstantMemory is not implemented for NVPTX, so we'll have fix that.

This is a function of AliasAnalysis and if I understand correctly there is no NVPTX-specific implementation of AA.
AMDGPU implementation has a set of rules based on pointer address spaces (e.g. pointers in global and private AS do not alias).
Does it make sense to add similar logic for NVPTX?

[tra]

In D112466#3122479, @asavonic wrote:

Does it make sense to add similar logic for NVPTX?

Yes. We can make similar no-aliasing conclusions for shared/constant/global/local/param memories.

[asavonic]

Sorry, this got pushed back by other things, but I'd like to return to this conversation. I've uploaded a patch to add a simple AA that takes NVPTX address spaces into account and implements pointsToConstantMemory (D124787). This patch should be useful on its own, but I'm no longer sure that it is enough to safely (performace-wise) add mayLoad to LDU/LDG instructions.

The problem is that AA::pointsToConstantMemory works with LLVM IR Value (MemoryLocation), and optimizations in CodeGen may not be able to use it or it is just not enabled yet. There are also cases where heuristics treat loads separately (BranchFolding, RegAllocScore), and it is not clear if we can change this and treat constant loads as regular instructions for all targets.

I'll go over these cases in a couple of days, but for now it seems like the current patch (drop memory references) is the safest option.

[asavonic]

• Rebased

[asavonic]

I think the current patch is the way to go: we want to treat LDU and LDG as regular instructions and exclude them from memory analysis. Therefore memory references are not needed and we can drop them.

pointsToConstantMemory or other AA functions do not seem to cover this case, because they operate on memory locations (pointers), and not on instructions themselves. Pointer operands for LDU and LDG are from global address space, and if I'm reading the PTX spec right, LDU can also load from generic address space.

[tra]

LGTM in principle.

llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp

1674–1675 ↗

(On Diff #488280)

The function also handles regular loads/stores (e.g. NVPTXISD::LoadV4) and those should still mark the node as a memory reference.

[asavonic]

llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp

1674–1675 ↗

(On Diff #488280)

Can you elaborate why we need to treat these nodes separately? A regular load can be lowered to LDG if it matches canLowerToLDG, but the final instruction is still LDG (aka ld.global.nc).

[tra]

llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp

1674–1675 ↗

(On Diff #488280)

Sorry, my mistake. I only looked at the case constant and didn't pay attention that we're lowering it into an LDG instruction.

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td

 defm INT_PTX_SATOM_MIN : ATOM2_minmax_impl<"min">;
 defm INT_PTX_SATOM_OR  : ATOM2_bitwise_impl<"or">;
 defm INT_PTX_SATOM_XOR : ATOM2_bitwise_impl<"xor">;
 
 //-----------------------------------
 // Support for ldu on sm_20 or later
 //-----------------------------------
 
-// Don't annotate ldu instructions as mayLoad, as they load from memory that is
-// read-only in a kernel.
-
+let mayLoad = 1 in {
 // Scalar
 
 multiclass LDU_G<string TyStr, NVPTXRegClass regclass> {
   def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
                !strconcat("ldu.global.", TyStr),
                       []>, Requires<[hasLDU]>;
   def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
                !strconcat("ldu.global.", TyStr),
: VLDU_G_ELE_V4<"v4.b16 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
     Float16Regs>;
 defm INT_PTX_LDU_G_v4f16x2_ELE
   : VLDU_G_ELE_V4<"v4.b32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
     Float16x2Regs>;
 defm INT_PTX_LDU_G_v4f32_ELE
   : VLDU_G_ELE_V4<"v4.f32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
     Float32Regs>;
 
+} // mayLoad
 
 //-----------------------------------
 // Support for ldg on sm_35 or later
 //-----------------------------------
 
-// Don't annotate ld.global.nc as mayLoad, because these loads go through the
-// non-coherent texture cache, and therefore the values read must be read-only
-// during the lifetime of the kernel.
-
+let mayLoad = 1 in {
 multiclass LDG_G<string TyStr, NVPTXRegClass regclass> {
   def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
                !strconcat("ld.global.nc.", TyStr),
                       []>, Requires<[hasLDG]>;
   def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
                !strconcat("ld.global.nc.", TyStr),
                         []>, Requires<[hasLDG]>;
  def avar:  NVPTXInst<(outs regclass:$result), (ins imemAny:$src),
defm INT_PTX_LDG_G_v4i32_ELE
   : VLDG_G_ELE_V4<"v4.u32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int32Regs>;
 defm INT_PTX_LDG_G_v4f16_ELE
   : VLDG_G_ELE_V4<"v4.b16 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Float16Regs>;
 defm INT_PTX_LDG_G_v4f16x2_ELE
   : VLDG_G_ELE_V4<"v4.b32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Float16x2Regs>;
 defm INT_PTX_LDG_G_v4f32_ELE
   : VLDG_G_ELE_V4<"v4.f32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Float32Regs>;
 
+} // mayLoad
 
 multiclass NG_TO_G<string Str, Intrinsic Intrin> {
    def _yes : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
           !strconcat("cvta.", Str, ".u32 \t$result, $src;"),
       [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>;
    def _yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
           !strconcat("cvta.", Str, ".u64 \t$result, $src;"),
       [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>;

llvm/test/CodeGen/NVPTX/bug26185-2.ll

-; RUN: llc < %s -march=nvptx -mcpu=sm_35 | FileCheck %s
+; RUN: llc < %s -march=nvptx -mcpu=sm_35 -verify-machineinstrs | FileCheck %s
 
 ; Verify that we correctly emit code for extending ldg/ldu. We do not expose
 ; extending variants in the backend, but the ldg/ldu selection code may pick
 ; extending loads as candidates. We do want to support this, so make sure we
 ; emit the necessary cvt.* instructions to implement the extension and let ptxas
 ; emit the real extending loads.
 
 target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"

llvm/test/CodeGen/NVPTX/bug26185.ll

-; RUN: llc < %s -march=nvptx -mcpu=sm_35 | FileCheck %s
+; RUN: llc < %s -march=nvptx -mcpu=sm_35 -verify-machineinstrs | FileCheck %s
 
 ; Verify that we correctly emit code for i8 ldg/ldu. We do not expose 8-bit
 ; registers in the backend, so these loads need special handling.
 
 target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
 target triple = "nvptx64-unknown-unknown"
 
 ; CHECK-LABEL: ex_zext

llvm/test/CodeGen/NVPTX/ldg-invariant.ll

-; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 | FileCheck %s
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 -verify-machineinstrs | FileCheck %s
 
 ; Check that invariant loads from the global addrspace are lowered to
 ; ld.global.nc.
 
 ; CHECK-LABEL: @ld_global
 define i32 @ld_global(i32 addrspace(1)* %ptr) {
 ; CHECK: ld.global.nc.{{[a-z]}}32
   %a = load i32, i32 addrspace(1)* %ptr, !invariant.load !0

llvm/test/CodeGen/NVPTX/ldu-i8.ll

-; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
+; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
 
 target datalayout = "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 declare i8 @llvm.nvvm.ldu.global.i.i8.p0i8(i8*, i32)
 
 define i8 @foo(i8* %a) {
 ; Ensure we properly truncate off the high-order 24 bits
 ; CHECK:        ldu.global.u8
 ; CHECK:        cvt.u32.u16
 ; CHECK:        and.b32         %r{{[0-9]+}}, %r{{[0-9]+}}, 255
   %val = tail call i8 @llvm.nvvm.ldu.global.i.i8.p0i8(i8* %a, i32 4)
   ret i8 %val
 }

llvm/test/CodeGen/NVPTX/ldu-ldg.ll

-; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
+; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
 
 
 declare i8 @llvm.nvvm.ldu.global.i.i8.p1i8(i8 addrspace(1)* %ptr, i32 %align)
 declare i32 @llvm.nvvm.ldu.global.i.i32.p1i32(i32 addrspace(1)* %ptr, i32 %align)
 declare i8 @llvm.nvvm.ldg.global.i.i8.p1i8(i8 addrspace(1)* %ptr, i32 %align)
 declare i32 @llvm.nvvm.ldg.global.i.i32.p1i32(i32 addrspace(1)* %ptr, i32 %align)
 
 

llvm/test/CodeGen/NVPTX/ldu-reg-plus-offset.ll

-; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
+; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
 
 target datalayout = "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 
 define void @reg_plus_offset(i32* %a) {
 ; CHECK:        ldu.global.u32  %r{{[0-9]+}}, [%r{{[0-9]+}}+32];
 ; CHECK:        ldu.global.u32  %r{{[0-9]+}}, [%r{{[0-9]+}}+36];
   %p2 = getelementptr i32, i32* %a, i32 8

llvm/test/CodeGen/NVPTX/load-with-non-coherent-cache.ll

-; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck -check-prefix=SM20 %s
-; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 | FileCheck -check-prefix=SM35 %s
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck -check-prefix=SM20 %s
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 -verify-machineinstrs | FileCheck -check-prefix=SM35 %s
 
 target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 target triple = "nvptx64-unknown-unknown"
 
 ; SM20-LABEL: .visible .entry foo1(
 ; SM20: ld.global.f32
 ; SM35-LABEL: .visible .entry foo1(
 ; SM35: ld.global.nc.f32

llvm/test/CodeGen/NVPTX/read-global-variable-constant.ll

-; RUN: llc < %s -march=nvptx -mcpu=sm_35 | FileCheck %s
+; RUN: llc < %s -march=nvptx -mcpu=sm_35 -verify-machineinstrs | FileCheck %s
 
 ; Check load from constant global variables.  These loads should be
 ; ld.global.nc (aka ldg).
 
 @gv_float = external constant float
 @gv_float2 = external constant <2 x float>
 @gv_float4 = external constant <4 x float>