From 85c2bd2a0e0e2c1706bbf50203d5bbbeedbbd741 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Wed, 19 Jan 2022 12:19:06 -0500 Subject: [PATCH] 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 --- clang/lib/CodeGen/CodeGenModule.cpp | 4 +++- clang/test/CodeGenCUDA/amdgpu-asan-printf.cu | 18 ++++++++++++++++++ 2 files changed, 21 insertions(+), 1 deletion(-) create mode 100644 clang/test/CodeGenCUDA/amdgpu-asan-printf.cu diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index f710b9a..e91da73 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -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(); diff --git a/clang/test/CodeGenCUDA/amdgpu-asan-printf.cu b/clang/test/CodeGenCUDA/amdgpu-asan-printf.cu new file mode 100644 index 0000000..54dbe5b --- /dev/null +++ b/clang/test/CodeGenCUDA/amdgpu-asan-printf.cu @@ -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(); +} + -- 2.7.4