shylie/llvm-project
1
0
Fork 0
Code Issues Pull Requests Actions 6 Packages Projects Releases Wiki Activity

[CHERI] Allow @llvm.returnaddress to return a pointer in any address space. (#188464)

Browse Source 
Clang now constructs calls to it using the default program address space from the DataLayout.

Co-authored-by: Alex Richardson <alexrichardson@google.com>
This commit is contained in:
Owen Anderson 2026-03-25 14:59:38 +01:00 committed by GitHub
parent 6b6d1573a8
commit ca9ac0e24a
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
56 changed files with 147 additions and 128 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
clang/lib/CodeGen/CGBuiltin.cpp
View File

@ -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: {

2
clang/lib/CodeGen/CodeGenModule.cpp
View File

@ -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();

3
clang/lib/CodeGen/CodeGenTypeCache.h
View File

@ -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;

2
clang/test/CodeGen/ms-intrinsics.c
View File

@ -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__)

4
clang/test/Headers/hip-header.hip
View File

@ -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);

2
llvm/include/llvm/IR/Intrinsics.td
View File

@ -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],

11
llvm/lib/Target/AMDGPU/AMDGPUSwLowerLDS.cpp
View File

@ -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);
@ -867,8 +868,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(DL.getProgramAddressSpace()),
{IRB.getInt32(0)});
FunctionCallee MallocFunc = M.getOrInsertFunction(
StringRef("__asan_malloc_impl"),
FunctionType::get(Int64Ty, {Int64Ty, Int64Ty}, false));
@ -933,8 +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.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});

6
llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp
View File

@ -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);

12
llvm/lib/Transforms/Utils/EntryExitInstrumenter.cpp
View File

@ -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);

2
llvm/test/Bitcode/compatibility-3.6.ll
View File

@ -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)

2
llvm/test/Bitcode/compatibility-3.7.ll
View File

@ -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)

2
llvm/test/Bitcode/compatibility-3.8.ll
View File

@ -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)

2
llvm/test/Bitcode/compatibility-3.9.ll
View File

@ -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)

2
llvm/test/Bitcode/compatibility-4.0.ll
View File

@ -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)

2
llvm/test/Bitcode/compatibility-5.0.ll
View File

@ -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)

2
llvm/test/Bitcode/compatibility-6.0.ll
View File

@ -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)

2
llvm/test/Bitcode/compatibility.ll
View File

@ -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)

4
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-dynamic-indirect-access-asan.ll
View File

@ -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]])

4
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-dynamic-indirect-access.ll
View File

@ -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]])

4
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-dynamic-lds-test-asan.ll
View File

@ -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]])

4
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-dynamic-lds-test.ll
View File

@ -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]])

8
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-lower-all.ll
View File

@ -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]])

8
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-multi-static-dynamic-indirect-access-asan.ll
View File

@ -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]])

8
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-multi-static-dynamic-indirect-access.ll
View File

@ -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]])

4
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-multiple-blocks-return-asan.ll
View File

@ -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]])

5
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-multiple-blocks-return.ll
View File

@ -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}
;.

4
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-non-kernel-declaration.ll
View File

@ -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]])

4
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-dynamic-indirect-access-asan.ll
View File

@ -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]])

4
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-dynamic-indirect-access.ll
View File

@ -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]])

4
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-dynamic-lds-test-asan.ll
View File

@ -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]])

5
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-dynamic-lds-test.ll
View File

@ -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}
;.

4
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-asan.ll
View File

@ -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]])

4
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-function-param-asan.ll
View File

@ -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]])

4
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-function-param.ll
View File

@ -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]])

4
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-lower-all.ll
View File

@ -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]])

17
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-nested-asan.ll
View File

@ -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]])

16
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-nested.ll
View File

@ -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]])

4
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access-no-kernel-lds-id.ll
View File

@ -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]])

4
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-indirect-access.ll
View File

@ -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]])

4
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-O0.ll
View File

@ -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]])

4
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-no-heap-ptr.ll
View File

@ -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]])

4
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-test-asan.ll
View File

@ -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]])

4
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-test-atomic-cmpxchg-asan.ll
View File

@ -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]])

4
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-test-atomicrmw-asan.ll
View File

@ -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]])

4
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-test.ll
View File

@ -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]])

4
llvm/test/CodeGen/AMDGPU/amdgpu-sw-lower-lds-static-lds-vector-ptrs.ll
View File

@ -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]])

12
llvm/test/Instrumentation/ThreadSanitizer/atomic-non-integer.ll
View File

@ -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)

10
llvm/test/Instrumentation/ThreadSanitizer/eh.ll
View File

@ -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()

2
llvm/test/Instrumentation/ThreadSanitizer/no_sanitize_thread.ll
View File

@ -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

2
llvm/test/Instrumentation/ThreadSanitizer/sanitize-thread-no-checking.ll
View File

@ -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

4
llvm/test/Transforms/EntryExitInstrumenter/debug-info.ll
View File

@ -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]]
}

2
llvm/test/Transforms/EntryExitInstrumenter/mcount-with-frompc.ll
View File

@ -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
;

16
llvm/test/Transforms/EntryExitInstrumenter/mcount.ll
View File

@ -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]]

4
llvm/test/Transforms/EntryExitInstrumenter/pre-inliner-instrumentation.ll
View File

@ -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

2
llvm/test/Verifier/LoongArch/intrinsic-immarg.ll
View File

@ -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
}

2
llvm/test/Verifier/intrinsic-immarg.ll
View File

@ -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
}