OSchip
/
llvm-project
mirror of https://github.com/THU-DSP-LAB/llvm-project.git
13
0
Fork 6
Code Issues Packages Projects Releases Wiki Activity
llvm-project/clang/test/CodeGenOpenCL
History
Austin Kerbow f5b21680d1 [AMDGPU] Add amdgcn_sched_group_barrier builtin 
This builtin allows the creation of custom scheduling pipelines on a per-region
basis. Like the sched_barrier builtin this is intended to be used either for
testing, in situations where the default scheduler heuristics cannot be
improved, or in critical kernels where users are trying to get performance that
is close to handwritten assembly. Obviously using these builtins will require
extra work from the kernel writer to maintain the desired behavior.

The builtin can be used to create groups of instructions called "scheduling
groups" where ordering between the groups is enforced by the scheduler.
__builtin_amdgcn_sched_group_barrier takes three parameters. The first parameter
is a mask that determines the types of instructions that you would like to
synchronize around and add to a scheduling group. These instructions will be
selected from the bottom up starting from the sched_group_barrier's location
during instruction scheduling. The second parameter is the number of matching
instructions that will be associated with this sched_group_barrier. The third
parameter is an identifier which is used to describe what other
sched_group_barriers should be synchronized with. Note that multiple
sched_group_barriers must be added in order for them to be useful since they
only synchronize with other sched_group_barriers. Only "scheduling groups" with
a matching third parameter will have any enforced ordering between them.

As an example, the code below tries to create a pipeline of 1 VMEM_READ
instruction followed by 1 VALU instruction followed by 5 MFMA instructions...
// 1 VMEM_READ
__builtin_amdgcn_sched_group_barrier(32, 1, 0)
// 1 VALU
__builtin_amdgcn_sched_group_barrier(2, 1, 0)
// 5 MFMA
__builtin_amdgcn_sched_group_barrier(8, 5, 0)
// 1 VMEM_READ
__builtin_amdgcn_sched_group_barrier(32, 1, 0)
// 3 VALU
__builtin_amdgcn_sched_group_barrier(2, 3, 0)
// 2 VMEM_WRITE
__builtin_amdgcn_sched_group_barrier(64, 2, 0)

Reviewed By: jrbyrnes

Differential Revision: https://reviews.llvm.org/D128158
 		2022-07-28 10:43:14 -07:00
