clang/AMDGPU: Force disable block enqueue arguments for HIP

This is a dirty, dirty hack to workaround bot failures at
-O0. Currently these fields are only used by OpenCL features and
evidently the HIP runtime isn't expecting to see them in HIP
programs. The code objects should be language agnostic, so just force
optimize these out until the runtime is fixed.

diff --git a/clang/lib/CodeGen/TargetInfo.cpp b/clang/lib/CodeGen/TargetInfo.cpp
index aec170ae557066cdeafa6e45ffc7421d0e8e1cbb..ee8852903eda52de14dad408558716974b8a7efa 100644 (file)
--- a/clang/lib/CodeGen/TargetInfo.cpp
+++ b/clang/lib/CodeGen/TargetInfo.cpp
@@ -9520,6 +9520,15 @@ void AMDGPUTargetCodeGenInfo::setFunctionDeclAttributes(
     if (NumVGPR != 0)
       F->addFnAttr("amdgpu-num-vgpr", llvm::utostr(NumVGPR));
   }
+
+  if (IsHIPKernel) {
+    // FIXME: This is a dirty, dirty hack to fix bot failures at -O0 and should
+    // be removed. The HIP runtime currently fails to handle the case where one
+    // of these fields fails to optimize out. The runtime should tolerate all
+    // requested implicit inputs regardless of language.
+    F->addFnAttr("amdgpu-no-default-queue");
+    F->addFnAttr("amdgpu-no-completion-action");
+  }
 }
 
 void AMDGPUTargetCodeGenInfo::setTargetAttributes(

diff --git a/clang/test/CodeGenHIP/default-attributes.hip b/clang/test/CodeGenHIP/default-attributes.hip
new file mode 100644 (file)
index 0000000..b4f4a62
--- /dev/null
+++ b/clang/test/CodeGenHIP/default-attributes.hip
@@ -0,0 +1,47 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --check-attributes --check-globals
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \
+// RUN:    -emit-llvm -o - %s | FileCheck -check-prefix=OPTNONE %s
+
+// RUN: %clang_cc1 -O3 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \
+// RUN:    -emit-llvm -o - %s | FileCheck -check-prefix=OPT %s
+
+#define __device__ __attribute__((device))
+#define __global__ __attribute__((global))
+
+// OPTNONE: Function Attrs: convergent mustprogress noinline nounwind optnone
+// OPTNONE-LABEL: define {{[^@]+}}@_Z4funcv
+// OPTNONE-SAME: () #[[ATTR0:[0-9]+]] {
+// OPTNONE-NEXT:  entry:
+// OPTNONE-NEXT:    ret void
+//
+// OPT: Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(none)
+// OPT-LABEL: define {{[^@]+}}@_Z4funcv
+// OPT-SAME: () local_unnamed_addr #[[ATTR0:[0-9]+]] {
+// OPT-NEXT:  entry:
+// OPT-NEXT:    ret void
+//
+__device__ void func() {
+
+}
+
+// OPTNONE: Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
+// OPTNONE-LABEL: define {{[^@]+}}@_Z6kernelv
+// OPTNONE-SAME: () #[[ATTR1:[0-9]+]] {
+// OPTNONE-NEXT:  entry:
+// OPTNONE-NEXT:    ret void
+//
+// OPT: Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(none)
+// OPT-LABEL: define {{[^@]+}}@_Z6kernelv
+// OPT-SAME: () local_unnamed_addr #[[ATTR1:[0-9]+]] {
+// OPT-NEXT:  entry:
+// OPT-NEXT:    ret void
+//
+__global__ void kernel() {
+
+}
+//.
+// OPTNONE: attributes #0 = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// OPTNONE: attributes #1 = { convergent mustprogress noinline norecurse nounwind optnone "amdgpu-flat-work-group-size"="1,1024" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" }
+//.
+// OPT: attributes #0 = { mustprogress nofree norecurse nosync nounwind willreturn memory(none) "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// OPT: attributes #1 = { mustprogress nofree norecurse nosync nounwind willreturn memory(none) "amdgpu-flat-work-group-size"="1,1024" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" }