Prevent adding module flag - amdgpu_hostcall multiple times.

[pvellien]

Summary

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.

[yaxunl]

LGTM. Thanks.

[lebedev.ri]

test coverage?

[pvellien]

The testcases related to this patch are already added via this patch https://reviews.llvm.org/D112820.

[lebedev.ri]

In D116216#3208096, @pvellien wrote:

The testcases related to this patch are already added via this patch https://reviews.llvm.org/D112820.

If this patch does not result in test changes, then there is no test coverage.

[pvellien]

@yaxunl It would be very much helpful to know how to write test coverage for this particular patch? thanks

[yaxunl]

In D116216#3209335, @pvellien wrote:

@yaxunl It would be very much helpful to know how to write test coverage for this particular patch? thanks

there is a lit test amdgpu-asan.cu. You can add a call of printf to that test and make sure it compiles.

[pvellien]

@lebedev.ri updated with test-cases.

[yaxunl]

I think it will be cleaner to keep the original amdgpu-asan.cu unchanged whereas add amdgpu-asan-printf.cu which tests asan with printf.

[pvellien]

Removed amdgpu-asan-noprintf.cu and added amdgpu-asan-printf.cu testcase.

[yaxunl]

LGTM. Thanks.

[pvellien]

In D116216#3251066, @yaxunl wrote:

LGTM. Thanks.

Could you please commit on my behalf? I don't have a commit access to llvm trunk

clang/lib/CodeGen/CodeGenModule.cpp

         FT, llvm::GlobalValue::ExternalLinkage,
         "__amdgpu_device_library_preserve_asan_functions", &getModule());
     auto *Var = new llvm::GlobalVariable(
         getModule(), FT->getPointerTo(),
         /*isConstant=*/true, llvm::GlobalValue::WeakAnyLinkage, F,
         "__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();
   if (SanStats)
     SanStats->finish();
 
   if (CodeGenOpts.Autolink &&
       (Context.getLangOpts().Modules || !LinkerOptionsMetadata.empty())) {

clang/test/CodeGenCUDA/amdgpu-asan-printf.cu

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