[PATCH] D127668: [NVPTX] Fix pointer type for short 32-bit pointers

[Andrew Savonichev via Phabricator via llvm-commits]

asavonic added inline comments.


================
Comment at: llvm/test/CodeGen/NVPTX/short-ptr.ll:18
+; CHECK-SHORT-LOCAL: .visible .local .align 8 .u32 l
+ at l = local_unnamed_addr addrspace(5) global i32 addrspace(5)* null, align 8
----------------
tra wrote:
> asavonic wrote:
> > Ptxas doesn't like global variables in local address space, so I will remove this case.
> > `ptxas error: Module-scoped variables in .local state space are not allowed with ABI`
> We could use `__const__` AS instead. 
> 
> For testing local vars we should probably add a function with a local variable.
> For testing local vars we should probably add a function with a local variable.

Local variables are lowered as `__local_depot` followed by as a sequence of casts:
```
	mov.u64 	%SPL, __local_depot0;
	cvta.local.u64 	%SP, %SPL;
	add.u64 	%rd1, %SP, 0;
	{ .reg .b64 %tmp;
	  cvta.to.local.u64 	%tmp, %rd1;
	  cvt.u32.u64 	%r1, %tmp;             <--- r1 is our 32-bit local pointer
        }
```
So I added a test, but it fails because of a different issue:
```
declare void @use(i8 addrspace(5)* %local)
define void @test() {
  %v = alloca i8
  %cast = addrspacecast i8* %v to i8 addrspace(5)*
  call void @use(i8 addrspace(5)* %cast)
  ret void
}

PTX:
.extern .func use
(
	.param .b64 use_param_0   <--- pointer is 64-bit instead of 32
)
[...]
.param .b32 param0;
st.param.b32 [param0+0], %r1;    <--- call assumes a 32-bit pointer
call.uni 
use, 
(
param0
);

ptxas: Type of argument does not match formal parameter 'use_param_0'
```
So either the call should pass the pointer as 64-bit, or the `use` function declaration is wrong. I'm not sure. Should `--nvptx-short-ptr` change function ABI?


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D127668/new/

https://reviews.llvm.org/D127668