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

[praveen velliengiri via Phabricator via cfe-commits]

pvellien updated this revision to Diff 398519.
pvellien added a comment.
Herald added subscribers: kerbowa, nhaehnle, jvesely.

@lebedev.ri updated with test-cases.


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-noprintf.cu
  clang/test/CodeGenCUDA/amdgpu-asan.cu


Index: clang/test/CodeGenCUDA/amdgpu-asan.cu
===================================================================
--- clang/test/CodeGenCUDA/amdgpu-asan.cu
+++ clang/test/CodeGenCUDA/amdgpu-asan.cu
@@ -30,5 +30,16 @@
 // MFCHECK: !llvm.module.flags = !{![[FLAG1:[0-9]+]], ![[FLAG2:[0-9]+]]}
 // MFCHECK: ![[FLAG1]] = !{i32 4, !"amdgpu_hostcall", i32 1}
 
+// CHECK: !llvm.module.flags = !{![[FLAG1:[0-9]+]], ![[FLAG2:[0-9]+]]}
+// CHECK: ![[FLAG1]] = !{i32 4, !"amdgpu_hostcall", i32 1}
 // CHECK-NOT: @__amdgpu_device_library_preserve_asan_functions
 // CHECK-NOT: @__asan_report_load1
+
+#include "Inputs/cuda.h"
+__device__ void non_kernel() {
+  printf("sanitized device function");
+}
+
+__global__ void kernel() {
+  non_kernel();
+}
Index: clang/test/CodeGenCUDA/amdgpu-asan-noprintf.cu
===================================================================
--- /dev/null
+++ clang/test/CodeGenCUDA/amdgpu-asan-noprintf.cu
@@ -0,0 +1,13 @@
+// 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
+// without a call to printf when compiled with -fsanitize=address.
+#include "Inputs/cuda.h"
+__global__ void sanitize_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();



Index: clang/test/CodeGenCUDA/amdgpu-asan.cu
===================================================================
--- clang/test/CodeGenCUDA/amdgpu-asan.cu
+++ clang/test/CodeGenCUDA/amdgpu-asan.cu
@@ -30,5 +30,16 @@
 // MFCHECK: !llvm.module.flags = !{![[FLAG1:[0-9]+]], ![[FLAG2:[0-9]+]]}
 // MFCHECK: ![[FLAG1]] = !{i32 4, !"amdgpu_hostcall", i32 1}
 
+// CHECK: !llvm.module.flags = !{![[FLAG1:[0-9]+]], ![[FLAG2:[0-9]+]]}
+// CHECK: ![[FLAG1]] = !{i32 4, !"amdgpu_hostcall", i32 1}
 // CHECK-NOT: @__amdgpu_device_library_preserve_asan_functions
 // CHECK-NOT: @__asan_report_load1
+
+#include "Inputs/cuda.h"
+__device__ void non_kernel() {
+  printf("sanitized device function");
+}
+
+__global__ void kernel() {
+  non_kernel();
+}
Index: clang/test/CodeGenCUDA/amdgpu-asan-noprintf.cu
===================================================================
--- /dev/null
+++ clang/test/CodeGenCUDA/amdgpu-asan-noprintf.cu
@@ -0,0 +1,13 @@
+// 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
+// without a call to printf when compiled with -fsanitize=address.
+#include "Inputs/cuda.h"
+__global__ void sanitize_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();

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits