Author: Yaxun (Sam) Liu Date: 2022-01-19T12:52:33-05:00 New Revision: 85c2bd2a0e0e2c1706bbf50203d5bbbeedbbd741
URL: https://github.com/llvm/llvm-project/commit/85c2bd2a0e0e2c1706bbf50203d5bbbeedbbd741 DIFF: https://github.com/llvm/llvm-project/commit/85c2bd2a0e0e2c1706bbf50203d5bbbeedbbd741.diff LOG: 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 Added: clang/test/CodeGenCUDA/amdgpu-asan-printf.cu Modified: clang/lib/CodeGen/CodeGenModule.cpp Removed: ################################################################################ diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index f710b9ad067d4..e91da73d2f03c 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 0000000000000..54dbe5bed4475 --- /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(); +} + _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
