Skip to content

[AMDGPU] Defaults for missing dimensions in SYCL required wg size #72652

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Closed
wants to merge 2 commits into from
Closed

[AMDGPU] Defaults for missing dimensions in SYCL required wg size #72652

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
c2c9e5e
[AMDGPU] Defaults for missing dimensions in SYCL required wg size
jchlanda Nov 17, 2023
4cfe82c
PR test fixes
jchlanda Dec 4, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
12 changes: 8 additions & 4 deletions llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -173,14 +173,18 @@ std::string MetadataStreamerMsgPackV4::getTypeName(Type *Ty,
}

msgpack::ArrayDocNode
MetadataStreamerMsgPackV4::getWorkGroupDimensions(MDNode *Node) const {
MetadataStreamerMsgPackV4::getWorkGroupDimensions(const Function &Func,
MDNode *Node) const {
auto Dims = HSAMetadataDoc->getArrayNode();
if (Node->getNumOperands() != 3)
if (Node->getNumOperands() != 3 && !Func.hasFnAttribute("sycl-module-id"))
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The attribute check should not be necessary. The metadata should have a standalone interpretation and not depend on any other attributes

return Dims;

for (auto &Op : Node->operands())
Dims.push_back(Dims.getDocument()->getNode(
uint64_t(mdconst::extract<ConstantInt>(Op)->getZExtValue())));
for (unsigned I = Dims.size(); I < 3; ++I)
Dims.push_back(Dims.getDocument()->getNode(1));

return Dims;
}

Expand Down Expand Up @@ -233,9 +237,9 @@ void MetadataStreamerMsgPackV4::emitKernelAttrs(const Function &Func,
msgpack::MapDocNode Kern) {

if (auto Node = Func.getMetadata("reqd_work_group_size"))
Kern[".reqd_workgroup_size"] = getWorkGroupDimensions(Node);
Kern[".reqd_workgroup_size"] = getWorkGroupDimensions(Func, Node);
if (auto Node = Func.getMetadata("work_group_size_hint"))
Kern[".workgroup_size_hint"] = getWorkGroupDimensions(Node);
Kern[".workgroup_size_hint"] = getWorkGroupDimensions(Func, Node);
if (auto Node = Func.getMetadata("vec_type_hint")) {
Kern[".vec_type_hint"] = Kern.getDocument()->getNode(
getTypeName(
Expand Down
3 changes: 2 additions & 1 deletion llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -81,7 +81,8 @@ class MetadataStreamerMsgPackV4 : public MetadataStreamer {

std::string getTypeName(Type *Ty, bool Signed) const;

msgpack::ArrayDocNode getWorkGroupDimensions(MDNode *Node) const;
msgpack::ArrayDocNode getWorkGroupDimensions(const Function &Func,
MDNode *Node) const;

msgpack::MapDocNode getHSAKernelProps(const MachineFunction &MF,
const SIProgramInfo &ProgramInfo,
Expand Down
42 changes: 42 additions & 0 deletions llvm/test/CodeGen/AMDGPU/required_work_group_size_sycl.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,42 @@
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a < %s | FileCheck %s

; Make sure that SYCL kernels with less than 3 dimensions specified in required
; work group size, have those dimensions padded up with 1.

; CHECK-LABEL: .name: sycl_kernel_1dim
; CHECK: .reqd_workgroup_size:
; CHECK-NEXT: - 3
; CHECK-NEXT: - 1
; CHECK-NEXT: - 1
define protected amdgpu_kernel void @sycl_kernel_1dim() #1 !reqd_work_group_size !0 {
entry:
ret void
}

; CHECK-LABEL: .name: sycl_kernel_2dim
; CHECK: .reqd_workgroup_size:
; CHECK-NEXT: - 5
; CHECK-NEXT: - 7
; CHECK-NEXT: - 1
define protected amdgpu_kernel void @sycl_kernel_2dim() #1 !reqd_work_group_size !1 {
entry:
ret void
}

; CHECK-LABEL: .name: sycl_kernel_3dim
; CHECK: .reqd_workgroup_size:
; CHECK-NEXT: - 11
; CHECK-NEXT: - 13
; CHECK-NEXT: - 17
define protected amdgpu_kernel void @sycl_kernel_3dim() #1 !reqd_work_group_size !2 {
entry:
ret void
}

attributes #0 = { nounwind speculatable memory(none) }
attributes #1 = { "sycl-module-id"="reqd_work_group_size_check_exception.cpp" }


!0 = !{i32 3}
!1 = !{i32 5, i32 7}
!2 = !{i32 11, i32 13, i32 17}