clang/AMDGPU: Restore O3 checks in default-attributes.hip #115238
Merged
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
These were dropped in b1bcb7c to avoid some bot failures.
|
This stack of pull requests is managed by Graphite. Learn more about stacking. |
|
@llvm/pr-subscribers-clang Author: Matt Arsenault (arsenm) ChangesThese were dropped in b1bcb7c to Full diff: https://github.com/llvm/llvm-project/pull/115238.diff 1 Files Affected:
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}
+//.
|
arsenm
deleted the
users/arsenm/restore-o3-checks-default-attributes-test
branch
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/144/builds/11132 Here is the relevant piece of the build log for the reference |
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/12/builds/9175 Here is the relevant piece of the build log for the reference |
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/3/builds/7326 Here is the relevant piece of the build log for the reference |
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/65/builds/7426 Here is the relevant piece of the build log for the reference |
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/27/builds/1708 Here is the relevant piece of the build log for the reference |
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/133/builds/6457 Here is the relevant piece of the build log for the reference |
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/154/builds/7141 Here is the relevant piece of the build log for the reference |
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/46/builds/7541 Here is the relevant piece of the build log for the reference |
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/190/builds/9077 Here is the relevant piece of the build log for the reference |
TomWeaver18
pushed a commit
that referenced
this pull request
Nov 8, 2024
…15238)" This reverts commit 889d677. Caused the following build bot failures, author has failed to address: https://lab.llvm.org/buildbot/#/builders/144/builds/11132 https://lab.llvm.org/buildbot/#/builders/46/builds/7541 https://lab.llvm.org/buildbot/#/builders/190/builds/9077 Please fix before recommitting.
|
Hello and greetings from the UK. My apologies but I've had to revert this change as it has turned several build-bots red for a good many hours this morning. https://lab.llvm.org/buildbot/#/builders/144/builds/11132 The revert is in the following commit: 4bcd4d8 Warmest regards, |
Groverkss
pushed a commit
to iree-org/llvm-project
that referenced
this pull request
Nov 15, 2024
These were dropped in b1bcb7c to avoid some bot failures.
Groverkss
pushed a commit
to iree-org/llvm-project
that referenced
this pull request
Nov 15, 2024
…vm#115238)" This reverts commit 889d677. Caused the following build bot failures, author has failed to address: https://lab.llvm.org/buildbot/#/builders/144/builds/11132 https://lab.llvm.org/buildbot/#/builders/46/builds/7541 https://lab.llvm.org/buildbot/#/builders/190/builds/9077 Please fix before recommitting.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
These were dropped in b1bcb7c to
avoid some bot failures.