[clang] clang/AMDGPU: Restore O3 checks in default-attributes.hip (PR #115238)

[Matt Arsenault via cfe-commits]

https://github.com/arsenm created https://github.com/llvm/llvm-project/pull/115238

These were dropped in b1bcb7ca460fcd317bbc8309e14c8761bf8394e0 to
avoid some bot failures.

>From 3a5d957b5fe0d36df2273693c7c865c39715d192 Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Mon, 15 Jul 2024 11:48:03 +0400
Subject: [PATCH] clang/AMDGPU: Restore O3 checks in default-attributes.hip

These were dropped in b1bcb7ca460fcd317bbc8309e14c8761bf8394e0 to
avoid some bot failures.
---
 clang/test/CodeGenHIP/default-attributes.hip | 30 ++++++++++++++++++++
 1 file changed, 30 insertions(+)

diff --git a/clang/test/CodeGenHIP/default-attributes.hip b/clang/test/CodeGenHIP/default-attributes.hip
index 1b53ebec9b5821..ee16ecd134bfee 100644
--- a/clang/test/CodeGenHIP/default-attributes.hip
+++ b/clang/test/CodeGenHIP/default-attributes.hip
@@ -2,6 +2,9 @@
 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fno-ident -fcuda-is-device \
 // RUN:    -emit-llvm -o - %s | FileCheck -check-prefix=OPTNONE %s
 
+// RUN: %clang_cc1 -O3 -triple amdgcn-amd-amdhsa -x hip -fno-ident -fcuda-is-device \
+// RUN:    -emit-llvm -o - %s | FileCheck -check-prefix=OPT %s
+
 #define __device__ __attribute__((device))
 #define __global__ __attribute__((global))
 
@@ -10,6 +13,10 @@
 // OPTNONE: @llvm.compiler.used = appending addrspace(1) global [1 x ptr] [ptr addrspacecast (ptr addrspace(1) @__hip_cuid_ to ptr)], section "llvm.metadata"
 // OPTNONE: @__oclc_ABI_version = weak_odr hidden local_unnamed_addr addrspace(4) constant i32 500
 //.
+// OPT: @__hip_cuid_ = addrspace(1) global i8 0
+// OPT: @__oclc_ABI_version = weak_odr hidden local_unnamed_addr addrspace(4) constant i32 500
+// OPT: @llvm.compiler.used = appending addrspace(1) global [1 x ptr] [ptr addrspacecast (ptr addrspace(1) @__hip_cuid_ to ptr)], section "llvm.metadata"
+//.
 __device__ void extern_func();
 
 // OPTNONE: Function Attrs: convergent mustprogress noinline nounwind optnone
@@ -19,6 +26,13 @@ __device__ void extern_func();
 // OPTNONE-NEXT:    call void @_Z11extern_funcv() #[[ATTR3:[0-9]+]]
 // OPTNONE-NEXT:    ret void
 //
+// OPT: Function Attrs: convergent mustprogress nounwind
+// OPT-LABEL: define {{[^@]+}}@_Z4funcv
+// OPT-SAME: () local_unnamed_addr #[[ATTR0:[0-9]+]] {
+// OPT-NEXT:  entry:
+// OPT-NEXT:    tail call void @_Z11extern_funcv() #[[ATTR3:[0-9]+]]
+// OPT-NEXT:    ret void
+//
 __device__ void func() {
  extern_func();
 }
@@ -30,6 +44,13 @@ __device__ void func() {
 // OPTNONE-NEXT:    call void @_Z11extern_funcv() #[[ATTR3]]
 // OPTNONE-NEXT:    ret void
 //
+// OPT: Function Attrs: convergent mustprogress norecurse nounwind
+// OPT-LABEL: define {{[^@]+}}@_Z6kernelv
+// OPT-SAME: () local_unnamed_addr #[[ATTR2:[0-9]+]] {
+// OPT-NEXT:  entry:
+// OPT-NEXT:    tail call void @_Z11extern_funcv() #[[ATTR3]]
+// OPT-NEXT:    ret void
+//
 __global__ void kernel() {
  extern_func();
 }
@@ -39,7 +60,16 @@ __global__ void kernel() {
 // OPTNONE: attributes #[[ATTR2]] = { convergent mustprogress noinline norecurse nounwind optnone "amdgpu-flat-work-group-size"="1,1024" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" }
 // OPTNONE: attributes #[[ATTR3]] = { convergent nounwind }
 //.
+// OPT: attributes #[[ATTR0]] = { convergent mustprogress nounwind "amdgpu-waves-per-eu"="4,10" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="false" }
+// OPT: attributes #[[ATTR1:[0-9]+]] = { convergent nounwind "amdgpu-waves-per-eu"="4,10" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="false" }
+// OPT: attributes #[[ATTR2]] = { convergent mustprogress norecurse nounwind "amdgpu-flat-work-group-size"="1,1024" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" }
+// OPT: attributes #[[ATTR3]] = { convergent nounwind }
+//.
 // OPTNONE: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 500}
 // OPTNONE: [[META1:![0-9]+]] = !{i32 1, !"amdgpu_printf_kind", !"hostcall"}
 // OPTNONE: [[META2:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}
 //.
+// OPT: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 500}
+// OPT: [[META1:![0-9]+]] = !{i32 1, !"amdgpu_printf_kind", !"hostcall"}
+// OPT: [[META2:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}
+//.