[PATCH] D116216: Prevent adding module flag - amdgpu_hostcall multiple times.

[Yaxun Liu via Phabricator via cfe-commits]

This revision was not accepted when it landed; it landed in state "Needs Review".
This revision was landed with ongoing or failed builds.
This revision was automatically updated to reflect the committed changes.
Closed by commit rG85c2bd2a0e0e: Prevent adding module flag amdgpu_hostcall multiple times (authored by yaxunl).

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D116216/new/

https://reviews.llvm.org/D116216

Files:
  clang/lib/CodeGen/CodeGenModule.cpp
  clang/test/CodeGenCUDA/amdgpu-asan-printf.cu


Index: clang/test/CodeGenCUDA/amdgpu-asan-printf.cu
===================================================================
--- /dev/null
+++ 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();
+}
+
Index: clang/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenModule.cpp
+++ clang/lib/CodeGen/CodeGenModule.cpp
@@ -565,7 +565,9 @@
         "__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();


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D116216.401297.patch
Type: text/x-patch
Size: 1447 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20220119/627a5cc8/attachment.bin>