[clang] 85c2bd2 - Prevent adding module flag amdgpu_hostcall multiple times

[Yaxun Liu via cfe-commits]

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();
+}
+