Page MenuHomePhabricator

[AMDGPU] Add __builtin_amdgcn_workgroup_size_x/y/z
ClosedPublic

Authored by yaxunl on Mar 25 2020, 6:57 AM.

Details

Summary

The main purpose of introducing these builtins is to add a range metadata [1, 1025) on the work group size loaded from dispatch ptr, which cannot be done by source code.

Diff Detail

Event Timeline

yaxunl created this revision.Mar 25 2020, 6:57 AM
arsenm added inline comments.Mar 25 2020, 8:40 AM
clang/lib/CodeGen/CGBuiltin.cpp
13428

Why is this necessary? The builtin always has the same return type?

13435

Comment that this is indexing the hsa_kernel_dispatch_packet sstruct?

13442

I thought I had a patch to include the maximum group size in AMDGPUTargetInfo to avoid hardcoding it, but I guess it was never committed

13443

Also set it's invariant

clang/test/CodeGenOpenCL/builtins-amdgcn.cl
539

Also run in a hip test, or some case where the addrspacecast is needed?

yaxunl updated this revision to Diff 252621.Mar 25 2020, 10:42 AM
yaxunl marked 9 inline comments as done.

Revised by Matt's comments

clang/lib/CodeGen/CGBuiltin.cpp
13428

due to https://github.com/llvm/llvm-project/commit/c65f966d76aa5412920b3f14d199e764135bd5ec

pointers returned by builtin functions are in default address space for HIP.

13435

done

13442

Added getMaxOpenCLWorkGroupSize() to TargetInfo

13443

done

arsenm accepted this revision.Mar 26 2020, 6:50 AM
arsenm added inline comments.
clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu
3

I assume the addrspacecast got optimized out? Should this disable llvm passes?

This revision is now accepted and ready to land.Mar 26 2020, 6:50 AM
yaxunl marked 2 inline comments as done.Mar 26 2020, 10:21 AM
yaxunl added inline comments.
clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu
3

We did not emit addrspacecast here since we only need return the loaded value.

HIP by default uses -O0, therefore no need to disable llvm passes.

This revision was automatically updated to reflect the committed changes.
yaxunl marked an inline comment as done.
Herald added a project: Restricted Project. · View Herald TranscriptMar 27 2020, 10:32 PM