..
2011-04-15-vec-init-from-vec.cl
…
addr-space-struct-arg.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
address-space-constant-initializers.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
address-spaces-conversions.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
address-spaces-mangling.cl [OpenCL] Add support of __opencl_c_device_enqueue feature macro. 2022-01-27 14:25:59 +03:00
address-spaces.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
amdgcn-automatic-variable.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
amdgcn-flat-scratch-name.cl
…
amdgcn-large-globals.cl
…
amdgcn-non-temporal-store.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
amdgpu-abi-struct-coerce.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
amdgpu-alignment.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
amdgpu-attrs.cl clang/AMDGPU: Don't set implicit arg attribute to default size 2022-01-14 18:43:30 -05:00
amdgpu-call-kernel.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
amdgpu-calling-conv.cl
…
amdgpu-debug-info-pointer-address-space.cl
…
amdgpu-debug-info-variable-expression.cl [Clang] Add -no-opaque-pointers to more tests (NFC) 2022-04-07 12:53:29 +02:00
amdgpu-enqueue-kernel.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
amdgpu-env-amdgcn.cl
…
amdgpu-features.cl [AMDGPU] Support for gfx940 fp8 conversions 2022-07-18 11:48:43 -07:00
amdgpu-ieee.cl [AMDGPU] Add options -mamdgpu-ieee -mno-amdgpu-ieee 2021-05-01 09:02:55 -04:00
amdgpu-nullptr.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
amdgpu-printf.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
amdgpu-sizeof-alignof.cl [OpenCL] Add support of __opencl_c_generic_address_space feature macro 2021-07-13 13:14:10 +03:00
arm-integer-dot-product.cl [OpenCL] Remove pragma requirement from Arm dot extension. 2021-05-12 16:25:33 +01:00
as_type.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
atomic-ops-libcall.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
atomic-ops.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
atomics-cas-remarks-gfx90a.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
atomics-unsafe-hw-remarks-gfx90a.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
backend-unsupported-warning.ll
…
blocks.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
bool_cast.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
builtins-amdgcn-ci.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
builtins-amdgcn-dl-insts-err-clamp.cl
…
builtins-amdgcn-dl-insts-err.cl [AMDGPU] gfx11 new dot instruction codegen support 2022-06-16 14:19:34 -04:00
builtins-amdgcn-dl-insts-gfx11.cl [AMDGPU] gfx11 new dot instruction codegen support 2022-06-16 14:19:34 -04:00
builtins-amdgcn-dl-insts.cl
…
builtins-amdgcn-fp-atomics-gfx7-err.cl [clang] Add clang builtins support for gfx90a 2021-08-05 02:08:06 -06:00
builtins-amdgcn-fp-atomics-gfx90a-err.cl [AMDGPU] new gfx940 fp atomics 2022-03-07 12:32:02 -08:00
builtins-amdgcn-fp-atomics-gfx908-err.cl [clang] Add clang builtins support for gfx90a 2021-08-05 02:08:06 -06:00
builtins-amdgcn-fp8.cl [AMDGPU] Support for gfx940 fp8 conversions 2022-07-18 11:48:43 -07:00
builtins-amdgcn-gfx9.cl [HIP] Add __builtin_amdgcn_groupstaticsize 2021-05-13 15:50:08 +00:00
builtins-amdgcn-gfx10.cl [HIP] Add __builtin_amdgcn_groupstaticsize 2021-05-13 15:50:08 +00:00
builtins-amdgcn-interp.cl
…
builtins-amdgcn-mfma.cl [AMDGPU] Support for gfx940 fp8 smfmac 2022-07-18 12:12:41 -07:00
builtins-amdgcn-raytracing.cl [AMDGPU] Change llvm.amdgcn.image.bvh.intersect.ray to take vec3 args 2021-12-04 10:32:11 +00:00
builtins-amdgcn-vi.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
builtins-amdgcn-wmma-w32.cl [AMDGPU] Add WMMA clang builtins 2022-07-01 08:55:25 +02:00
builtins-amdgcn-wmma-w64.cl [AMDGPU] Add WMMA clang builtins 2022-07-01 08:55:25 +02:00
builtins-amdgcn.cl [AMDGPU] Add amdgcn_sched_group_barrier builtin 2022-07-28 10:43:14 -07:00
builtins-f16.cl
…
builtins-fp-atomics-gfx8.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
builtins-fp-atomics-gfx90a.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
builtins-fp-atomics-gfx940.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
builtins-fp-atomics-gfx1030.cl [clang] Add clang builtins support for gfx90a 2021-08-05 02:08:06 -06:00
builtins-generic-amdgcn.cl Strengthen -Wint-conversion to default to an error 2022-07-22 15:24:54 -04:00
builtins-r600.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
builtins.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
byval.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
cast_image.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
cl-strict-aliasing.cl
…
cl-uniform-wg-size.cl
…
cl20-device-side-enqueue.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
const-str-array-decay.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
constant-addr-space-globals.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
convergent.cl [Clang/Test]: Rename enable_noundef_analysis to disable-noundef-analysis and turn it off by default 2022-01-16 18:54:17 +09:00
debug-info-programming-language.cl [OpenCL] Use DW_LANG_OpenCL language tag for OpenCL C 2021-06-25 11:48:42 +01:00
enqueue-kernel-non-entry-block.cl
…
event_t.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
ext-int-shift.cl [Tests] Use %clang_cc1 instead of %clang -cc1 in codegen tests (NFC) 2022-04-05 13:21:44 +02:00
ext-vector-shuffle.cl
…
fdeclare-opencl-builtins.cl [OpenCL] Make generic addrspace optional for -fdeclare-opencl-builtins 2022-01-31 10:21:05 +00:00
fpmath.cl [Clang/Test]: Rename enable_noundef_analysis to disable-noundef-analysis and turn it off by default 2022-01-16 18:54:17 +09:00
func-call-dbg-loc.cl A significant number of our tests in C accidentally use functions 2022-02-13 08:03:40 -05:00
half.cl [Clang/Test]: Rename enable_noundef_analysis to disable-noundef-analysis and turn it off by default 2022-01-16 18:54:17 +09:00
images.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
inline-asm-amdgcn.cl
…
intel-subgroups-avc-ext-types.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
kernel-arg-info-single-as.cl
…
kernel-arg-info.cl [OpenCL] Fix types with signed prefix in arginfo metadata. 2021-02-09 15:13:19 +00:00
kernel-attributes.cl
…
kernel-metadata.cl
…
kernel-param-alignment.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
kernels-have-spir-cc-by-default.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
lifetime.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
local-initializer-undef.cl
…
local.cl
…
logical-ops.cl
…
memcpy.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
no-half.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
no-signed-zeros.cl Reduce the number of attributes attached to each function 2021-02-16 16:19:54 +01:00
norecurse.cl
…
null_queue.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
numbered-address-space.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
opencl_types.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
overload.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
partial_initializer.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
pipe_builtin.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
pipe_types.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
pipe_types_mangling.cl [OpenCL] Add support of __opencl_c_device_enqueue feature macro. 2022-01-27 14:25:59 +03:00
preserve_vec3.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
printf.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
private-array-initialization.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
ptx-calls.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
ptx-kernels.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
relaxed-fpmath.cl Reduce the number of attributes attached to each function 2021-02-16 16:19:54 +01:00
sampler.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
shifts.cl
…
single-precision-constant.cl
…
size_t.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
spir-calling-conv.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
spir-debug-info-pointer-address-space.cl [OpenCL] Add DWARF address spaces mapping for SPIR 2021-06-04 18:10:54 +01:00
spir32_target.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
spir64_target.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
spir_version.cl [OpenCL] Change default standard version to CL1.2 2021-07-26 15:04:34 +01:00
spirv_target.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
str_literals.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
to_addr_builtin.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
unroll-hint.cl
…
vectorLoadStore.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
vector_literals.cl [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02:00
vector_logops.cl
…
vector_odd.cl
…
vector_shufflevector.cl
…
visibility.cl
…
vla.cl
…