Changpeng Fang dd5895cc39 AMDGPU: Use the implicit kernargs for code object version 5 ... Summary: Specifically, for trap handling, for targets that do not support getDoorbellID, we load the queue_ptr from the implicit kernarg, and move queue_ptr to s[0:1]. To get aperture bases when targets do not have aperture registers, we load private_base or shared_base directly from the implicit kernarg. In clang, we use implicitarg_ptr + offsets to implement __builtin_amdgcn_workgroup_size_{xyz}. Reviewers: arsenm, sameerds, yaxunl Differential Revision: https://reviews.llvm.org/D120265		2022-03-17 14:12:36 -07:00
..
Inputs	[HIP] Support `-fgpu-default-stream`	2022-02-23 22:28:29 -05:00
address-spaces.cu	[Clang/Test]: Rename enable_noundef_analysis to disable-noundef-analysis and turn it off by default	2022-01-16 18:54:17 +09:00
alias.cu	…
amdgpu-alias-undef-symbols.cu	[HIP] [AlwaysInliner] Disable AlwaysInliner to eliminate undefined symbols	2021-10-18 16:53:15 -06:00
amdgpu-asan.cu	[AMDGPU] replace hostcall module flag with function attribute	2022-02-11 22:51:56 +05:30
amdgpu-atomic-ops.cu	…
amdgpu-code-object-version.cu	[HIP] Emit amdgpu_code_object_version module flag	2022-02-08 21:58:40 -05:00
amdgpu-func-attrs.cu	…
amdgpu-kernel-arg-pointer-type.cu	[AMDGPU] Set noclobber metadata on loads instead of cast to constant	2022-03-07 23:13:02 -08:00
amdgpu-kernel-attrs.cu	…
amdgpu-visibility.cu	…
amdgpu-workgroup-size.cu	AMDGPU: Use the implicit kernargs for code object version 5	2022-03-17 14:12:36 -07:00
atomic-ops.cu	[HIP] Add atomic load, atomic store and atomic cmpxchng_weak builtin support in HIP-clang	2021-11-29 12:07:13 -07:00
atomics-remarks-gfx90a.cu	[Remarks] Emit optimization remarks for atomics generating CAS loop	2021-08-16 14:56:01 -06:00
builtin-mangled-name.cu	…
builtins-amdgcn.cu	[Clang/Test]: Rename enable_noundef_analysis to disable-noundef-analysis and turn it off by default	2022-01-16 18:54:17 +09:00
builtins-unsafe-atomics-gfx90a.cu	[HIP] Relax conditions for address space cast in builtin args	2021-10-15 15:35:52 -06:00
const-var.cu	[CUDA][HIP] Do not treat host var address as constant in device compilation	2022-01-28 16:04:52 -05:00
constexpr-variables.cu	…
convergent.cu	[Clang] [NFC] fix CHECK lines for convergent attribute tests	2021-06-29 00:21:07 +05:30
correctly-rounded-div.cu	…
cuda-builtin-vars.cu	[Clang/Test]: Rename enable_noundef_analysis to disable-noundef-analysis and turn it off by default	2022-01-16 18:54:17 +09:00
cxx-call-kernel.cpp	…
debug-info-address-class.cu	…
debug-info-template.cu	…
deferred-diag.cu	…
dependent-libs.cu	…
device-fun-linkage.cu	…
device-init-fun.cu	…
device-stub.cu	[HIP] Fix -fgpu-rdc for Windows	2021-12-06 16:42:23 -05:00
device-use-host-var.cu	CUDA/HIP: Change device-use-host-var.cu's NOT "external" check to include variable name	2021-06-04 13:10:00 -04:00
device-var-init.cu	…
device-var-linkage.cu	[CUDA] Fix static device variables with -fgpu-rdc	2021-08-25 09:31:22 +02:00
device-vtable.cu	…
dft-func-attr-skip-intrinsic.hip	[Clang] [NFC] fix CHECK lines for convergent attribute tests	2021-06-29 00:21:07 +05:30
filter-decl.cu	…
flush-denormals.cu	…
fp-contract.cu	…
function-overload.cu	…
host-used-device-var.cu	[Internalize] Preserve variables externally initialized.	2021-07-08 10:48:19 -04:00
kernel-amdgcn.cu	…
kernel-args-alignment.cu	[Clang/Test]: Rename enable_noundef_analysis to disable-noundef-analysis and turn it off by default	2022-01-16 18:54:17 +09:00
kernel-args.cu	[Clang/Test]: Rename enable_noundef_analysis to disable-noundef-analysis and turn it off by default	2022-01-16 18:54:17 +09:00
kernel-call.cu	[HIP] Support `-fgpu-default-stream`	2022-02-23 22:28:29 -05:00
kernel-dbg-info.cu	…
kernel-stub-name.cu	[Clang/Test]: Rename enable_noundef_analysis to disable-noundef-analysis and turn it off by default	2022-01-16 18:54:17 +09:00
lambda-reference-var.cu	…
lambda.cu	[Clang/Test]: Rename enable_noundef_analysis to disable-noundef-analysis and turn it off by default	2022-01-16 18:54:17 +09:00
launch-bounds.cu	…
library-builtin.cu	…
link-device-bitcode.cu	…
llvm-used.cu	…
managed-var.cu	[CUDA] Fix static device variables with -fgpu-rdc	2021-08-25 09:31:22 +02:00
ms-linker-options.cu	…
norecurse.cu	…
nothrow.cu	…
openmp-target.cu	…
printf-aggregate.cu	…
printf.cu	…
profile-coverage-mapping.cu	…
propagate-metadata.cu	…
ptx-kernels.cu	…
redux-builtins.cu	[Clang/Test]: Rename enable_noundef_analysis to disable-noundef-analysis and turn it off by default	2022-01-16 18:54:17 +09:00
static-device-var-no-rdc.cu	…
static-device-var-rdc.cu	[CUDA] Fix static device variables with -fgpu-rdc	2021-08-25 09:31:22 +02:00
surface.cu	[Clang/Test]: Rename enable_noundef_analysis to disable-noundef-analysis and turn it off by default	2022-01-16 18:54:17 +09:00
texture.cu	[Clang/Test]: Rename enable_noundef_analysis to disable-noundef-analysis and turn it off by default	2022-01-16 18:54:17 +09:00
types.cu	…
unnamed-types.cu	[Clang/Test]: Rename enable_noundef_analysis to disable-noundef-analysis and turn it off by default	2022-01-16 18:54:17 +09:00
unused-global-var.cu	[Internalize] Preserve variables externally initialized.	2021-07-08 10:48:19 -04:00
usual-deallocators.cu	[Clang/Test]: Rename enable_noundef_analysis to disable-noundef-analysis and turn it off by default	2022-01-16 18:54:17 +09:00
vtbl.cu	[IRGen] Do not overwrite existing attributes in CGCall.	2022-01-20 13:45:19 +00:00