llvm-project/clang/test/CodeGenOpenCL
Austin Kerbow 2db700215a [AMDGPU] Add llvm.amdgcn.sched.barrier intrinsic
Adds an intrinsic/builtin that can be used to fine tune scheduler behavior. If
there is a need to have highly optimized codegen and kernel developers have
knowledge of inter-wave runtime behavior which is unknown to the compiler this
builtin can be used to tune scheduling.

This intrinsic creates a barrier between scheduling regions. The immediate
parameter is a mask to determine the types of instructions that should be
prevented from crossing the sched_barrier. In this initial patch, there are only
two variations. A mask of 0 means that no instructions may be scheduled across
the sched_barrier. A mask of 1 means that non-memory, non-side-effect inducing
instructions may cross the sched_barrier.

Note that this intrinsic is only meant to work with the scheduling passes. Any
other transformations that may move code will not be impacted in the ways
described above.

Reviewed By: rampitec

Differential Revision: https://reviews.llvm.org/D124700
2022-05-11 13:22:51 -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 [test] Add {{.*}} to make ELF tests immune to dso_local/dso_preemptable/(none) differences 2020-12-31 00:27:11 -08:00
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 [test] Add {{.*}} to make ELF tests immune to dso_local/dso_preemptable/(none) differences 2020-12-31 00:27:11 -08:00
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][clang] Definition of gfx11 subtarget 2022-04-29 13:55:56 -04: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] Split dot2-insts feature 2021-03-17 09:42:21 +00: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-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 gfx940 smfmac instructions 2022-03-24 12:40:42 -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.cl [AMDGPU] Add llvm.amdgcn.sched.barrier intrinsic 2022-05-11 13:22:51 -07:00
builtins-f16.cl [test] Add {{.*}} to make tests immune to dso_local/dso_preemptable/(none) differences 2020-12-30 20:52:01 -08:00
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 [OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC) 2022-04-07 12:09:47 +02: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 [test] Add {{.*}} to make ELF tests immune to dso_local/dso_preemptable/(none) differences 2020-12-31 00:27:11 -08:00
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 [OpenCL][NFC] Improve OpenCL test file naming 2021-01-06 14:16:44 +00:00
visibility.cl [test] Add {{.*}} to make ELF tests immune to dso_local/dso_preemptable/(none) differences 2020-12-31 00:27:11 -08:00
vla.cl [test] Add {{.*}} to make ELF tests immune to dso_local/dso_preemptable/(none) differences 2020-12-31 00:27:11 -08:00