maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

Prevent adding module flag amdgpu_hostcall multiple times 

HIP program with printf call fails to compile with -fsanitize=address
option, because of appending module flag - amdgpu_hostcall twice, one
for printf and one for sanitize option. This patch fixes that issue.

Patch by: Praveen Velliengiri

Reviewed by: Yaxun Liu, Roman Lebedev

Differential Revision: https://reviews.llvm.org/D116216
This commit is contained in:
Yaxun (Sam) Liu 2022-01-19 12:19:06 -05:00
parent 32842c1bc9
commit 85c2bd2a0e
2 changed files with 21 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -565,7 +565,9 @@ void CodeGenModule::Release() {
"__amdgpu_device_library_preserve_asan_functions_ptr", nullptr,
llvm::GlobalVariable::NotThreadLocal);
addCompilerUsedGlobal(Var);
getModule().addModuleFlag(llvm::Module::Override, "amdgpu_hostcall", 1);
if (!getModule().getModuleFlag("amdgpu_hostcall")) {
getModule().addModuleFlag(llvm::Module::Override, "amdgpu_hostcall", 1);
}
}
emitLLVMUsed();

18
clang/test/CodeGenCUDA/amdgpu-asan-printf.cu Normal file
View File

@ -0,0 +1,18 @@
// RUN: %clang_cc1 %s -emit-llvm -o - -triple=amdgcn-amd-amdhsa \
// RUN: -fcuda-is-device -target-cpu gfx906 -fsanitize=address \
// RUN: -O3 -x hip | FileCheck -check-prefixes=MFCHECK %s
// MFCHECK: !llvm.module.flags = !{![[FLAG1:[0-9]+]], ![[FLAG2:[0-9]+]]}
// MFCHECK: ![[FLAG1]] = !{i32 4, !"amdgpu_hostcall", i32 1}
// Test to check hostcall module flag metadata is generated correctly
// when a program has printf call and compiled with -fsanitize=address.
#include "Inputs/cuda.h"
__device__ void non_kernel() {
printf("sanitized device function");
}
__global__ void kernel() {
non_kernel();
}