[clang] [llvm] [CHERI] Allow @llvm.returnaddress to return a pointer in any address space. (PR #188464)
Owen Anderson via cfe-commits
cfe-commits at lists.llvm.org
Wed Mar 25 05:07:30 PDT 2026
https://github.com/resistor updated https://github.com/llvm/llvm-project/pull/188464
>From b351671a337fb4d29519630b67192e0c355bbd10 Mon Sep 17 00:00:00 2001
From: Owen Anderson <resistor at mac.com>
Date: Wed, 25 Mar 2026 12:26:04 +0100
Subject: [PATCH 1/2] [CHERI] Allow @llvm.returnaddress to return a pointer in
any address space.
Clang now constructs calls to it using the default program address space from the DataLayout.
Co-authored-by: Alex Richardson <alexrichardson at google.com>
---
clang/lib/CodeGen/CGBuiltin.cpp | 6 ++++--
clang/lib/CodeGen/CodeGenModule.cpp | 2 ++
clang/lib/CodeGen/CodeGenTypeCache.h | 3 +++
clang/test/CodeGen/ms-intrinsics.c | 2 +-
clang/test/Headers/hip-header.hip | 4 ++--
llvm/include/llvm/IR/Intrinsics.td | 2 +-
llvm/lib/Target/AMDGPU/AMDGPUSwLowerLDS.cpp | 10 ++++++----
.../Instrumentation/ThreadSanitizer.cpp | 6 ++++--
.../Transforms/Utils/EntryExitInstrumenter.cpp | 12 +++++++++---
llvm/test/Bitcode/compatibility-3.6.ll | 2 +-
llvm/test/Bitcode/compatibility-3.7.ll | 2 +-
llvm/test/Bitcode/compatibility-3.8.ll | 2 +-
llvm/test/Bitcode/compatibility-3.9.ll | 2 +-
llvm/test/Bitcode/compatibility-4.0.ll | 2 +-
llvm/test/Bitcode/compatibility-5.0.ll | 2 +-
llvm/test/Bitcode/compatibility-6.0.ll | 2 +-
llvm/test/Bitcode/compatibility.ll | 2 +-
...sw-lower-lds-dynamic-indirect-access-asan.ll | 4 ++--
...dgpu-sw-lower-lds-dynamic-indirect-access.ll | 4 ++--
...amdgpu-sw-lower-lds-dynamic-lds-test-asan.ll | 4 ++--
.../amdgpu-sw-lower-lds-dynamic-lds-test.ll | 4 ++--
.../AMDGPU/amdgpu-sw-lower-lds-lower-all.ll | 8 ++++----
...multi-static-dynamic-indirect-access-asan.ll | 8 ++++----
...-lds-multi-static-dynamic-indirect-access.ll | 8 ++++----
...-sw-lower-lds-multiple-blocks-return-asan.ll | 4 ++--
...mdgpu-sw-lower-lds-multiple-blocks-return.ll | 5 +++--
...mdgpu-sw-lower-lds-non-kernel-declaration.ll | 4 ++--
...r-lds-static-dynamic-indirect-access-asan.ll | 4 ++--
...-lower-lds-static-dynamic-indirect-access.ll | 4 ++--
...sw-lower-lds-static-dynamic-lds-test-asan.ll | 4 ++--
...dgpu-sw-lower-lds-static-dynamic-lds-test.ll | 5 +++--
...-sw-lower-lds-static-indirect-access-asan.ll | 4 ++--
...tatic-indirect-access-function-param-asan.ll | 4 ++--
...lds-static-indirect-access-function-param.ll | 4 ++--
...ower-lds-static-indirect-access-lower-all.ll | 4 ++--
...er-lds-static-indirect-access-nested-asan.ll | 17 ++++++++---------
...w-lower-lds-static-indirect-access-nested.ll | 16 ++++++++--------
...s-static-indirect-access-no-kernel-lds-id.ll | 4 ++--
...mdgpu-sw-lower-lds-static-indirect-access.ll | 4 ++--
.../AMDGPU/amdgpu-sw-lower-lds-static-lds-O0.ll | 4 ++--
...mdgpu-sw-lower-lds-static-lds-no-heap-ptr.ll | 4 ++--
.../amdgpu-sw-lower-lds-static-lds-test-asan.ll | 4 ++--
...r-lds-static-lds-test-atomic-cmpxchg-asan.ll | 4 ++--
...-lower-lds-static-lds-test-atomicrmw-asan.ll | 4 ++--
.../amdgpu-sw-lower-lds-static-lds-test.ll | 4 ++--
...mdgpu-sw-lower-lds-static-lds-vector-ptrs.ll | 4 ++--
.../ThreadSanitizer/atomic-non-integer.ll | 12 ++++++------
llvm/test/Instrumentation/ThreadSanitizer/eh.ll | 10 +++++-----
.../ThreadSanitizer/no_sanitize_thread.ll | 2 +-
.../sanitize-thread-no-checking.ll | 2 +-
.../EntryExitInstrumenter/debug-info.ll | 4 ++--
.../EntryExitInstrumenter/mcount-with-frompc.ll | 2 +-
.../Transforms/EntryExitInstrumenter/mcount.ll | 16 ++++++++--------
.../pre-inliner-instrumentation.ll | 4 ++--
.../test/Verifier/LoongArch/intrinsic-immarg.ll | 2 +-
llvm/test/Verifier/intrinsic-immarg.ll | 2 +-
56 files changed, 146 insertions(+), 128 deletions(-)
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index df03e84ce9f81..51c5d970a0f84 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -4824,11 +4824,13 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
case Builtin::BI__builtin_return_address: {
Value *Depth = ConstantEmitter(*this).emitAbstract(E->getArg(0),
getContext().UnsignedIntTy);
- Function *F = CGM.getIntrinsic(Intrinsic::returnaddress);
+ Function *F =
+ CGM.getIntrinsic(Intrinsic::returnaddress, {CGM.ProgramPtrTy});
return RValue::get(Builder.CreateCall(F, Depth));
}
case Builtin::BI_ReturnAddress: {
- Function *F = CGM.getIntrinsic(Intrinsic::returnaddress);
+ Function *F =
+ CGM.getIntrinsic(Intrinsic::returnaddress, {CGM.ProgramPtrTy});
return RValue::get(Builder.CreateCall(F, Builder.getInt32(0)));
}
case Builtin::BI__builtin_frame_address: {
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index daaa846bf42bc..3fcd6f5f904db 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -451,6 +451,8 @@ CodeGenModule::CodeGenModule(ASTContext &C,
llvm::PointerType::get(LLVMContext, DL.getAllocaAddrSpace());
GlobalsInt8PtrTy =
llvm::PointerType::get(LLVMContext, DL.getDefaultGlobalsAddressSpace());
+ ProgramPtrTy =
+ llvm::PointerType::get(LLVMContext, DL.getProgramAddressSpace());
ConstGlobalsPtrTy = llvm::PointerType::get(
LLVMContext, C.getTargetAddressSpace(GetGlobalConstantAddressSpace()));
ASTAllocaAddressSpace = getTargetCodeGenInfo().getASTAllocaAddressSpace();
diff --git a/clang/lib/CodeGen/CodeGenTypeCache.h b/clang/lib/CodeGen/CodeGenTypeCache.h
index 015306bb97373..39ea8a681dc42 100644
--- a/clang/lib/CodeGen/CodeGenTypeCache.h
+++ b/clang/lib/CodeGen/CodeGenTypeCache.h
@@ -72,6 +72,9 @@ struct CodeGenTypeCache {
llvm::PointerType *GlobalsInt8PtrTy;
};
+ /// Pointer in program address space
+ llvm::PointerType *ProgramPtrTy;
+
/// void* in the address space for constant globals
llvm::PointerType *ConstGlobalsPtrTy;
diff --git a/clang/test/CodeGen/ms-intrinsics.c b/clang/test/CodeGen/ms-intrinsics.c
index 6528a63e380c2..271aced5e0b7c 100644
--- a/clang/test/CodeGen/ms-intrinsics.c
+++ b/clang/test/CodeGen/ms-intrinsics.c
@@ -134,7 +134,7 @@ void *test_ReturnAddress(void) {
return _ReturnAddress();
}
// CHECK-LABEL: define{{.*}}ptr @test_ReturnAddress()
-// CHECK: = tail call ptr @llvm.returnaddress(i32 0)
+// CHECK: = tail call ptr @llvm.returnaddress.p0(i32 0)
// CHECK: ret ptr
#if defined(__i386__) || defined(__x86_64__) || defined (__aarch64__)
diff --git a/clang/test/Headers/hip-header.hip b/clang/test/Headers/hip-header.hip
index 146a43b643dba..b4bd32dc9e3b0 100644
--- a/clang/test/Headers/hip-header.hip
+++ b/clang/test/Headers/hip-header.hip
@@ -169,7 +169,7 @@ __device__ double test_isnan() {
// MALLOC: call i64 @__ockl_dm_alloc
// NOMALLOC: call void @llvm.trap
// MALLOC-ASAN-LABEL: define weak {{.*}}ptr @malloc(i64
-// MALLOC-ASAN: call ptr @llvm.returnaddress(i32 0)
+// MALLOC-ASAN: call ptr @llvm.returnaddress.p0(i32 0)
// MALLOC-ASAN: call i64 @__asan_malloc_impl(i64 {{.*}}, i64 {{.*}})
__device__ void test_malloc(void *a) {
a = malloc(42);
@@ -183,7 +183,7 @@ __device__ void test_malloc(void *a) {
// MALLOC: call void @__ockl_dm_dealloc
// NOMALLOC: call void @llvm.trap
// MALLOC-ASAN-LABEL: define weak {{.*}}void @free(ptr
-// MALLOC-ASAN: call ptr @llvm.returnaddress(i32 0)
+// MALLOC-ASAN: call ptr @llvm.returnaddress.p0(i32 0)
// MALLOC-ASAN: call void @__asan_free_impl(i64 {{.*}}, i64 {{.*}})
__device__ void test_free(void *a) {
free(a);
diff --git a/llvm/include/llvm/IR/Intrinsics.td b/llvm/include/llvm/IR/Intrinsics.td
index 4469ff155b854..6d4b9bd4415ae 100644
--- a/llvm/include/llvm/IR/Intrinsics.td
+++ b/llvm/include/llvm/IR/Intrinsics.td
@@ -885,7 +885,7 @@ def int_swift_async_context_addr : Intrinsic<[llvm_ptr_ty], [], []>;
//===--------------------- Code Generator Intrinsics ----------------------===//
//
-def int_returnaddress : DefaultAttrsIntrinsic<[llvm_ptr_ty], [llvm_i32_ty],
+def int_returnaddress : DefaultAttrsIntrinsic<[llvm_anyptr_ty], [llvm_i32_ty],
[IntrNoMem, ImmArg<ArgIndex<0>>]>;
def int_addressofreturnaddress : DefaultAttrsIntrinsic<[llvm_anyptr_ty], [], [IntrNoMem]>;
def int_frameaddress : DefaultAttrsIntrinsic<[llvm_anyptr_ty], [llvm_i32_ty],
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUSwLowerLDS.cpp b/llvm/lib/Target/AMDGPU/AMDGPUSwLowerLDS.cpp
index 362c221aa1392..04383855b946b 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUSwLowerLDS.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUSwLowerLDS.cpp
@@ -867,8 +867,9 @@ void AMDGPUSwLowerLDS::lowerKernelLDSAccesses(Function *Func,
// Create a call to malloc function which does device global memory allocation
// with size equals to all LDS global accesses size in this kernel.
- Value *ReturnAddress =
- IRB.CreateIntrinsic(Intrinsic::returnaddress, {IRB.getInt32(0)});
+ Value *ReturnAddress = IRB.CreateIntrinsic(
+ Intrinsic::returnaddress, IRB.getPtrTy(AMDGPUAS::FLAT_ADDRESS),
+ {IRB.getInt32(0)});
FunctionCallee MallocFunc = M.getOrInsertFunction(
StringRef("__asan_malloc_impl"),
FunctionType::get(Int64Ty, {Int64Ty, Int64Ty}, false));
@@ -933,8 +934,9 @@ void AMDGPUSwLowerLDS::lowerKernelLDSAccesses(Function *Func,
FunctionCallee AsanFreeFunc = M.getOrInsertFunction(
StringRef("__asan_free_impl"),
FunctionType::get(IRB.getVoidTy(), {Int64Ty, Int64Ty}, false));
- Value *ReturnAddr =
- IRB.CreateIntrinsic(Intrinsic::returnaddress, IRB.getInt32(0));
+ Value *ReturnAddr = IRB.CreateIntrinsic(Intrinsic::returnaddress,
+ IRB.getPtrTy(AMDGPUAS::FLAT_ADDRESS),
+ IRB.getInt32(0));
Value *RAPToInt = IRB.CreatePtrToInt(ReturnAddr, Int64Ty);
Value *MallocPtrToInt = IRB.CreatePtrToInt(LoadMallocPtr, Int64Ty);
IRB.CreateCall(AsanFreeFunc, {MallocPtrToInt, RAPToInt});
diff --git a/llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp b/llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp
index 811911644106b..f05efd863fb74 100644
--- a/llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp
+++ b/llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp
@@ -577,8 +577,10 @@ bool ThreadSanitizer::sanitizeFunction(Function &F,
if ((Res || HasCalls) && ClInstrumentFuncEntryExit) {
InstrumentationIRBuilder IRB(&F.getEntryBlock(),
F.getEntryBlock().getFirstNonPHIIt());
- Value *ReturnAddress =
- IRB.CreateIntrinsic(Intrinsic::returnaddress, IRB.getInt32(0));
+ auto ProgramAsPtrTy = PointerType::get(F.getParent()->getContext(),
+ DL.getProgramAddressSpace());
+ Value *ReturnAddress = IRB.CreateIntrinsic(
+ Intrinsic::returnaddress, {ProgramAsPtrTy}, IRB.getInt32(0));
IRB.CreateCall(TsanFuncEntry, ReturnAddress);
EscapeEnumerator EE(F, "tsan_cleanup", ClHandleCxxExceptions);
diff --git a/llvm/lib/Transforms/Utils/EntryExitInstrumenter.cpp b/llvm/lib/Transforms/Utils/EntryExitInstrumenter.cpp
index 29c17ffc41a74..71a32664c7e77 100644
--- a/llvm/lib/Transforms/Utils/EntryExitInstrumenter.cpp
+++ b/llvm/lib/Transforms/Utils/EntryExitInstrumenter.cpp
@@ -53,8 +53,11 @@ static void insertCall(Function &CurFn, StringRef Func,
// On RISC-V, AArch64, and LoongArch, the `_mcount` function takes
// `__builtin_return_address(0)` as an argument since
// `__builtin_return_address(1)` is not available on these platforms.
+ auto ProgASPtr =
+ PointerType::get(C, M.getDataLayout().getProgramAddressSpace());
Instruction *RetAddr = CallInst::Create(
- Intrinsic::getOrInsertDeclaration(&M, Intrinsic::returnaddress),
+ Intrinsic::getOrInsertDeclaration(&M, Intrinsic::returnaddress,
+ {ProgASPtr}),
ConstantInt::get(Type::getInt32Ty(C), 0), "", InsertionPt);
RetAddr->setDebugLoc(DL);
@@ -77,13 +80,16 @@ static void insertCall(Function &CurFn, StringRef Func,
}
if (Func == "__cyg_profile_func_enter" || Func == "__cyg_profile_func_exit") {
- Type *ArgTypes[] = {PointerType::getUnqual(C), PointerType::getUnqual(C)};
+ auto ProgASPtr =
+ PointerType::get(C, M.getDataLayout().getProgramAddressSpace());
+ Type *ArgTypes[] = {ProgASPtr, ProgASPtr};
FunctionCallee Fn = M.getOrInsertFunction(
Func, FunctionType::get(Type::getVoidTy(C), ArgTypes, false));
Instruction *RetAddr = CallInst::Create(
- Intrinsic::getOrInsertDeclaration(&M, Intrinsic::returnaddress),
+ Intrinsic::getOrInsertDeclaration(&M, Intrinsic::returnaddress,
+ {ProgASPtr}),
ArrayRef<Value *>(ConstantInt::get(Type::getInt32Ty(C), 0)), "",
InsertionPt);
RetAddr->setDebugLoc(DL);
diff --git a/llvm/test/Bitcode/compatibility-3.6.ll b/llvm/test/Bitcode/compatibility-3.6.ll
index 2148e013126b3..62b5a88d085c6 100644
--- a/llvm/test/Bitcode/compatibility-3.6.ll
+++ b/llvm/test/Bitcode/compatibility-3.6.ll
@@ -1112,7 +1112,7 @@ declare void @llvm.instrprof_increment(i8*, i64, i32, i32)
!10 = !{!"rax"}
define void @intrinsics.codegen() {
call i8* @llvm.returnaddress(i32 1)
- ; CHECK: call ptr @llvm.returnaddress(i32 1)
+ ; CHECK: call ptr @llvm.returnaddress.p0(i32 1)
call i8* @llvm.frameaddress(i32 1)
; CHECK: call ptr @llvm.frameaddress.p0(i32 1)
diff --git a/llvm/test/Bitcode/compatibility-3.7.ll b/llvm/test/Bitcode/compatibility-3.7.ll
index fed9cce2a0091..61cc50ef4dead 100644
--- a/llvm/test/Bitcode/compatibility-3.7.ll
+++ b/llvm/test/Bitcode/compatibility-3.7.ll
@@ -1143,7 +1143,7 @@ declare void @llvm.instrprof_increment(i8*, i64, i32, i32)
!10 = !{!"rax"}
define void @intrinsics.codegen() {
call i8* @llvm.returnaddress(i32 1)
- ; CHECK: call ptr @llvm.returnaddress(i32 1)
+ ; CHECK: call ptr @llvm.returnaddress.p0(i32 1)
call i8* @llvm.frameaddress(i32 1)
; CHECK: call ptr @llvm.frameaddress.p0(i32 1)
diff --git a/llvm/test/Bitcode/compatibility-3.8.ll b/llvm/test/Bitcode/compatibility-3.8.ll
index 92695b9a41b80..19a5c0f7a4e1f 100644
--- a/llvm/test/Bitcode/compatibility-3.8.ll
+++ b/llvm/test/Bitcode/compatibility-3.8.ll
@@ -1298,7 +1298,7 @@ declare void @llvm.instrprof_increment(i8*, i64, i32, i32)
!10 = !{!"rax"}
define void @intrinsics.codegen() {
call i8* @llvm.returnaddress(i32 1)
- ; CHECK: call ptr @llvm.returnaddress(i32 1)
+ ; CHECK: call ptr @llvm.returnaddress.p0(i32 1)
call i8* @llvm.frameaddress(i32 1)
; CHECK: call ptr @llvm.frameaddress.p0(i32 1)
diff --git a/llvm/test/Bitcode/compatibility-3.9.ll b/llvm/test/Bitcode/compatibility-3.9.ll
index aa11917332e11..b29463940424a 100644
--- a/llvm/test/Bitcode/compatibility-3.9.ll
+++ b/llvm/test/Bitcode/compatibility-3.9.ll
@@ -1369,7 +1369,7 @@ declare void @llvm.instrprof_increment(i8*, i64, i32, i32)
!10 = !{!"rax"}
define void @intrinsics.codegen() {
call i8* @llvm.returnaddress(i32 1)
- ; CHECK: call ptr @llvm.returnaddress(i32 1)
+ ; CHECK: call ptr @llvm.returnaddress.p0(i32 1)
call i8* @llvm.frameaddress(i32 1)
; CHECK: call ptr @llvm.frameaddress.p0(i32 1)
diff --git a/llvm/test/Bitcode/compatibility-4.0.ll b/llvm/test/Bitcode/compatibility-4.0.ll
index cefccdc02c08c..0d3a024af511d 100644
--- a/llvm/test/Bitcode/compatibility-4.0.ll
+++ b/llvm/test/Bitcode/compatibility-4.0.ll
@@ -1369,7 +1369,7 @@ declare void @llvm.instrprof_increment(i8*, i64, i32, i32)
!10 = !{!"rax"}
define void @intrinsics.codegen() {
call i8* @llvm.returnaddress(i32 1)
- ; CHECK: call ptr @llvm.returnaddress(i32 1)
+ ; CHECK: call ptr @llvm.returnaddress.p0(i32 1)
call i8* @llvm.frameaddress(i32 1)
; CHECK: call ptr @llvm.frameaddress.p0(i32 1)
diff --git a/llvm/test/Bitcode/compatibility-5.0.ll b/llvm/test/Bitcode/compatibility-5.0.ll
index ae3e2e8ffbb0f..c59701c5915aa 100644
--- a/llvm/test/Bitcode/compatibility-5.0.ll
+++ b/llvm/test/Bitcode/compatibility-5.0.ll
@@ -1381,7 +1381,7 @@ declare void @llvm.instrprof_increment(i8*, i64, i32, i32)
!10 = !{!"rax"}
define void @intrinsics.codegen() {
call i8* @llvm.returnaddress(i32 1)
- ; CHECK: call ptr @llvm.returnaddress(i32 1)
+ ; CHECK: call ptr @llvm.returnaddress.p0(i32 1)
call i8* @llvm.frameaddress(i32 1)
; CHECK: call ptr @llvm.frameaddress.p0(i32 1)
diff --git a/llvm/test/Bitcode/compatibility-6.0.ll b/llvm/test/Bitcode/compatibility-6.0.ll
index cfb5ff7b350a2..f0b18a8c8145e 100644
--- a/llvm/test/Bitcode/compatibility-6.0.ll
+++ b/llvm/test/Bitcode/compatibility-6.0.ll
@@ -1391,7 +1391,7 @@ declare void @llvm.instrprof_increment(i8*, i64, i32, i32)
!10 = !{!"rax"}
define void @intrinsics.codegen() {
call i8* @llvm.returnaddress(i32 1)
- ; CHECK: call ptr @llvm.returnaddress(i32 1)
+ ; CHECK: call ptr @llvm.returnaddress.p0(i32 1)
call i8* @llvm.frameaddress(i32 1)
; CHECK: call ptr @llvm.frameaddress.p0(i32 1)
diff --git a/llvm/test/Bitcode/compatibility.ll b/llvm/test/Bitcode/compatibility.ll
index c87159fe960f3..f2df4c68404fb 100644
--- a/llvm/test/Bitcode/compatibility.ll
+++ b/llvm/test/Bitcode/compatibility.ll
@@ -1887,7 +1887,7 @@ declare void @llvm.instrprof_increment(ptr, i64, i32, i32)
!10 = !{!"rax"}
define void @intrinsics.codegen() {
call ptr @llvm.returnaddress(i32 1)
- ; CHECK: call ptr @llvm.returnaddress(i32 1)
+ ; CHECK: call ptr @llvm.returnaddress.p0(i32 1)
call ptr @llvm.frameaddress(i32 1)
; CHECK: call ptr @llvm.frameaddress.p0(i32 1)
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-dynamic-indirect-access-asan.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-dynamic-indirect-access-asan.ll
index 4e53df3924985..e3a28b6379077 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-dynamic-indirect-access-asan.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-dynamic-indirect-access-asan.ll
@@ -122,7 +122,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: store i32 [[TMP19]], ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_K0_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.k0.md, i32 0, i32 4, i32 2), align 4
; CHECK-NEXT: [[TMP28:%.*]] = add i32 [[TMP15]], [[TMP19]]
; CHECK-NEXT: [[TMP26:%.*]] = zext i32 [[TMP28]] to i64
-; CHECK-NEXT: [[TMP22:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP22:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP23:%.*]] = ptrtoint ptr [[TMP22]] to i64
; CHECK-NEXT: [[TMP35:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP26]], i64 [[TMP23]])
; CHECK-NEXT: [[TMP20:%.*]] = inttoptr i64 [[TMP35]] to ptr addrspace(1)
@@ -227,7 +227,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label [[FREE:%.*]], label [[END:%.*]]
; CHECK: Free:
-; CHECK-NEXT: [[TMP32:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP32:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP33:%.*]] = ptrtoint ptr [[TMP32]] to i64
; CHECK-NEXT: [[TMP34:%.*]] = ptrtoint ptr addrspace(1) [[TMP31]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP34]], i64 [[TMP33]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-dynamic-indirect-access.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-dynamic-indirect-access.ll
index 8cbeb80d62335..c155a99ccca80 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-dynamic-indirect-access.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-dynamic-indirect-access.ll
@@ -76,7 +76,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: store i32 [[TMP19]], ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_K0_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.k0.md, i32 0, i32 4, i32 2), align 4
; CHECK-NEXT: [[TMP28:%.*]] = add i32 [[TMP15]], [[TMP19]]
; CHECK-NEXT: [[TMP26:%.*]] = zext i32 [[TMP28]] to i64
-; CHECK-NEXT: [[TMP22:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP22:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP23:%.*]] = ptrtoint ptr [[TMP22]] to i64
; CHECK-NEXT: [[TMP35:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP26]], i64 [[TMP23]])
; CHECK-NEXT: [[TMP20:%.*]] = inttoptr i64 [[TMP35]] to ptr addrspace(1)
@@ -112,7 +112,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label [[FREE:%.*]], label [[END:%.*]]
; CHECK: Free:
-; CHECK-NEXT: [[TMP32:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP32:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP33:%.*]] = ptrtoint ptr [[TMP32]] to i64
; CHECK-NEXT: [[TMP34:%.*]] = ptrtoint ptr addrspace(1) [[TMP31]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP34]], i64 [[TMP33]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-dynamic-lds-test-asan.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-dynamic-lds-test-asan.ll
index 32601422c7e67..12dcc92f49dc6 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-dynamic-lds-test-asan.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-dynamic-lds-test-asan.ll
@@ -36,7 +36,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: store i32 [[TMP16]], ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_K0_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.k0.md, i32 0, i32 1, i32 2), align 4
; CHECK-NEXT: [[TMP17:%.*]] = add i32 [[TMP24]], [[TMP16]]
; CHECK-NEXT: [[TMP21:%.*]] = zext i32 [[TMP17]] to i64
-; CHECK-NEXT: [[TMP22:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP22:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP23:%.*]] = ptrtoint ptr [[TMP22]] to i64
; CHECK-NEXT: [[TMP19:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP21]], i64 [[TMP23]])
; CHECK-NEXT: [[TMP6:%.*]] = inttoptr i64 [[TMP19]] to ptr addrspace(1)
@@ -82,7 +82,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label [[FREE:%.*]], label [[END:%.*]]
; CHECK: Free:
-; CHECK-NEXT: [[TMP25:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP25:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP26:%.*]] = ptrtoint ptr [[TMP25]] to i64
; CHECK-NEXT: [[TMP27:%.*]] = ptrtoint ptr addrspace(1) [[TMP28]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP27]], i64 [[TMP26]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-dynamic-lds-test.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-dynamic-lds-test.ll
index 5e90eb0b95219..f6876702dc0bb 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-dynamic-lds-test.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-dynamic-lds-test.ll
@@ -36,7 +36,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: store i32 [[TMP16]], ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_K0_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.k0.md, i32 0, i32 1, i32 2), align 4
; CHECK-NEXT: [[TMP17:%.*]] = add i32 [[TMP24]], [[TMP16]]
; CHECK-NEXT: [[TMP21:%.*]] = zext i32 [[TMP17]] to i64
-; CHECK-NEXT: [[TMP22:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP22:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP23:%.*]] = ptrtoint ptr [[TMP22]] to i64
; CHECK-NEXT: [[TMP19:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP21]], i64 [[TMP23]])
; CHECK-NEXT: [[TMP6:%.*]] = inttoptr i64 [[TMP19]] to ptr addrspace(1)
@@ -60,7 +60,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label [[FREE:%.*]], label [[END:%.*]]
; CHECK: Free:
-; CHECK-NEXT: [[TMP25:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP25:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP26:%.*]] = ptrtoint ptr [[TMP25]] to i64
; CHECK-NEXT: [[TMP27:%.*]] = ptrtoint ptr addrspace(1) [[TMP28]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP27]], i64 [[TMP26]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-lower-all.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-lower-all.ll
index f30a382a62c6b..bb8e762bafe6e 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-lower-all.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-lower-all.ll
@@ -22,7 +22,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: [[TMP7:%.*]] = load i32, ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_K0_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.k0.md, i32 0, i32 2, i32 2), align 4
; CHECK-NEXT: [[TMP8:%.*]] = add i32 [[TMP6]], [[TMP7]]
; CHECK-NEXT: [[TMP9:%.*]] = zext i32 [[TMP8]] to i64
-; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP11:%.*]] = ptrtoint ptr [[TMP10]] to i64
; CHECK-NEXT: [[TMP12:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP9]], i64 [[TMP11]])
; CHECK-NEXT: [[TMP13:%.*]] = inttoptr i64 [[TMP12]] to ptr addrspace(1)
@@ -56,7 +56,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label %[[FREE:.*]], label %[[END:.*]]
; CHECK: [[FREE]]:
-; CHECK-NEXT: [[TMP30:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP30:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP31:%.*]] = ptrtoint ptr [[TMP30]] to i64
; CHECK-NEXT: [[TMP32:%.*]] = ptrtoint ptr addrspace(1) [[TMP21]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP32]], i64 [[TMP31]])
@@ -85,7 +85,7 @@ define amdgpu_kernel void @k1() {
; CHECK-NEXT: [[TMP7:%.*]] = load i32, ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_K1_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.k1.md, i32 0, i32 1, i32 2), align 4
; CHECK-NEXT: [[TMP8:%.*]] = add i32 [[TMP6]], [[TMP7]]
; CHECK-NEXT: [[TMP9:%.*]] = zext i32 [[TMP8]] to i64
-; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP11:%.*]] = ptrtoint ptr [[TMP10]] to i64
; CHECK-NEXT: [[TMP12:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP9]], i64 [[TMP11]])
; CHECK-NEXT: [[TMP13:%.*]] = inttoptr i64 [[TMP12]] to ptr addrspace(1)
@@ -111,7 +111,7 @@ define amdgpu_kernel void @k1() {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label %[[FREE:.*]], label %[[END:.*]]
; CHECK: [[FREE]]:
-; CHECK-NEXT: [[TMP24:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP24:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP25:%.*]] = ptrtoint ptr [[TMP24]] to i64
; CHECK-NEXT: [[TMP26:%.*]] = ptrtoint ptr addrspace(1) [[TMP19]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP26]], i64 [[TMP25]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-multi-static-dynamic-indirect-access-asan.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-multi-static-dynamic-indirect-access-asan.ll
index 91e0a9fc5018b..ca41da7ec42e9 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-multi-static-dynamic-indirect-access-asan.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-multi-static-dynamic-indirect-access-asan.ll
@@ -210,7 +210,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: store i32 [[TMP19]], ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_K0_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.k0.md, i32 0, i32 3, i32 2), align 4
; CHECK-NEXT: [[TMP26:%.*]] = add i32 [[TMP15]], [[TMP19]]
; CHECK-NEXT: [[TMP27:%.*]] = zext i32 [[TMP26]] to i64
-; CHECK-NEXT: [[TMP28:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP28:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP33:%.*]] = ptrtoint ptr [[TMP28]] to i64
; CHECK-NEXT: [[TMP24:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP27]], i64 [[TMP33]])
; CHECK-NEXT: [[TMP20:%.*]] = inttoptr i64 [[TMP24]] to ptr addrspace(1)
@@ -260,7 +260,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label [[FREE:%.*]], label [[END:%.*]]
; CHECK: Free:
-; CHECK-NEXT: [[TMP30:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP30:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP31:%.*]] = ptrtoint ptr [[TMP30]] to i64
; CHECK-NEXT: [[TMP32:%.*]] = ptrtoint ptr addrspace(1) [[TMP29]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP32]], i64 [[TMP31]])
@@ -307,7 +307,7 @@ define amdgpu_kernel void @k1() sanitize_address {
; CHECK-NEXT: store i32 [[TMP24]], ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_K1_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.k1.md, i32 0, i32 4, i32 2), align 4
; CHECK-NEXT: [[TMP25:%.*]] = add i32 [[TMP20]], [[TMP24]]
; CHECK-NEXT: [[TMP26:%.*]] = zext i32 [[TMP25]] to i64
-; CHECK-NEXT: [[TMP27:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP27:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP28:%.*]] = ptrtoint ptr [[TMP27]] to i64
; CHECK-NEXT: [[TMP34:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP26]], i64 [[TMP28]])
; CHECK-NEXT: [[TMP13:%.*]] = inttoptr i64 [[TMP34]] to ptr addrspace(1)
@@ -361,7 +361,7 @@ define amdgpu_kernel void @k1() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label [[FREE:%.*]], label [[END:%.*]]
; CHECK: Free:
-; CHECK-NEXT: [[TMP35:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP35:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP36:%.*]] = ptrtoint ptr [[TMP35]] to i64
; CHECK-NEXT: [[TMP37:%.*]] = ptrtoint ptr addrspace(1) [[TMP29]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP37]], i64 [[TMP36]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-multi-static-dynamic-indirect-access.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-multi-static-dynamic-indirect-access.ll
index d0caddb7934a7..f10bdd9c05ef0 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-multi-static-dynamic-indirect-access.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-multi-static-dynamic-indirect-access.ll
@@ -97,7 +97,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: store i32 [[TMP19]], ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_K0_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.k0.md, i32 0, i32 3, i32 2), align 4
; CHECK-NEXT: [[TMP26:%.*]] = add i32 [[TMP15]], [[TMP19]]
; CHECK-NEXT: [[TMP27:%.*]] = zext i32 [[TMP26]] to i64
-; CHECK-NEXT: [[TMP28:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP28:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP33:%.*]] = ptrtoint ptr [[TMP28]] to i64
; CHECK-NEXT: [[TMP24:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP27]], i64 [[TMP33]])
; CHECK-NEXT: [[TMP20:%.*]] = inttoptr i64 [[TMP24]] to ptr addrspace(1)
@@ -125,7 +125,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label [[FREE:%.*]], label [[END:%.*]]
; CHECK: Free:
-; CHECK-NEXT: [[TMP30:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP30:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP31:%.*]] = ptrtoint ptr [[TMP30]] to i64
; CHECK-NEXT: [[TMP32:%.*]] = ptrtoint ptr addrspace(1) [[TMP29]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP32]], i64 [[TMP31]])
@@ -172,7 +172,7 @@ define amdgpu_kernel void @k1() sanitize_address {
; CHECK-NEXT: store i32 [[TMP24]], ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_K1_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.k1.md, i32 0, i32 4, i32 2), align 4
; CHECK-NEXT: [[TMP25:%.*]] = add i32 [[TMP20]], [[TMP24]]
; CHECK-NEXT: [[TMP26:%.*]] = zext i32 [[TMP25]] to i64
-; CHECK-NEXT: [[TMP27:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP27:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP28:%.*]] = ptrtoint ptr [[TMP27]] to i64
; CHECK-NEXT: [[TMP34:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP26]], i64 [[TMP28]])
; CHECK-NEXT: [[TMP13:%.*]] = inttoptr i64 [[TMP34]] to ptr addrspace(1)
@@ -204,7 +204,7 @@ define amdgpu_kernel void @k1() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label [[FREE:%.*]], label [[END:%.*]]
; CHECK: Free:
-; CHECK-NEXT: [[TMP35:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP35:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP36:%.*]] = ptrtoint ptr [[TMP35]] to i64
; CHECK-NEXT: [[TMP37:%.*]] = ptrtoint ptr addrspace(1) [[TMP29]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP37]], i64 [[TMP36]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-multiple-blocks-return-asan.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-multiple-blocks-return-asan.ll
index 07baf90e370d1..44522c1d8ebbe 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-multiple-blocks-return-asan.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-multiple-blocks-return-asan.ll
@@ -27,7 +27,7 @@ define amdgpu_kernel void @test_kernel() sanitize_address {
; CHECK-NEXT: [[TMP16:%.*]] = load i32, ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_TEST_KERNEL_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.test_kernel.md, i32 0, i32 2, i32 2), align 4
; CHECK-NEXT: [[TMP18:%.*]] = add i32 [[TMP15]], [[TMP16]]
; CHECK-NEXT: [[TMP17:%.*]] = zext i32 [[TMP18]] to i64
-; CHECK-NEXT: [[TMP14:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP14:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP19:%.*]] = ptrtoint ptr [[TMP14]] to i64
; CHECK-NEXT: [[TMP20:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP17]], i64 [[TMP19]])
; CHECK-NEXT: [[TMP6:%.*]] = inttoptr i64 [[TMP20]] to ptr addrspace(1)
@@ -70,7 +70,7 @@ define amdgpu_kernel void @test_kernel() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label [[FREE:%.*]], label [[END:%.*]]
; CHECK: Free:
-; CHECK-NEXT: [[TMP22:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP22:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP23:%.*]] = ptrtoint ptr [[TMP22]] to i64
; CHECK-NEXT: [[TMP24:%.*]] = ptrtoint ptr addrspace(1) [[TMP21]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP24]], i64 [[TMP23]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-multiple-blocks-return.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-multiple-blocks-return.ll
index 6848e2c06c1e1..2cc5a2af75859 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-multiple-blocks-return.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-multiple-blocks-return.ll
@@ -27,7 +27,7 @@ define amdgpu_kernel void @test_kernel() sanitize_address {
; CHECK-NEXT: [[TMP16:%.*]] = load i32, ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_TEST_KERNEL_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.test_kernel.md, i32 0, i32 2, i32 2), align 4
; CHECK-NEXT: [[TMP18:%.*]] = add i32 [[TMP15]], [[TMP16]]
; CHECK-NEXT: [[TMP17:%.*]] = zext i32 [[TMP18]] to i64
-; CHECK-NEXT: [[TMP14:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP14:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP19:%.*]] = ptrtoint ptr [[TMP14]] to i64
; CHECK-NEXT: [[TMP20:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP17]], i64 [[TMP19]])
; CHECK-NEXT: [[TMP6:%.*]] = inttoptr i64 [[TMP20]] to ptr addrspace(1)
@@ -70,7 +70,7 @@ define amdgpu_kernel void @test_kernel() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label [[FREE:%.*]], label [[END:%.*]]
; CHECK: Free:
-; CHECK-NEXT: [[TMP22:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP22:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP23:%.*]] = ptrtoint ptr [[TMP22]] to i64
; CHECK-NEXT: [[TMP24:%.*]] = ptrtoint ptr addrspace(1) [[TMP21]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP24]], i64 [[TMP23]])
@@ -109,4 +109,5 @@ ret void
; CHECK: attributes #[[ATTR3:[0-9]+]] = { convergent nocallback nofree nounwind willreturn }
;.
; CHECK: [[META0]] = !{i32 0, i32 1}
+; CHECK: [[META1:![0-9]+]] = !{i32 4, !"nosanitize_address", i32 1}
;.
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-non-kernel-declaration.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-non-kernel-declaration.ll
index a6e6b84bba304..fb50a74ec55b5 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-non-kernel-declaration.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-non-kernel-declaration.ll
@@ -50,7 +50,7 @@ define amdgpu_kernel void @k1() sanitize_address {
; CHECK-NEXT: [[TMP7:%.*]] = load i32, ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_K1_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.k1.md, i32 0, i32 1, i32 2), align 4
; CHECK-NEXT: [[TMP8:%.*]] = add i32 [[TMP6]], [[TMP7]]
; CHECK-NEXT: [[TMP9:%.*]] = zext i32 [[TMP8]] to i64
-; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP11:%.*]] = ptrtoint ptr [[TMP10]] to i64
; CHECK-NEXT: [[TMP12:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP9]], i64 [[TMP11]])
; CHECK-NEXT: [[TMP13:%.*]] = inttoptr i64 [[TMP12]] to ptr addrspace(1)
@@ -73,7 +73,7 @@ define amdgpu_kernel void @k1() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label %[[FREE:.*]], label %[[END:.*]]
; CHECK: [[FREE]]:
-; CHECK-NEXT: [[TMP20:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP20:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP21:%.*]] = ptrtoint ptr [[TMP20]] to i64
; CHECK-NEXT: [[TMP22:%.*]] = ptrtoint ptr addrspace(1) [[TMP19]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP22]], i64 [[TMP21]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-dynamic-indirect-access-asan.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-dynamic-indirect-access-asan.ll
index bad2d8e0fb5f4..310398ad6948c 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-dynamic-indirect-access-asan.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-dynamic-indirect-access-asan.ll
@@ -123,7 +123,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: store i32 [[TMP19]], ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_K0_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.k0.md, i32 0, i32 4, i32 2), align 4
; CHECK-NEXT: [[TMP28:%.*]] = add i32 [[TMP15]], [[TMP19]]
; CHECK-NEXT: [[TMP26:%.*]] = zext i32 [[TMP28]] to i64
-; CHECK-NEXT: [[TMP22:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP22:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP23:%.*]] = ptrtoint ptr [[TMP22]] to i64
; CHECK-NEXT: [[TMP35:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP26]], i64 [[TMP23]])
; CHECK-NEXT: [[TMP20:%.*]] = inttoptr i64 [[TMP35]] to ptr addrspace(1)
@@ -228,7 +228,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label [[FREE:%.*]], label [[END:%.*]]
; CHECK: Free:
-; CHECK-NEXT: [[TMP32:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP32:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP33:%.*]] = ptrtoint ptr [[TMP32]] to i64
; CHECK-NEXT: [[TMP34:%.*]] = ptrtoint ptr addrspace(1) [[TMP31]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP34]], i64 [[TMP33]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-dynamic-indirect-access.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-dynamic-indirect-access.ll
index 0cc49c94e2279..3b4785872a675 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-dynamic-indirect-access.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-dynamic-indirect-access.ll
@@ -77,7 +77,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: store i32 [[TMP19]], ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_K0_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.k0.md, i32 0, i32 4, i32 2), align 4
; CHECK-NEXT: [[TMP28:%.*]] = add i32 [[TMP15]], [[TMP19]]
; CHECK-NEXT: [[TMP26:%.*]] = zext i32 [[TMP28]] to i64
-; CHECK-NEXT: [[TMP22:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP22:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP23:%.*]] = ptrtoint ptr [[TMP22]] to i64
; CHECK-NEXT: [[TMP35:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP26]], i64 [[TMP23]])
; CHECK-NEXT: [[TMP20:%.*]] = inttoptr i64 [[TMP35]] to ptr addrspace(1)
@@ -113,7 +113,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label [[FREE:%.*]], label [[END:%.*]]
; CHECK: Free:
-; CHECK-NEXT: [[TMP32:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP32:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP33:%.*]] = ptrtoint ptr [[TMP32]] to i64
; CHECK-NEXT: [[TMP34:%.*]] = ptrtoint ptr addrspace(1) [[TMP31]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP34]], i64 [[TMP33]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-dynamic-lds-test-asan.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-dynamic-lds-test-asan.ll
index c5985e5cc4df8..e5c606addb3b0 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-dynamic-lds-test-asan.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-dynamic-lds-test-asan.ll
@@ -46,7 +46,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: store i32 [[TMP19]], ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_K0_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.k0.md, i32 0, i32 4, i32 2), align 4
; CHECK-NEXT: [[TMP32:%.*]] = add i32 [[TMP15]], [[TMP19]]
; CHECK-NEXT: [[TMP30:%.*]] = zext i32 [[TMP32]] to i64
-; CHECK-NEXT: [[TMP22:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP22:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP23:%.*]] = ptrtoint ptr [[TMP22]] to i64
; CHECK-NEXT: [[TMP39:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP30]], i64 [[TMP23]])
; CHECK-NEXT: [[TMP20:%.*]] = inttoptr i64 [[TMP39]] to ptr addrspace(1)
@@ -180,7 +180,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label [[FREE:%.*]], label [[END:%.*]]
; CHECK: Free:
-; CHECK-NEXT: [[TMP36:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP36:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP37:%.*]] = ptrtoint ptr [[TMP36]] to i64
; CHECK-NEXT: [[TMP38:%.*]] = ptrtoint ptr addrspace(1) [[TMP35]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP38]], i64 [[TMP37]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-dynamic-lds-test.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-dynamic-lds-test.ll
index e0bfca0f63ca7..53a651e11bccc 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-dynamic-lds-test.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-dynamic-lds-test.ll
@@ -46,7 +46,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: store i32 [[TMP19]], ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_K0_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.k0.md, i32 0, i32 4, i32 2), align 4
; CHECK-NEXT: [[TMP32:%.*]] = add i32 [[TMP15]], [[TMP19]]
; CHECK-NEXT: [[TMP30:%.*]] = zext i32 [[TMP32]] to i64
-; CHECK-NEXT: [[TMP22:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP22:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP23:%.*]] = ptrtoint ptr [[TMP22]] to i64
; CHECK-NEXT: [[TMP39:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP30]], i64 [[TMP23]])
; CHECK-NEXT: [[TMP20:%.*]] = inttoptr i64 [[TMP39]] to ptr addrspace(1)
@@ -91,7 +91,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label [[FREE:%.*]], label [[END:%.*]]
; CHECK: Free:
-; CHECK-NEXT: [[TMP36:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP36:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP37:%.*]] = ptrtoint ptr [[TMP36]] to i64
; CHECK-NEXT: [[TMP38:%.*]] = ptrtoint ptr addrspace(1) [[TMP35]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP38]], i64 [[TMP37]])
@@ -117,4 +117,5 @@ define amdgpu_kernel void @k0() sanitize_address {
;.
; CHECK: [[META0]] = !{i32 0, i32 1}
; CHECK: [[META1]] = !{i32 8, i32 9}
+; CHECK: [[META2:![0-9]+]] = !{i32 4, !"nosanitize_address", i32 1}
;.
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-asan.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-asan.ll
index fa888a35cb8ba..9f2287bb54924 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-asan.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-asan.ll
@@ -86,7 +86,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: [[TMP14:%.*]] = load i32, ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_K0_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.k0.md, i32 0, i32 4, i32 2), align 4
; CHECK-NEXT: [[TMP16:%.*]] = add i32 [[TMP13]], [[TMP14]]
; CHECK-NEXT: [[TMP15:%.*]] = zext i32 [[TMP16]] to i64
-; CHECK-NEXT: [[TMP23:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP23:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP24:%.*]] = ptrtoint ptr [[TMP23]] to i64
; CHECK-NEXT: [[TMP12:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP15]], i64 [[TMP24]])
; CHECK-NEXT: [[TMP6:%.*]] = inttoptr i64 [[TMP12]] to ptr addrspace(1)
@@ -196,7 +196,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label %[[FREE:.*]], label %[[END:.*]]
; CHECK: [[FREE]]:
-; CHECK-NEXT: [[TMP20:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP20:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP21:%.*]] = ptrtoint ptr [[TMP20]] to i64
; CHECK-NEXT: [[TMP22:%.*]] = ptrtoint ptr addrspace(1) [[TMP19]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP22]], i64 [[TMP21]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-function-param-asan.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-function-param-asan.ll
index a521d9d9d436b..6ca68376befb4 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-function-param-asan.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-function-param-asan.ll
@@ -96,7 +96,7 @@ define amdgpu_kernel void @my_kernel() sanitize_address {
; CHECK-NEXT: [[TMP12:%.*]] = load i32, ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_MY_KERNEL_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.my_kernel.md, i32 0, i32 1, i32 2), align 4
; CHECK-NEXT: [[TMP14:%.*]] = add i32 [[TMP11]], [[TMP12]]
; CHECK-NEXT: [[TMP13:%.*]] = zext i32 [[TMP14]] to i64
-; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP15:%.*]] = ptrtoint ptr [[TMP10]] to i64
; CHECK-NEXT: [[TMP16:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP13]], i64 [[TMP15]])
; CHECK-NEXT: [[TMP6:%.*]] = inttoptr i64 [[TMP16]] to ptr addrspace(1)
@@ -121,7 +121,7 @@ define amdgpu_kernel void @my_kernel() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label [[FREE:%.*]], label [[END:%.*]]
; CHECK: Free:
-; CHECK-NEXT: [[TMP18:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP18:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP19:%.*]] = ptrtoint ptr [[TMP18]] to i64
; CHECK-NEXT: [[TMP20:%.*]] = ptrtoint ptr addrspace(1) [[TMP17]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP20]], i64 [[TMP19]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-function-param.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-function-param.ll
index 55a36f85dc73a..07d4764a1852a 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-function-param.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-function-param.ll
@@ -49,7 +49,7 @@ define amdgpu_kernel void @my_kernel() sanitize_address {
; CHECK-NEXT: [[TMP12:%.*]] = load i32, ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_MY_KERNEL_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.my_kernel.md, i32 0, i32 1, i32 2), align 4
; CHECK-NEXT: [[TMP14:%.*]] = add i32 [[TMP11]], [[TMP12]]
; CHECK-NEXT: [[TMP13:%.*]] = zext i32 [[TMP14]] to i64
-; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP15:%.*]] = ptrtoint ptr [[TMP10]] to i64
; CHECK-NEXT: [[TMP16:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP13]], i64 [[TMP15]])
; CHECK-NEXT: [[TMP6:%.*]] = inttoptr i64 [[TMP16]] to ptr addrspace(1)
@@ -74,7 +74,7 @@ define amdgpu_kernel void @my_kernel() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label [[FREE:%.*]], label [[END:%.*]]
; CHECK: Free:
-; CHECK-NEXT: [[TMP18:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP18:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP19:%.*]] = ptrtoint ptr [[TMP18]] to i64
; CHECK-NEXT: [[TMP20:%.*]] = ptrtoint ptr addrspace(1) [[TMP17]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP20]], i64 [[TMP19]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-lower-all.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-lower-all.ll
index 4625a7f626f9b..8acb379606ecc 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-lower-all.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-lower-all.ll
@@ -62,7 +62,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: [[TMP7:%.*]] = load i32, ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_K0_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.k0.md, i32 0, i32 4, i32 2), align 4
; CHECK-NEXT: [[TMP8:%.*]] = add i32 [[TMP6]], [[TMP7]]
; CHECK-NEXT: [[TMP9:%.*]] = zext i32 [[TMP8]] to i64
-; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP11:%.*]] = ptrtoint ptr [[TMP10]] to i64
; CHECK-NEXT: [[TMP12:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP9]], i64 [[TMP11]])
; CHECK-NEXT: [[TMP13:%.*]] = inttoptr i64 [[TMP12]] to ptr addrspace(1)
@@ -103,7 +103,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label %[[FREE:.*]], label %[[END:.*]]
; CHECK: [[FREE]]:
-; CHECK-NEXT: [[TMP34:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP34:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP35:%.*]] = ptrtoint ptr [[TMP34]] to i64
; CHECK-NEXT: [[TMP36:%.*]] = ptrtoint ptr addrspace(1) [[TMP25]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP36]], i64 [[TMP35]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-nested-asan.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-nested-asan.ll
index 255dda562c1ea..0e12aae72320d 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-nested-asan.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-nested-asan.ll
@@ -6,7 +6,6 @@
@A = external addrspace(3) global [8 x ptr]
@B = external addrspace(3) global [0 x i32]
-;.
; @llvm.amdgcn.sw.lds.kernel_0 = internal addrspace(3) global ptr poison, no_sanitize_address, align 8, !absolute_symbol [[META0:![0-9]+]]
; @llvm.amdgcn.sw.lds.kernel_0.md = internal addrspace(1) global %llvm.amdgcn.sw.lds.kernel_0.md.type { %llvm.amdgcn.sw.lds.kernel_0.md.item { i32 0, i32 8, i32 32 }, %llvm.amdgcn.sw.lds.kernel_0.md.item { i32 32, i32 64, i32 96 } }, no_sanitize_address
; @llvm.amdgcn.sw.lds.kernel_2 = internal addrspace(3) global ptr poison, no_sanitize_address, align 8, !absolute_symbol [[META0]]
@@ -36,7 +35,7 @@ define amdgpu_kernel void @kernel_0() sanitize_address {
; CHECK-NEXT: [[TMP7:%.*]] = load i32, ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_KERNEL_0_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.kernel_0.md, i32 0, i32 1, i32 2), align 4
; CHECK-NEXT: [[TMP8:%.*]] = add i32 [[TMP6]], [[TMP7]]
; CHECK-NEXT: [[TMP9:%.*]] = zext i32 [[TMP8]] to i64
-; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP11:%.*]] = ptrtoint ptr [[TMP10]] to i64
; CHECK-NEXT: [[TMP12:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP9]], i64 [[TMP11]])
; CHECK-NEXT: [[TMP13:%.*]] = inttoptr i64 [[TMP12]] to ptr addrspace(1)
@@ -58,7 +57,7 @@ define amdgpu_kernel void @kernel_0() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label %[[FREE:.*]], label %[[END:.*]]
; CHECK: [[FREE]]:
-; CHECK-NEXT: [[TMP20:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP20:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP21:%.*]] = ptrtoint ptr [[TMP20]] to i64
; CHECK-NEXT: [[TMP22:%.*]] = ptrtoint ptr addrspace(1) [[TMP19]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP22]], i64 [[TMP21]])
@@ -96,7 +95,7 @@ define amdgpu_kernel void @kernel_1() sanitize_address {
; CHECK-NEXT: store i32 [[TMP14]], ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_KERNEL_1_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.kernel_1.md, i32 0, i32 1, i32 2), align 4
; CHECK-NEXT: [[TMP15:%.*]] = add i32 [[TMP8]], [[TMP14]]
; CHECK-NEXT: [[TMP16:%.*]] = zext i32 [[TMP15]] to i64
-; CHECK-NEXT: [[TMP17:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP17:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP18:%.*]] = ptrtoint ptr [[TMP17]] to i64
; CHECK-NEXT: [[TMP19:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP16]], i64 [[TMP18]])
; CHECK-NEXT: [[TMP20:%.*]] = inttoptr i64 [[TMP19]] to ptr addrspace(1)
@@ -116,7 +115,7 @@ define amdgpu_kernel void @kernel_1() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label %[[FREE:.*]], label %[[END:.*]]
; CHECK: [[FREE]]:
-; CHECK-NEXT: [[TMP25:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP25:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP26:%.*]] = ptrtoint ptr [[TMP25]] to i64
; CHECK-NEXT: [[TMP27:%.*]] = ptrtoint ptr addrspace(1) [[TMP24]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP27]], i64 [[TMP26]])
@@ -144,7 +143,7 @@ define amdgpu_kernel void @kernel_2() sanitize_address {
; CHECK-NEXT: [[TMP7:%.*]] = load i32, ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_KERNEL_2_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.kernel_2.md, i32 0, i32 1, i32 2), align 4
; CHECK-NEXT: [[TMP8:%.*]] = add i32 [[TMP6]], [[TMP7]]
; CHECK-NEXT: [[TMP9:%.*]] = zext i32 [[TMP8]] to i64
-; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP11:%.*]] = ptrtoint ptr [[TMP10]] to i64
; CHECK-NEXT: [[TMP12:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP9]], i64 [[TMP11]])
; CHECK-NEXT: [[TMP13:%.*]] = inttoptr i64 [[TMP12]] to ptr addrspace(1)
@@ -166,7 +165,7 @@ define amdgpu_kernel void @kernel_2() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label %[[FREE:.*]], label %[[END:.*]]
; CHECK: [[FREE]]:
-; CHECK-NEXT: [[TMP20:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP20:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP21:%.*]] = ptrtoint ptr [[TMP20]] to i64
; CHECK-NEXT: [[TMP22:%.*]] = ptrtoint ptr addrspace(1) [[TMP19]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP22]], i64 [[TMP21]])
@@ -204,7 +203,7 @@ define amdgpu_kernel void @kernel_3() sanitize_address {
; CHECK-NEXT: store i32 [[TMP14]], ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_KERNEL_3_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.kernel_3.md, i32 0, i32 1, i32 2), align 4
; CHECK-NEXT: [[TMP15:%.*]] = add i32 [[TMP8]], [[TMP14]]
; CHECK-NEXT: [[TMP16:%.*]] = zext i32 [[TMP15]] to i64
-; CHECK-NEXT: [[TMP17:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP17:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP18:%.*]] = ptrtoint ptr [[TMP17]] to i64
; CHECK-NEXT: [[TMP19:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP16]], i64 [[TMP18]])
; CHECK-NEXT: [[TMP20:%.*]] = inttoptr i64 [[TMP19]] to ptr addrspace(1)
@@ -224,7 +223,7 @@ define amdgpu_kernel void @kernel_3() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label %[[FREE:.*]], label %[[END:.*]]
; CHECK: [[FREE]]:
-; CHECK-NEXT: [[TMP25:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP25:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP26:%.*]] = ptrtoint ptr [[TMP25]] to i64
; CHECK-NEXT: [[TMP27:%.*]] = ptrtoint ptr addrspace(1) [[TMP24]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP27]], i64 [[TMP26]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-nested.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-nested.ll
index 7184ebbb8faa3..ada116e5bbda7 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-nested.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-nested.ll
@@ -36,7 +36,7 @@ define amdgpu_kernel void @kernel_0() sanitize_address {
; CHECK-NEXT: [[TMP10:%.*]] = load i32, ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_KERNEL_0_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.kernel_0.md, i32 0, i32 1, i32 2), align 4
; CHECK-NEXT: [[TMP12:%.*]] = add i32 [[TMP9]], [[TMP10]]
; CHECK-NEXT: [[TMP11:%.*]] = zext i32 [[TMP12]] to i64
-; CHECK-NEXT: [[TMP13:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP13:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP14:%.*]] = ptrtoint ptr [[TMP13]] to i64
; CHECK-NEXT: [[TMP19:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP11]], i64 [[TMP14]])
; CHECK-NEXT: [[TMP6:%.*]] = inttoptr i64 [[TMP19]] to ptr addrspace(1)
@@ -58,7 +58,7 @@ define amdgpu_kernel void @kernel_0() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label %[[FREE:.*]], label %[[END:.*]]
; CHECK: [[FREE]]:
-; CHECK-NEXT: [[TMP16:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP16:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP17:%.*]] = ptrtoint ptr [[TMP16]] to i64
; CHECK-NEXT: [[TMP18:%.*]] = ptrtoint ptr addrspace(1) [[TMP15]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP18]], i64 [[TMP17]])
@@ -96,7 +96,7 @@ define amdgpu_kernel void @kernel_1() sanitize_address {
; CHECK-NEXT: store i32 [[TMP11]], ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_KERNEL_1_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.kernel_1.md, i32 0, i32 1, i32 2), align 4
; CHECK-NEXT: [[TMP15:%.*]] = add i32 [[TMP21]], [[TMP11]]
; CHECK-NEXT: [[TMP16:%.*]] = zext i32 [[TMP15]] to i64
-; CHECK-NEXT: [[TMP17:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP17:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP18:%.*]] = ptrtoint ptr [[TMP17]] to i64
; CHECK-NEXT: [[TMP19:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP16]], i64 [[TMP18]])
; CHECK-NEXT: [[TMP13:%.*]] = inttoptr i64 [[TMP19]] to ptr addrspace(1)
@@ -116,7 +116,7 @@ define amdgpu_kernel void @kernel_1() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label %[[FREE:.*]], label %[[END:.*]]
; CHECK: [[FREE]]:
-; CHECK-NEXT: [[TMP23:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP23:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP24:%.*]] = ptrtoint ptr [[TMP23]] to i64
; CHECK-NEXT: [[TMP25:%.*]] = ptrtoint ptr addrspace(1) [[TMP22]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP25]], i64 [[TMP24]])
@@ -144,7 +144,7 @@ define amdgpu_kernel void @kernel_2() sanitize_address {
; CHECK-NEXT: [[TMP10:%.*]] = load i32, ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_KERNEL_2_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.kernel_2.md, i32 0, i32 1, i32 2), align 4
; CHECK-NEXT: [[TMP12:%.*]] = add i32 [[TMP9]], [[TMP10]]
; CHECK-NEXT: [[TMP11:%.*]] = zext i32 [[TMP12]] to i64
-; CHECK-NEXT: [[TMP13:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP13:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP14:%.*]] = ptrtoint ptr [[TMP13]] to i64
; CHECK-NEXT: [[TMP19:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP11]], i64 [[TMP14]])
; CHECK-NEXT: [[TMP6:%.*]] = inttoptr i64 [[TMP19]] to ptr addrspace(1)
@@ -166,7 +166,7 @@ define amdgpu_kernel void @kernel_2() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label %[[FREE:.*]], label %[[END:.*]]
; CHECK: [[FREE]]:
-; CHECK-NEXT: [[TMP16:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP16:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP17:%.*]] = ptrtoint ptr [[TMP16]] to i64
; CHECK-NEXT: [[TMP18:%.*]] = ptrtoint ptr addrspace(1) [[TMP15]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP18]], i64 [[TMP17]])
@@ -204,7 +204,7 @@ define amdgpu_kernel void @kernel_3() sanitize_address {
; CHECK-NEXT: store i32 [[TMP11]], ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_KERNEL_3_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.kernel_3.md, i32 0, i32 1, i32 2), align 4
; CHECK-NEXT: [[TMP15:%.*]] = add i32 [[TMP21]], [[TMP11]]
; CHECK-NEXT: [[TMP16:%.*]] = zext i32 [[TMP15]] to i64
-; CHECK-NEXT: [[TMP17:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP17:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP18:%.*]] = ptrtoint ptr [[TMP17]] to i64
; CHECK-NEXT: [[TMP19:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP16]], i64 [[TMP18]])
; CHECK-NEXT: [[TMP13:%.*]] = inttoptr i64 [[TMP19]] to ptr addrspace(1)
@@ -224,7 +224,7 @@ define amdgpu_kernel void @kernel_3() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label %[[FREE:.*]], label %[[END:.*]]
; CHECK: [[FREE]]:
-; CHECK-NEXT: [[TMP23:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP23:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP24:%.*]] = ptrtoint ptr [[TMP23]] to i64
; CHECK-NEXT: [[TMP25:%.*]] = ptrtoint ptr addrspace(1) [[TMP22]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP25]], i64 [[TMP24]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-no-kernel-lds-id.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-no-kernel-lds-id.ll
index 704bc9e635294..13e4f28c4a88f 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-no-kernel-lds-id.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-no-kernel-lds-id.ll
@@ -63,7 +63,7 @@ define amdgpu_kernel void @k0() sanitize_address #1 {
; CHECK-NEXT: [[TMP7:%.*]] = load i32, ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_K0_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.k0.md, i32 0, i32 4, i32 2), align 4
; CHECK-NEXT: [[TMP8:%.*]] = add i32 [[TMP6]], [[TMP7]]
; CHECK-NEXT: [[TMP9:%.*]] = zext i32 [[TMP8]] to i64
-; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP11:%.*]] = ptrtoint ptr [[TMP10]] to i64
; CHECK-NEXT: [[TMP12:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP9]], i64 [[TMP11]])
; CHECK-NEXT: [[TMP13:%.*]] = inttoptr i64 [[TMP12]] to ptr addrspace(1)
@@ -104,7 +104,7 @@ define amdgpu_kernel void @k0() sanitize_address #1 {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label %[[FREE:.*]], label %[[END:.*]]
; CHECK: [[FREE]]:
-; CHECK-NEXT: [[TMP34:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP34:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP35:%.*]] = ptrtoint ptr [[TMP34]] to i64
; CHECK-NEXT: [[TMP36:%.*]] = ptrtoint ptr addrspace(1) [[TMP25]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP36]], i64 [[TMP35]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access.ll
index 8f5abe962f8eb..24bce7406e1b0 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access.ll
@@ -62,7 +62,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: [[TMP14:%.*]] = load i32, ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_K0_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.k0.md, i32 0, i32 4, i32 2), align 4
; CHECK-NEXT: [[TMP16:%.*]] = add i32 [[TMP13]], [[TMP14]]
; CHECK-NEXT: [[TMP15:%.*]] = zext i32 [[TMP16]] to i64
-; CHECK-NEXT: [[TMP23:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP23:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP24:%.*]] = ptrtoint ptr [[TMP23]] to i64
; CHECK-NEXT: [[TMP12:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP15]], i64 [[TMP24]])
; CHECK-NEXT: [[TMP6:%.*]] = inttoptr i64 [[TMP12]] to ptr addrspace(1)
@@ -103,7 +103,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label %[[FREE:.*]], label %[[END:.*]]
; CHECK: [[FREE]]:
-; CHECK-NEXT: [[TMP20:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP20:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP21:%.*]] = ptrtoint ptr [[TMP20]] to i64
; CHECK-NEXT: [[TMP22:%.*]] = ptrtoint ptr addrspace(1) [[TMP19]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP22]], i64 [[TMP21]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-O0.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-O0.ll
index 1973a0acf4659..934ae1634d722 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-O0.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-O0.ll
@@ -22,7 +22,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: [[TMP7:%.*]] = load i32, ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_K0_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.k0.md, i32 0, i32 1, i32 2), align 4
; CHECK-NEXT: [[TMP8:%.*]] = add i32 [[TMP6]], [[TMP7]]
; CHECK-NEXT: [[TMP9:%.*]] = zext i32 [[TMP8]] to i64
-; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP11:%.*]] = ptrtoint ptr [[TMP10]] to i64
; CHECK-NEXT: [[TMP12:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP9]], i64 [[TMP11]])
; CHECK-NEXT: [[TMP13:%.*]] = inttoptr i64 [[TMP12]] to ptr addrspace(1)
@@ -50,7 +50,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label %[[FREE:.*]], label %[[END:.*]]
; CHECK: [[FREE]]:
-; CHECK-NEXT: [[TMP25:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP25:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP26:%.*]] = ptrtoint ptr [[TMP25]] to i64
; CHECK-NEXT: [[TMP27:%.*]] = ptrtoint ptr addrspace(1) [[TMP19]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP27]], i64 [[TMP26]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-no-heap-ptr.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-no-heap-ptr.ll
index ab3300ea659b8..eca24adad5258 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-no-heap-ptr.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-no-heap-ptr.ll
@@ -26,7 +26,7 @@ define amdgpu_kernel void @k0() sanitize_address #1 {
; CHECK-NEXT: [[TMP7:%.*]] = load i32, ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_K0_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.k0.md, i32 0, i32 2, i32 2), align 4
; CHECK-NEXT: [[TMP8:%.*]] = add i32 [[TMP6]], [[TMP7]]
; CHECK-NEXT: [[TMP9:%.*]] = zext i32 [[TMP8]] to i64
-; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP11:%.*]] = ptrtoint ptr [[TMP10]] to i64
; CHECK-NEXT: [[TMP12:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP9]], i64 [[TMP11]])
; CHECK-NEXT: [[TMP13:%.*]] = inttoptr i64 [[TMP12]] to ptr addrspace(1)
@@ -129,7 +129,7 @@ define amdgpu_kernel void @k0() sanitize_address #1 {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label %[[FREE:.*]], label %[[END:.*]]
; CHECK: [[FREE]]:
-; CHECK-NEXT: [[TMP78:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP78:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP79:%.*]] = ptrtoint ptr [[TMP78]] to i64
; CHECK-NEXT: [[TMP80:%.*]] = ptrtoint ptr addrspace(1) [[TMP21]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP80]], i64 [[TMP79]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-test-asan.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-test-asan.ll
index c7550dd9576ec..aa54a96d3658c 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-test-asan.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-test-asan.ll
@@ -25,7 +25,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: [[TMP14:%.*]] = load i32, ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_K0_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.k0.md, i32 0, i32 2, i32 2), align 4
; CHECK-NEXT: [[TMP16:%.*]] = add i32 [[TMP13]], [[TMP14]]
; CHECK-NEXT: [[TMP15:%.*]] = zext i32 [[TMP16]] to i64
-; CHECK-NEXT: [[TMP23:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP23:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP11:%.*]] = ptrtoint ptr [[TMP23]] to i64
; CHECK-NEXT: [[TMP12:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP15]], i64 [[TMP11]])
; CHECK-NEXT: [[TMP6:%.*]] = inttoptr i64 [[TMP12]] to ptr addrspace(1)
@@ -128,7 +128,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label [[FREE:%.*]], label [[END:%.*]]
; CHECK: Free:
-; CHECK-NEXT: [[TMP20:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP20:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP21:%.*]] = ptrtoint ptr [[TMP20]] to i64
; CHECK-NEXT: [[TMP22:%.*]] = ptrtoint ptr addrspace(1) [[TMP19]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP22]], i64 [[TMP21]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-test-atomic-cmpxchg-asan.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-test-atomic-cmpxchg-asan.ll
index 15b074c2d9c11..fe3e677485c54 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-test-atomic-cmpxchg-asan.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-test-atomic-cmpxchg-asan.ll
@@ -23,7 +23,7 @@ define amdgpu_kernel void @atomic_xchg_kernel(ptr addrspace(1) %out, [8 x i32],
; CHECK-NEXT: [[TMP9:%.*]] = load i32, ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_ATOMIC_XCHG_KERNEL_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.atomic_xchg_kernel.md, i32 0, i32 1, i32 2), align 4
; CHECK-NEXT: [[TMP10:%.*]] = add i32 [[TMP8]], [[TMP9]]
; CHECK-NEXT: [[TMP11:%.*]] = zext i32 [[TMP10]] to i64
-; CHECK-NEXT: [[TMP12:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP12:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP13:%.*]] = ptrtoint ptr [[TMP12]] to i64
; CHECK-NEXT: [[TMP14:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP11]], i64 [[TMP13]])
; CHECK-NEXT: [[TMP15:%.*]] = inttoptr i64 [[TMP14]] to ptr addrspace(1)
@@ -99,7 +99,7 @@ define amdgpu_kernel void @atomic_xchg_kernel(ptr addrspace(1) %out, [8 x i32],
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label [[FREE:%.*]], label [[END:%.*]]
; CHECK: Free:
-; CHECK-NEXT: [[TMP43:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP43:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP44:%.*]] = ptrtoint ptr [[TMP43]] to i64
; CHECK-NEXT: [[TMP45:%.*]] = ptrtoint ptr addrspace(1) [[TMP21]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP45]], i64 [[TMP44]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-test-atomicrmw-asan.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-test-atomicrmw-asan.ll
index 1b3664bf1e4e7..502b1cc9ec5ef 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-test-atomicrmw-asan.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-test-atomicrmw-asan.ll
@@ -25,7 +25,7 @@ define amdgpu_kernel void @atomicrmw_kernel(ptr addrspace(1) %arg0) sanitize_add
; CHECK-NEXT: [[TMP7:%.*]] = load i32, ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_ATOMICRMW_KERNEL_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.atomicrmw_kernel.md, i32 0, i32 2, i32 2), align 4
; CHECK-NEXT: [[TMP8:%.*]] = add i32 [[TMP6]], [[TMP7]]
; CHECK-NEXT: [[TMP9:%.*]] = zext i32 [[TMP8]] to i64
-; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP11:%.*]] = ptrtoint ptr [[TMP10]] to i64
; CHECK-NEXT: [[TMP12:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP9]], i64 [[TMP11]])
; CHECK-NEXT: [[TMP13:%.*]] = inttoptr i64 [[TMP12]] to ptr addrspace(1)
@@ -181,7 +181,7 @@ define amdgpu_kernel void @atomicrmw_kernel(ptr addrspace(1) %arg0) sanitize_add
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label [[FREE:%.*]], label [[END:%.*]]
; CHECK: Free:
-; CHECK-NEXT: [[TMP84:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP84:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP85:%.*]] = ptrtoint ptr [[TMP84]] to i64
; CHECK-NEXT: [[TMP86:%.*]] = ptrtoint ptr addrspace(1) [[TMP21]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP86]], i64 [[TMP85]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-test.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-test.ll
index 066b9429425ac..6c5f32f0a4e3a 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-test.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-test.ll
@@ -26,7 +26,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: [[TMP14:%.*]] = load i32, ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_K0_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.k0.md, i32 0, i32 2, i32 2), align 4
; CHECK-NEXT: [[TMP16:%.*]] = add i32 [[TMP13]], [[TMP14]]
; CHECK-NEXT: [[TMP15:%.*]] = zext i32 [[TMP16]] to i64
-; CHECK-NEXT: [[TMP23:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP23:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP11:%.*]] = ptrtoint ptr [[TMP23]] to i64
; CHECK-NEXT: [[TMP12:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP15]], i64 [[TMP11]])
; CHECK-NEXT: [[TMP6:%.*]] = inttoptr i64 [[TMP12]] to ptr addrspace(1)
@@ -60,7 +60,7 @@ define amdgpu_kernel void @k0() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label [[FREE:%.*]], label [[END:%.*]]
; CHECK: Free:
-; CHECK-NEXT: [[TMP20:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP20:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP21:%.*]] = ptrtoint ptr [[TMP20]] to i64
; CHECK-NEXT: [[TMP22:%.*]] = ptrtoint ptr addrspace(1) [[TMP19]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP22]], i64 [[TMP21]])
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-vector-ptrs.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-vector-ptrs.ll
index 34caf91def933..7146965b1c66d 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-vector-ptrs.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-vector-ptrs.ll
@@ -25,7 +25,7 @@ define amdgpu_kernel void @example() sanitize_address {
; CHECK-NEXT: [[TMP7:%.*]] = load i32, ptr addrspace(1) getelementptr inbounds ([[LLVM_AMDGCN_SW_LDS_EXAMPLE_MD_TYPE]], ptr addrspace(1) @llvm.amdgcn.sw.lds.example.md, i32 0, i32 2, i32 2), align 4
; CHECK-NEXT: [[TMP8:%.*]] = add i32 [[TMP6]], [[TMP7]]
; CHECK-NEXT: [[TMP9:%.*]] = zext i32 [[TMP8]] to i64
-; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP10:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP11:%.*]] = ptrtoint ptr [[TMP10]] to i64
; CHECK-NEXT: [[TMP12:%.*]] = call i64 @__asan_malloc_impl(i64 [[TMP9]], i64 [[TMP11]])
; CHECK-NEXT: [[TMP13:%.*]] = inttoptr i64 [[TMP12]] to ptr addrspace(1)
@@ -62,7 +62,7 @@ define amdgpu_kernel void @example() sanitize_address {
; CHECK-NEXT: call void @llvm.amdgcn.s.barrier()
; CHECK-NEXT: br i1 [[XYZCOND]], label %[[FREE:.*]], label %[[END:.*]]
; CHECK: [[FREE]]:
-; CHECK-NEXT: [[TMP33:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP33:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: [[TMP34:%.*]] = ptrtoint ptr [[TMP33]] to i64
; CHECK-NEXT: [[TMP35:%.*]] = ptrtoint ptr addrspace(1) [[TMP20]] to i64
; CHECK-NEXT: call void @__asan_free_impl(i64 [[TMP35]], i64 [[TMP34]])
diff --git a/llvm/test/Instrumentation/ThreadSanitizer/atomic-non-integer.ll b/llvm/test/Instrumentation/ThreadSanitizer/atomic-non-integer.ll
index 015ee2fe711e1..0ca6465814090 100644
--- a/llvm/test/Instrumentation/ThreadSanitizer/atomic-non-integer.ll
+++ b/llvm/test/Instrumentation/ThreadSanitizer/atomic-non-integer.ll
@@ -6,7 +6,7 @@ target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f3
define float @load_float(ptr %fptr) {
; CHECK-LABEL: define float @load_float(
; CHECK-SAME: ptr [[FPTR:%.*]]) {
-; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: call void @__tsan_func_entry(ptr [[TMP1]])
; CHECK-NEXT: [[TMP2:%.*]] = call i32 @__tsan_atomic32_load(ptr [[FPTR]], i32 0)
; CHECK-NEXT: [[TMP3:%.*]] = bitcast i32 [[TMP2]] to float
@@ -20,7 +20,7 @@ define float @load_float(ptr %fptr) {
define double @load_double(ptr %fptr) {
; CHECK-LABEL: define double @load_double(
; CHECK-SAME: ptr [[FPTR:%.*]]) {
-; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: call void @__tsan_func_entry(ptr [[TMP1]])
; CHECK-NEXT: [[TMP2:%.*]] = call i64 @__tsan_atomic64_load(ptr [[FPTR]], i32 0)
; CHECK-NEXT: [[TMP3:%.*]] = bitcast i64 [[TMP2]] to double
@@ -34,7 +34,7 @@ define double @load_double(ptr %fptr) {
define fp128 @load_fp128(ptr %fptr) {
; CHECK-LABEL: define fp128 @load_fp128(
; CHECK-SAME: ptr [[FPTR:%.*]]) {
-; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: call void @__tsan_func_entry(ptr [[TMP1]])
; CHECK-NEXT: [[TMP2:%.*]] = call i128 @__tsan_atomic128_load(ptr [[FPTR]], i32 0)
; CHECK-NEXT: [[TMP3:%.*]] = bitcast i128 [[TMP2]] to fp128
@@ -48,7 +48,7 @@ define fp128 @load_fp128(ptr %fptr) {
define void @store_float(ptr %fptr, float %v) {
; CHECK-LABEL: define void @store_float(
; CHECK-SAME: ptr [[FPTR:%.*]], float [[V:%.*]]) {
-; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: call void @__tsan_func_entry(ptr [[TMP1]])
; CHECK-NEXT: [[TMP2:%.*]] = bitcast float [[V]] to i32
; CHECK-NEXT: call void @__tsan_atomic32_store(ptr [[FPTR]], i32 [[TMP2]], i32 0)
@@ -62,7 +62,7 @@ define void @store_float(ptr %fptr, float %v) {
define void @store_double(ptr %fptr, double %v) {
; CHECK-LABEL: define void @store_double(
; CHECK-SAME: ptr [[FPTR:%.*]], double [[V:%.*]]) {
-; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: call void @__tsan_func_entry(ptr [[TMP1]])
; CHECK-NEXT: [[TMP2:%.*]] = bitcast double [[V]] to i64
; CHECK-NEXT: call void @__tsan_atomic64_store(ptr [[FPTR]], i64 [[TMP2]], i32 0)
@@ -76,7 +76,7 @@ define void @store_double(ptr %fptr, double %v) {
define void @store_fp128(ptr %fptr, fp128 %v) {
; CHECK-LABEL: define void @store_fp128(
; CHECK-SAME: ptr [[FPTR:%.*]], fp128 [[V:%.*]]) {
-; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: call void @__tsan_func_entry(ptr [[TMP1]])
; CHECK-NEXT: [[TMP2:%.*]] = bitcast fp128 [[V]] to i128
; CHECK-NEXT: call void @__tsan_atomic128_store(ptr [[FPTR]], i128 [[TMP2]], i32 0)
diff --git a/llvm/test/Instrumentation/ThreadSanitizer/eh.ll b/llvm/test/Instrumentation/ThreadSanitizer/eh.ll
index 9c08b8f077e94..aec0f8b0aab11 100644
--- a/llvm/test/Instrumentation/ThreadSanitizer/eh.ll
+++ b/llvm/test/Instrumentation/ThreadSanitizer/eh.ll
@@ -10,7 +10,7 @@ declare void @cannot_throw() nounwind
define i32 @func1() sanitize_thread {
; CHECK-EXC-LABEL: define i32 @func1
; CHECK-EXC-SAME: () #[[ATTR1:[0-9]+]] personality ptr @__gcc_personality_v0 {
-; CHECK-EXC-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-EXC-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-EXC-NEXT: call void @__tsan_func_entry(ptr [[TMP1]])
; CHECK-EXC-NEXT: invoke void @can_throw()
; CHECK-EXC-NEXT: to label [[DOTNOEXC:%.*]] unwind label [[TSAN_CLEANUP:%.*]]
@@ -25,7 +25,7 @@ define i32 @func1() sanitize_thread {
;
; CHECK-NOEXC-LABEL: define i32 @func1
; CHECK-NOEXC-SAME: () #[[ATTR1:[0-9]+]] {
-; CHECK-NOEXC-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NOEXC-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NOEXC-NEXT: call void @__tsan_func_entry(ptr [[TMP1]])
; CHECK-NOEXC-NEXT: call void @can_throw()
; CHECK-NOEXC-NEXT: call void @__tsan_func_exit()
@@ -38,7 +38,7 @@ define i32 @func1() sanitize_thread {
define i32 @func2() sanitize_thread {
; CHECK-LABEL: define i32 @func2
; CHECK-SAME: () #[[ATTR1:[0-9]+]] {
-; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: call void @__tsan_func_entry(ptr [[TMP1]])
; CHECK-NEXT: call void @cannot_throw()
; CHECK-NEXT: call void @__tsan_func_exit()
@@ -51,7 +51,7 @@ define i32 @func2() sanitize_thread {
define i32 @func3(ptr %p) sanitize_thread {
; CHECK-LABEL: define i32 @func3
; CHECK-SAME: (ptr [[P:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: call void @__tsan_func_entry(ptr [[TMP1]])
; CHECK-NEXT: call void @__tsan_read4(ptr [[P]])
; CHECK-NEXT: [[A:%.*]] = load i32, ptr [[P]], align 4
@@ -65,7 +65,7 @@ define i32 @func3(ptr %p) sanitize_thread {
define i32 @func4() sanitize_thread nounwind {
; CHECK-LABEL: define i32 @func4
; CHECK-SAME: () #[[ATTR2:[0-9]+]] {
-; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: call void @__tsan_func_entry(ptr [[TMP1]])
; CHECK-NEXT: call void @can_throw()
; CHECK-NEXT: call void @__tsan_func_exit()
diff --git a/llvm/test/Instrumentation/ThreadSanitizer/no_sanitize_thread.ll b/llvm/test/Instrumentation/ThreadSanitizer/no_sanitize_thread.ll
index 965704c24bd8a..9edf0cac3afca 100644
--- a/llvm/test/Instrumentation/ThreadSanitizer/no_sanitize_thread.ll
+++ b/llvm/test/Instrumentation/ThreadSanitizer/no_sanitize_thread.ll
@@ -25,7 +25,7 @@ entry:
; CHECK: define i32 @read_4_bytes_and_call(ptr %a) {
; CHECK-NEXT: entry:
-; CHECK-NEXT: %0 = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: %0 = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: call void @__tsan_func_entry(ptr %0)
; CHECK-NEXT: call void @foo()
; CHECK-NEXT: %tmp1 = load i32, ptr %a, align 4
diff --git a/llvm/test/Instrumentation/ThreadSanitizer/sanitize-thread-no-checking.ll b/llvm/test/Instrumentation/ThreadSanitizer/sanitize-thread-no-checking.ll
index dbee198736569..254103dfd3689 100644
--- a/llvm/test/Instrumentation/ThreadSanitizer/sanitize-thread-no-checking.ll
+++ b/llvm/test/Instrumentation/ThreadSanitizer/sanitize-thread-no-checking.ll
@@ -25,7 +25,7 @@ entry:
; CHECK: define i32 @"\01-[WithCalls dealloc]"(ptr %a)
; CHECK-NEXT: entry:
-; CHECK-NEXT: %0 = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: %0 = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: call void @__tsan_func_entry(ptr %0)
; CHECK-NEXT: call void @__tsan_ignore_thread_begin()
; CHECK-NEXT: %tmp1 = load i32, ptr %a, align 4
diff --git a/llvm/test/Transforms/EntryExitInstrumenter/debug-info.ll b/llvm/test/Transforms/EntryExitInstrumenter/debug-info.ll
index 32a60dd6a4a00..b219d1999c8e8 100644
--- a/llvm/test/Transforms/EntryExitInstrumenter/debug-info.ll
+++ b/llvm/test/Transforms/EntryExitInstrumenter/debug-info.ll
@@ -10,10 +10,10 @@ entry:
ret i32 42, !dbg !12
; CHECK-LABEL: define i32 @f(i32 %x)
-; CHECK: call ptr @llvm.returnaddress(i32 0), !dbg ![[ENTRYLOC:[0-9]+]]
+; CHECK: call ptr @llvm.returnaddress.p0(i32 0), !dbg ![[ENTRYLOC:[0-9]+]]
; CHECK: call void @__cyg_profile_func_enter{{.*}}, !dbg ![[ENTRYLOC]]
-; CHECK: call ptr @llvm.returnaddress(i32 0), !dbg ![[EXITLOC:[0-9]+]]
+; CHECK: call ptr @llvm.returnaddress.p0(i32 0), !dbg ![[EXITLOC:[0-9]+]]
; CHECK: call void @__cyg_profile_func_exit{{.*}}, !dbg ![[EXITLOC]]
; CHECK: ret i32 42, !dbg ![[EXITLOC]]
}
diff --git a/llvm/test/Transforms/EntryExitInstrumenter/mcount-with-frompc.ll b/llvm/test/Transforms/EntryExitInstrumenter/mcount-with-frompc.ll
index 0f8cf5c735453..f72c927c85bd4 100644
--- a/llvm/test/Transforms/EntryExitInstrumenter/mcount-with-frompc.ll
+++ b/llvm/test/Transforms/EntryExitInstrumenter/mcount-with-frompc.ll
@@ -9,7 +9,7 @@
define void @f1() "instrument-function-entry-inlined"="_mcount" {
; CHECK-LABEL: define void @f1() {
-; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: call void @_mcount(ptr [[TMP1]])
; CHECK-NEXT: ret void
;
diff --git a/llvm/test/Transforms/EntryExitInstrumenter/mcount.ll b/llvm/test/Transforms/EntryExitInstrumenter/mcount.ll
index 56ccfb9ed2e7e..658bfd654f5a1 100644
--- a/llvm/test/Transforms/EntryExitInstrumenter/mcount.ll
+++ b/llvm/test/Transforms/EntryExitInstrumenter/mcount.ll
@@ -10,9 +10,9 @@ target triple = "powerpc64le-unknown-linux"
define void @leaf_function() #0 {
; CHECK-LABEL: define void @leaf_function() {
; CHECK-NEXT: call void @mcount()
-; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: call void @__cyg_profile_func_enter(ptr @leaf_function, ptr [[TMP1]])
-; CHECK-NEXT: [[TMP2:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP2:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: call void @__cyg_profile_func_exit(ptr @leaf_function, ptr [[TMP2]])
; CHECK-NEXT: ret void
;
@@ -23,13 +23,13 @@ define void @leaf_function() #0 {
define void @root_function() #0 {
; CHECK-LABEL: define void @root_function() {
; CHECK-NEXT: call void @mcount()
-; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: call void @__cyg_profile_func_enter(ptr @root_function, ptr [[TMP1]])
-; CHECK-NEXT: [[TMP2:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP2:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: call void @__cyg_profile_func_enter(ptr @leaf_function, ptr [[TMP2]])
-; CHECK-NEXT: [[TMP3:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP3:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: call void @__cyg_profile_func_exit(ptr @leaf_function, ptr [[TMP3]])
-; CHECK-NEXT: [[TMP4:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP4:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: call void @__cyg_profile_func_exit(ptr @root_function, ptr [[TMP4]])
; CHECK-NEXT: ret void
;
@@ -100,7 +100,7 @@ define void @f7() #7 {
declare ptr @tailcallee()
define ptr @tailcaller() #8 {
; CHECK-LABEL: define ptr @tailcaller() {
-; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: call void @__cyg_profile_func_exit(ptr @tailcaller, ptr [[TMP1]])
; CHECK-NEXT: [[TMP2:%.*]] = musttail call ptr @tailcallee()
; CHECK-NEXT: ret ptr [[TMP2]]
@@ -110,7 +110,7 @@ define ptr @tailcaller() #8 {
}
define ptr @tailcaller2() #8 {
; CHECK-LABEL: define ptr @tailcaller2() {
-; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress(i32 0)
+; CHECK-NEXT: [[TMP1:%.*]] = call ptr @llvm.returnaddress.p0(i32 0)
; CHECK-NEXT: call void @__cyg_profile_func_exit(ptr @tailcaller2, ptr [[TMP1]])
; CHECK-NEXT: [[TMP2:%.*]] = musttail call ptr @tailcallee()
; CHECK-NEXT: ret ptr [[TMP2]]
diff --git a/llvm/test/Transforms/EntryExitInstrumenter/pre-inliner-instrumentation.ll b/llvm/test/Transforms/EntryExitInstrumenter/pre-inliner-instrumentation.ll
index 89e12a2971290..1ae096c48c4a1 100644
--- a/llvm/test/Transforms/EntryExitInstrumenter/pre-inliner-instrumentation.ll
+++ b/llvm/test/Transforms/EntryExitInstrumenter/pre-inliner-instrumentation.ll
@@ -15,7 +15,7 @@ define void @leaf_function() #0 {
entry:
ret void
; INSTRUMENT-LABEL: entry:
-; INSTRUMENT-NEXT: %0 ={{.*}} call ptr @llvm.returnaddress(i32 0)
+; INSTRUMENT-NEXT: %0 ={{.*}} call ptr @llvm.returnaddress.p0(i32 0)
; INSTRUMENT-NEXT: {{.* call void @__cyg_profile_func_enter\(ptr( nonnull)? @leaf_function, ptr %0\)}}
; NOINSTRUMENT-NOT: {{.*}} call void @__cyg_profile_func_enter
; INSTRUMENT: {{.*}} call void @__cyg_profile_func_exit
@@ -31,7 +31,7 @@ entry:
call void @leaf_function()
ret void
; INSTRUMENT-LABEL: entry:
-; INSTRUMENT-NEXT: %0 ={{.*}} call ptr @llvm.returnaddress(i32 0)
+; INSTRUMENT-NEXT: %0 ={{.*}} call ptr @llvm.returnaddress.p0(i32 0)
; INSTRUMENT-NEXT: {{.*}} call void @__cyg_profile_func_enter(ptr{{( nonnull)?}} @root_function, ptr %0)
; INSTRUMENT: {{.*}} call void @__cyg_profile_func_enter
; INSTRUMENT: {{.*}} call void @__cyg_profile_func_exit
diff --git a/llvm/test/Verifier/LoongArch/intrinsic-immarg.ll b/llvm/test/Verifier/LoongArch/intrinsic-immarg.ll
index 488f77ff55ed4..b8f506477ab57 100644
--- a/llvm/test/Verifier/LoongArch/intrinsic-immarg.ll
+++ b/llvm/test/Verifier/LoongArch/intrinsic-immarg.ll
@@ -14,7 +14,7 @@ define ptr @non_const_depth_frameaddress(i32 %x) nounwind {
define ptr @non_const_depth_returnaddress(i32 %x) nounwind {
; CHECK: immarg operand has non-immediate parameter
; CHECK-NEXT: i32 %x
- ; CHECK-NEXT: %1 = call ptr @llvm.returnaddress(i32 %x)
+ ; CHECK-NEXT: %1 = call ptr @llvm.returnaddress.p0(i32 %x)
%1 = call ptr @llvm.returnaddress(i32 %x)
ret ptr %1
}
diff --git a/llvm/test/Verifier/intrinsic-immarg.ll b/llvm/test/Verifier/intrinsic-immarg.ll
index 6e68dde62afae..c95c7214fde71 100644
--- a/llvm/test/Verifier/intrinsic-immarg.ll
+++ b/llvm/test/Verifier/intrinsic-immarg.ll
@@ -4,7 +4,7 @@ declare ptr @llvm.returnaddress(i32)
define void @return_address(i32 %var) {
; CHECK: immarg operand has non-immediate parameter
; CHECK-NEXT: i32 %var
- ; CHECK-NEXT: %result = call ptr @llvm.returnaddress(i32 %var)
+ ; CHECK-NEXT: %result = call ptr @llvm.returnaddress.p0(i32 %var)
%result = call ptr @llvm.returnaddress(i32 %var)
ret void
}
>From c819623023e46c6d4f255242a92498a2feb5f599 Mon Sep 17 00:00:00 2001
From: Owen Anderson <resistor at mac.com>
Date: Wed, 25 Mar 2026 13:06:46 +0100
Subject: [PATCH 2/2] Get program address space from DataLayout in AMDGPU
---
llvm/lib/Target/AMDGPU/AMDGPUSwLowerLDS.cpp | 9 +++++----
1 file changed, 5 insertions(+), 4 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUSwLowerLDS.cpp b/llvm/lib/Target/AMDGPU/AMDGPUSwLowerLDS.cpp
index 04383855b946b..b012b9e5f9935 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUSwLowerLDS.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUSwLowerLDS.cpp
@@ -774,6 +774,7 @@ void AMDGPUSwLowerLDS::lowerKernelLDSAccesses(Function *Func,
auto *PrevEntryBlock = &Func->getEntryBlock();
SetVector<Instruction *> LDSInstructions;
getLDSMemoryInstructions(Func, LDSInstructions);
+ const DataLayout &DL = M.getDataLayout();
// Create malloc block.
auto *MallocBlock = BasicBlock::Create(Ctx, "Malloc", Func, PrevEntryBlock);
@@ -868,7 +869,7 @@ void AMDGPUSwLowerLDS::lowerKernelLDSAccesses(Function *Func,
// Create a call to malloc function which does device global memory allocation
// with size equals to all LDS global accesses size in this kernel.
Value *ReturnAddress = IRB.CreateIntrinsic(
- Intrinsic::returnaddress, IRB.getPtrTy(AMDGPUAS::FLAT_ADDRESS),
+ Intrinsic::returnaddress, IRB.getPtrTy(DL.getProgramAddressSpace()),
{IRB.getInt32(0)});
FunctionCallee MallocFunc = M.getOrInsertFunction(
StringRef("__asan_malloc_impl"),
@@ -934,9 +935,9 @@ void AMDGPUSwLowerLDS::lowerKernelLDSAccesses(Function *Func,
FunctionCallee AsanFreeFunc = M.getOrInsertFunction(
StringRef("__asan_free_impl"),
FunctionType::get(IRB.getVoidTy(), {Int64Ty, Int64Ty}, false));
- Value *ReturnAddr = IRB.CreateIntrinsic(Intrinsic::returnaddress,
- IRB.getPtrTy(AMDGPUAS::FLAT_ADDRESS),
- IRB.getInt32(0));
+ Value *ReturnAddr = IRB.CreateIntrinsic(
+ Intrinsic::returnaddress, IRB.getPtrTy(DL.getProgramAddressSpace()),
+ IRB.getInt32(0));
Value *RAPToInt = IRB.CreatePtrToInt(ReturnAddr, Int64Ty);
Value *MallocPtrToInt = IRB.CreatePtrToInt(LoadMallocPtr, Int64Ty);
IRB.CreateCall(AsanFreeFunc, {MallocPtrToInt, RAPToInt});
More information about the cfe-commits
mailing list