Changpeng Fang 5035d20dcb AMDGPU: Implement ds_atomic_async_barrier_arrive_b64/ds_atomic_barrier_arrive_rtn_b64 (#146409) ...

These two instructions are supported by gfx1250. We define the
instructions and implement the corresponding intrinsic and builtin.

Co-authored-by: Stanislav Mekhanoshin <Stanislav.Mekhanoshin@amd.com>

2025-07-01 11:08:49 -07:00

..

2011-04-15-vec-init-from-vec.cl

…

addr-space-struct-arg.cl

clang: Fix broken implicit cast to generic address space (#138863)

2025-05-08 07:51:57 +02:00

address-space-constant-initializers.cl

…

address-spaces-conversions.cl

…

address-spaces-mangling.cl

…

address-spaces.cl

…

amdgcn-automatic-variable.cl

clang: Fix broken implicit cast to generic address space (#138863)

2025-05-08 07:51:57 +02:00

amdgcn-buffer-rsrc-type.cl

[ValueTracking] Return true for AddrSpaceCast in canCreateUndefOrPoison (#144686)

2025-06-24 08:43:47 +08:00

amdgcn-flat-scratch-name.cl

[clang][SPIR-V] Add support for AMDGCN flavoured SPIRV (#89796)

2024-06-07 11:50:23 +01:00

amdgcn-large-globals.cl

…

amdgcn-non-temporal-store.cl

…

amdgpu-abi-struct-arg-byref.cl

clang: Fix broken implicit cast to generic address space (#138863)

2025-05-08 07:51:57 +02:00

amdgpu-abi-struct-coerce.cl

[Clang][OpenCL][AMDGPU] Use byref for aggregate OpenCL kernel arguments (#134892)

2025-04-13 10:17:55 -04:00

amdgpu-alignment.cl

[Clang] Remove 3-element vector load and store special handling (#104661)

2025-01-21 09:18:16 -05:00

amdgpu-attrs.cl

…

amdgpu-call-kernel.cl

[IR] Convert from nocapture to captures(none) (#123181)

2025-01-29 16:56:47 +01:00

amdgpu-calling-conv.cl

…

amdgpu-debug-info-pointer-address-space.cl

…

amdgpu-debug-info-variable-expression.cl

[RemoveDIs] Print IR with debug records by default (#91724)

2024-06-14 15:07:27 +01:00

amdgpu-enqueue-kernel-linking.cl

AMDGPU: Move enqueued block handling into clang (#128519)

2025-03-10 19:54:04 +07:00

amdgpu-enqueue-kernel.cl

clang: Fix broken implicit cast to generic address space (#138863)

2025-05-08 07:51:57 +02:00

amdgpu-env-amdgcn.cl

[AMDGPU] Set AS8 address width to 48 bits

2025-05-19 17:26:05 -07:00

amdgpu-features-illegal.cl

[AMDGPU] Report error in clang if wave32 is requested where unsupported (#97633)

2024-07-09 14:25:58 -07:00

amdgpu-features-readonly.cl

…

amdgpu-features.cl

AMDGPU: Add MC layer support for load transpose instructions for gfx1250 (#146024)

2025-06-26 22:30:31 -07:00

amdgpu-ieee.cl

…

amdgpu-nullptr.cl

clang: Fix broken implicit cast to generic address space (#138863)

2025-05-08 07:51:57 +02:00

amdgpu-printf.cl

[Clang][OpenCL][AMDGPU] Allow a kernel to call another kernel (#115821)

2025-04-08 10:29:30 +05:30

amdgpu-sizeof-alignof.cl

…

arm-integer-dot-product.cl

…

as_type.cl

[ValueTracking] Return true for AddrSpaceCast in canCreateUndefOrPoison (#144686)

2025-06-24 08:43:47 +08:00

atomic-builtins-default-to-device-scope.cl

[IR] Convert from nocapture to captures(none) (#123181)

2025-01-29 16:56:47 +01:00

atomic-ops-libcall.cl

…

atomic-ops.cl

Add clang atomic control options and attribute (#114841)

2025-02-27 10:41:04 -05:00

atomics-cas-remarks-gfx90a.cl

[Clang] [NFC] Fix more -Wreturn-type warnings in tests everywhere (#123470)

2025-02-20 19:49:37 +01:00

atomics-unsafe-hw-remarks-gfx90a.cl

[clang][OpenCL][CodeGen][AMDGPU] Do not use private as the default AS for when generic is available (#112442)

2024-10-22 12:05:48 +01:00

backend-unsupported-warning.ll

…

blocks.cl

clang: Fix broken implicit cast to generic address space (#138863)

2025-05-08 07:51:57 +02:00

bool_cast.cl

[LLVM][IR] Use splat syntax when printing Constant[Data]Vector. (#112548)

2024-11-06 11:53:33 +00:00

builtins-alloca.cl

clang: Fix broken implicit cast to generic address space (#138863)

2025-05-08 07:51:57 +02:00

builtins-amdgcn-ci.cl

…

builtins-amdgcn-cvt-off-f32-i4.cl

[Clang][AMDGPU] Add __builtin_amdgcn_cvt_off_f32_i4 (#133741)

2025-04-02 19:51:40 +02:00

builtins-amdgcn-dl-insts-err-clamp.cl

…

builtins-amdgcn-dl-insts-err.cl

AMDGPU: Add support for v_dot2c_f32_bf16 instruction for gfx950 (#117598)

2024-11-25 19:51:01 -08:00

builtins-amdgcn-dl-insts-gfx11.cl

[AMDGPU] Run DL builtin tests for new GFX (#130054)

2025-03-06 14:24:49 +01:00

builtins-amdgcn-dl-insts-gfx12.cl

…

builtins-amdgcn-dl-insts.cl

…

builtins-amdgcn-flat-address-space.cl

…

builtins-amdgcn-fp8.cl

[AMDGPU][clang] Replace gfx940 and gfx941 with gfx942 in clang (#126762)

2025-02-19 10:11:48 +01:00

builtins-amdgcn-fp-atomics-gfx7-err.cl

…

builtins-amdgcn-fp-atomics-gfx11-err.cl

…

builtins-amdgcn-fp-atomics-gfx90a-err.cl

…

builtins-amdgcn-fp-atomics-gfx908-err.cl

…

builtins-amdgcn-gfx9-4-generic-err.cl

[AMDGPU] Introduce a new generic target gfx9-4-generic (#115190)

2024-11-12 23:11:05 -05:00

builtins-amdgcn-gfx9.cl

…

builtins-amdgcn-gfx10.cl

[AMDGPU] Allow overload of __builtin_amdgcn_mov_dpp8 (#113610)

2024-10-31 02:19:20 -07:00

builtins-amdgcn-gfx11-err.cl

[LLVM][AMDGPU] Add Intrinsic and Builtin for ds_bpermute_fi_b32 (#124616)

2025-01-29 14:04:10 +01:00

builtins-amdgcn-gfx11-param-err.cl

…

builtins-amdgcn-gfx11.cl

[AMDGPU] Add a new target for gfx1153 (#113138)

2024-10-23 12:56:58 +09:00

builtins-amdgcn-gfx12-err.cl

…

builtins-amdgcn-gfx12-param-err.cl

[AMDGPU] Remove unused s_barrier_{init,join,leave} instructions (#129548)

2025-03-04 17:52:43 +01:00

builtins-amdgcn-gfx12-wmma-w32.cl

…

builtins-amdgcn-gfx12-wmma-w64.cl

…

builtins-amdgcn-gfx12.cl

clang: Fix broken implicit cast to generic address space (#138863)

2025-05-08 07:51:57 +02:00

builtins-amdgcn-gfx950-err.cl

[AMDGPU][clang] Replace gfx940 and gfx941 with gfx942 in clang (#126762)

2025-02-19 10:11:48 +01:00

builtins-amdgcn-gfx950-read-tr.cl

[IR] Convert from nocapture to captures(none) (#123181)

2025-01-29 16:56:47 +01:00

builtins-amdgcn-gfx950.cl

[AMDGPU] Add a new amdgcn.load.to.lds intrinsic (#137425)

2025-05-19 07:15:04 -07:00

builtins-amdgcn-gfx1250-async-load-store-lds.cl

AMDGPU: Implement ds_atomic_async_barrier_arrive_b64/ds_atomic_barrier_arrive_rtn_b64 (#146409)

2025-07-01 11:08:49 -07:00

builtins-amdgcn-gfx1250-load-tr.cl

AMDGPU: Implement intrinsic/builtins for gfx1250 load transpose instructions (#146289)

2025-06-29 14:33:31 -07:00

builtins-amdgcn-gfx1250.cl

[AMDGPU] Add support for v_cvt_f16_bf8 on gfx1250 (#146305)

2025-06-30 07:54:55 -04:00

builtins-amdgcn-global-load-lds.cl

[AMDGPU] Use a target feature to enable __builtin_amdgcn_global_load_lds on gfx9/10 (#133055)

2025-04-02 20:00:09 +02:00

builtins-amdgcn-global-load-tr-gfx11-err.cl

AMDGPU: Add back half and bfloat support for global_load_tr16 pats (#99540)

2024-07-18 11:23:35 -07:00

builtins-amdgcn-global-load-tr-gfx12-w32-err.cl

AMDGPU: Add back half and bfloat support for global_load_tr16 pats (#99540)

2024-07-18 11:23:35 -07:00

builtins-amdgcn-global-load-tr-gfx12-w64-err.cl

AMDGPU: Add back half and bfloat support for global_load_tr16 pats (#99540)

2024-07-18 11:23:35 -07:00

builtins-amdgcn-global-load-tr-w32.cl

AMDGPU: Add back half and bfloat support for global_load_tr16 pats (#99540)

2024-07-18 11:23:35 -07:00

builtins-amdgcn-global-load-tr-w64.cl

AMDGPU: Add back half and bfloat support for global_load_tr16 pats (#99540)

2024-07-18 11:23:35 -07:00

builtins-amdgcn-gws-insts.cl

[AMDGPU][clang] Replace gfx940 and gfx941 with gfx942 in clang (#126762)

2025-02-19 10:11:48 +01:00

builtins-amdgcn-interp.cl

…

builtins-amdgcn-load-to-lds.cl

[AMDGPU] Add a new amdgcn.load.to.lds intrinsic (#137425)

2025-05-19 07:15:04 -07:00

builtins-amdgcn-make-buffer-rsrc.cl

[AMDGPU] Generalize amdgcn.make.buffer.rsrc to fat pointers (#126828)

2025-02-18 14:15:28 -06:00

builtins-amdgcn-mfma.cl

[AMDGPU][clang] Replace gfx940 and gfx941 with gfx942 in clang (#126762)

2025-02-19 10:11:48 +01:00

builtins-amdgcn-raw-buffer-load-lds.cl

[AMDGPU] vmem-to-lds-load-insts incoherence between TargetParser and AMDGPU.td (#135376)

2025-04-11 16:31:04 +02:00

builtins-amdgcn-raw-buffer-load.cl

[AMDGPU] vmem-to-lds-load-insts incoherence between TargetParser and AMDGPU.td (#135376)

2025-04-11 16:31:04 +02:00

builtins-amdgcn-raw-buffer-store.cl

[Clang][AMDGPU] Use unsigned data type for __builtin_amdgcn_raw_buffer_store_* (#99546)

2024-07-18 16:34:59 -04:00

builtins-amdgcn-raytracing.cl

[AMDGPU][Clang] Add builtins for gfx12 ray tracing intrinsics (#135224)

2025-04-11 09:33:32 -04:00

builtins-amdgcn-swmmac-w32.cl

[Clang][AMDGPU] Use 32-bit index for SWMMAC builtins (#129101)

2025-02-27 23:28:48 -05:00

builtins-amdgcn-swmmac-w64.cl

[Clang][AMDGPU] Use 32-bit index for SWMMAC builtins (#129101)

2025-02-27 23:28:48 -05:00

builtins-amdgcn-vi.cl

Fix regression tests with bad FileCheck checks (#140373)

2025-05-22 07:59:57 +03:00

builtins-amdgcn-wave32.cl

…

builtins-amdgcn-wave64.cl

…

builtins-amdgcn-wmma-w32-gfx10-err.cl

…

builtins-amdgcn-wmma-w32-param-err.cl

…

builtins-amdgcn-wmma-w32.cl

…

builtins-amdgcn-wmma-w64-gfx10-err.cl

…

builtins-amdgcn-wmma-w64-param-err.cl

…

builtins-amdgcn-wmma-w64.cl

…

builtins-amdgcn.cl

AMDGPU: Mark workitem ID intrinsics with range attribute (#136196)

2025-04-18 12:27:38 +02:00

builtins-f16.cl

Emit constrained atan2 intrinsic for clang builtin (#113636)

2024-11-12 13:34:29 -08:00

builtins-fp-atomics-gfx8.cl

[clang][OpenCL][CodeGen][AMDGPU] Do not use private as the default AS for when generic is available (#112442)

2024-10-22 12:05:48 +01:00

builtins-fp-atomics-gfx12.cl

[clang][OpenCL][CodeGen][AMDGPU] Do not use private as the default AS for when generic is available (#112442)

2024-10-22 12:05:48 +01:00

builtins-fp-atomics-gfx90a.cl

[clang][OpenCL][CodeGen][AMDGPU] Do not use private as the default AS for when generic is available (#112442)

2024-10-22 12:05:48 +01:00

builtins-fp-atomics-gfx942.cl

[AMDGPU][clang] Replace gfx940 and gfx941 with gfx942 in clang (#126762)

2025-02-19 10:11:48 +01:00

builtins-fp-atomics-gfx1030.cl

…

builtins-generic-amdgcn.cl

…

builtins-r600.cl

AMDGPU: Mark workitem ID intrinsics with range attribute (#136196)

2025-04-18 12:27:38 +02:00

builtins.cl

…

byval.cl

…

cast_image.cl

…

check-atomic-alignment.cl

[clang][CodeGen] Fix crash on non-natural type in CheckAtomicAlignment (#141053)

2025-05-22 16:45:46 +02:00

cl20-device-side-enqueue-attributes.cl

[Clang][OpenCL][AMDGPU] Allow a kernel to call another kernel (#115821)

2025-04-08 10:29:30 +05:30

cl20-device-side-enqueue.cl

[Clang][OpenCL][AMDGPU] OpenCL Kernel stubs should be assigned alwaysinline attribute (#137769)

2025-05-07 15:42:23 +05:30

cl-strict-aliasing.cl

…

cl-uniform-wg-size.cl

[Clang][OpenCL][AMDGPU] OpenCL Kernel stubs should be assigned alwaysinline attribute (#137769)

2025-05-07 15:42:23 +05:30

const-str-array-decay.cl

…

constant-addr-space-globals.cl

…

convergent.cl

[Clang][OpenCL][AMDGPU] OpenCL Kernel stubs should be assigned alwaysinline attribute (#137769)

2025-05-07 15:42:23 +05:30

debug-info-programming-language.cl

…

enqueue-kernel-non-entry-block.cl

[Clang][OpenCL][AMDGPU] OpenCL Kernel stubs should be assigned alwaysinline attribute (#137769)

2025-05-07 15:42:23 +05:30

event_t.cl

…

ext-int-shift.cl

…

ext-vector-shuffle.cl

…

fdeclare-opencl-builtins.cl

[Clang][OpenCL] Fix wait_for_event argument address space with -fdeclare-opencl-builtins (#134598)

2025-04-07 10:44:37 +02:00

fpmath.cl

…

func-call-dbg-loc.cl

…

half.cl

…

images.cl

…

implicit-addrspacecast-function-parameter.cl

clang: Fix broken implicit cast to generic address space (#138863)

2025-05-08 07:51:57 +02:00

inline-asm-amdgcn.cl

[clang][SPIR-V] Add support for AMDGCN flavoured SPIRV (#89796)

2024-06-07 11:50:23 +01:00

intel-subgroups-avc-ext-types.cl

…

kernel-arg-info-single-as.cl

…

kernel-arg-info.cl

[Clang][OpenCL][AMDGPU] Allow a kernel to call another kernel (#115821)

2025-04-08 10:29:30 +05:30

kernel-attributes.cl

[OpenCL] Add test for __kernel_exec macro (#138287)

2025-05-05 09:52:51 +02:00

kernel-metadata.cl

…

kernel-param-alignment.cl

[IR] Convert from nocapture to captures(none) (#123181)

2025-01-29 16:56:47 +01:00

kernels-have-spir-cc-by-default.cl

[IR] Convert from nocapture to captures(none) (#123181)

2025-01-29 16:56:47 +01:00

lifetime.cl

…

local-initializer-undef.cl

…

local.cl

…

logical-ops.cl

[LLVM][IR] Use splat syntax when printing Constant[Data]Vector. (#112548)

2024-11-06 11:53:33 +00:00

memcpy.cl

…

no-half.cl

…

no-signed-zeros.cl

…

norecurse.cl

…

null_queue.cl

…

numbered-address-space.cl

…

opencl_types.cl

AMDGPU: Fix libcall recognition of image array types (#119832)

2024-12-16 15:04:53 +09:00

opencl-kernel-call.cl

[Clang][OpenCL][AMDGPU] Add tests for optnone attribute assigned to OpenCL Kernels (#138849)

2025-05-07 18:09:13 +05:30

overload.cl

…

partial_initializer.cl

[LLVM][IR] Use splat syntax when printing Constant[Data]Vector. (#112548)

2024-11-06 11:53:33 +00:00

pipe_builtin.cl

[OpenCL] Restore addrspacecast for pipe builtins (#112514)

2024-10-16 13:58:12 +02:00

pipe_types_mangling.cl

…

pipe_types.cl

…

preserve_vec3.cl

[IR] Convert from nocapture to captures(none) (#123181)

2025-01-29 16:56:47 +01:00

printf.cl

…

private-array-initialization.cl

…

ptx-calls.cl

[NVPTX] Switch front-ends and tests to ptx_kernel cc (#120806)

2025-01-07 18:24:50 -08:00

ptx-kernels.cl

[NVPTX] Switch front-ends and tests to ptx_kernel cc (#120806)

2025-01-07 18:24:50 -08:00

reflect.cl

[Clang][OpenCL][AMDGPU] Allow a kernel to call another kernel (#115821)

2025-04-08 10:29:30 +05:30

relaxed-fpmath.cl

Remove FiniteMathOnly and use only NoHonorINFs and NoHonorNANs. (#97342)

2024-07-26 08:16:38 -04:00

sampler.cl

[Clang][OpenCL][AMDGPU] Allow a kernel to call another kernel (#115821)

2025-04-08 10:29:30 +05:30

shifts.cl

[LLVM][IR] Use splat syntax when printing Constant[Data]Vector. (#112548)

2024-11-06 11:53:33 +00:00

single-precision-constant.cl

…

size_t.cl

…

spir32_target.cl

[Clang][CodeGen][OpenCL] Add -fno-delete-null-pointer-checks to avoid UB. NFC. (#135602)

2025-04-17 14:29:09 +08:00

spir64_target.cl

[Clang][CodeGen][OpenCL] Add -fno-delete-null-pointer-checks to avoid UB. NFC. (#135602)

2025-04-17 14:29:09 +08:00

spir_version.cl

…

spir-calling-conv.cl

[Clang][OpenCL][AMDGPU] Allow a kernel to call another kernel (#115821)

2025-04-08 10:29:30 +05:30

spir-debug-info-pointer-address-space.cl

…

spirv_target.cl

[Clang][CodeGen][OpenCL] Add -fno-delete-null-pointer-checks to avoid UB. NFC. (#135602)

2025-04-17 14:29:09 +08:00

sqrt-fpmath.cl

…

str_literals.cl

…

to_addr_builtin.cl

…

unroll-hint.cl

…

vector_literals.cl

[LLVM][IR] Use splat syntax when printing Constant[Data]Vector. (#112548)

2024-11-06 11:53:33 +00:00

vector_logops.cl

…

vector_odd.cl

…

vector_shufflevector.cl

…

vectorLoadStore.cl

…

visibility.cl

[Clang][OpenCL][AMDGPU] Allow a kernel to call another kernel (#115821)

2025-04-08 10:29:30 +05:30

vla.cl

…