[clang] [llvm] [AMDGPU] Add initial support for VGPR as memory (PR #205435)

Matt Arsenault via llvm-commits llvm-commits at lists.llvm.org
Thu Jun 25 07:46:57 PDT 2026


================
@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx1200 \
+// RUN:   -fcuda-is-device -emit-llvm -O1 -disable-llvm-passes -o - %s \
+// RUN:   | FileCheck %s
+
+#define __global__ __attribute__((global))
+
+// A kernel-local variable marked amdgpu_vgpr is allocated in the "VGPR as
+// memory" address space (addrspace(13)), and its accesses stay in that space.
+
+// CHECK-LABEL: define {{.*}}@_Z6kernelPii(
+// CHECK: %buf = alloca [4 x i32], align 4, addrspace(13)
----------------
arsenm wrote:

This should be a global variable like LDS 

https://github.com/llvm/llvm-project/pull/205435


More information about the llvm-commits mailing list