Matt Arsenault bfea84946d

clang: Hack around opencl enqueue_block using wrong ABI for aggregrate (#130011) ...

EmitAggExprToLValue started wrapping the temporary alloca in an
addrspacecast
at some point. We take the direct type from this as the pointer argument
for the
runtime function type, but this isn't correct. Technically, we should be
querying
the target's ABI for what IR to produce for this sequence. The
assumption seems to
always have been that this will be indirectly passed with byval (or
byref).

I started working on a patch to go through the ABI handling, but it
seems to
require more time and/or clang expertise than I have at the moment.

2025-03-06 23:13:28 +07:00

..

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

…

addr-space-struct-arg.cl

[clang][CodeGen] sret args should always point to the alloca AS, so use that (#114062)

2025-02-14 11:20:45 +00:00

address-space-constant-initializers.cl

…

address-spaces-conversions.cl

…

address-spaces-mangling.cl

…

address-spaces.cl

…

amdgcn-automatic-variable.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

amdgcn-buffer-rsrc-type.cl

Reapply [CaptureTracking][FunctionAttrs] Add support for CaptureInfo (#125880) (#128020)

2025-02-27 09:38:29 +01: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

[test] %clang_cc1 -emit-llvm: remove redundant -S

2024-05-04 17:00:29 -07:00

amdgcn-non-temporal-store.cl

[test] %clang_cc1 -emit-llvm: remove redundant -S

2024-05-04 17:00:29 -07:00

amdgpu-abi-struct-arg-byref.cl

[clang][CodeGen] sret args should always point to the alloca AS, so use that (#114062)

2025-02-14 11:20:45 +00:00

amdgpu-abi-struct-coerce.cl

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

2025-01-29 16:56:47 +01: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] Adding the amdgpu_num_work_groups function attribute (#79035)

2024-03-12 10:30:39 -07:00

amdgpu-call-kernel.cl

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

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

amdgpu-calling-conv.cl

[test] %clang_cc1 -emit-llvm: remove redundant -S

2024-05-04 17:00:29 -07:00

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.cl

clang: Hack around opencl enqueue_block using wrong ABI for aggregrate (#130011)

2025-03-06 23:13:28 +07:00

amdgpu-env-amdgcn.cl

[AMDGPU] - Add address space for strided buffers (#74471)

2023-12-15 15:49:25 +01: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] Add target feature gws to clang

2023-08-25 11:50:47 -04:00

amdgpu-features.cl

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

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

amdgpu-ieee.cl

…

amdgpu-nullptr.cl

[InstCombine] Infer nusw + nneg -> nuw for getelementptr (#111144)

2024-12-05 14:36:40 +01:00

amdgpu-printf.cl

[clang][CodeGen] Allow memcpy replace with trivial auto var init

2024-03-21 09:55:04 +01:00

amdgpu-sizeof-alignof.cl

…

arm-integer-dot-product.cl

…

as_type.cl

Reapply [CaptureTracking][FunctionAttrs] Add support for CaptureInfo (#125880) (#128020)

2025-02-27 09:38:29 +01: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

[clang][CodeGen] Emit atomic IR in place of optimized libcalls. (#73176)

2024-02-12 09:33:09 -08:00

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][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

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][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-amdgcn-ci.cl

[test] %clang_cc1 -emit-llvm: remove redundant -S

2024-05-04 17:00:29 -07:00

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

[test] %clang_cc1 -emit-llvm: remove redundant -S

2024-05-04 17:00:29 -07:00

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

[test] %clang_cc1 -emit-llvm: remove redundant -S

2024-05-04 17:00:29 -07:00

builtins-amdgcn-dl-insts.cl

[test] %clang_cc1 -emit-llvm: remove redundant -S

2024-05-04 17:00:29 -07:00

builtins-amdgcn-flat-address-space.cl

[test] %clang_cc1 -emit-llvm: remove redundant -S

2024-05-04 17:00:29 -07:00

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

[AMDGPU] Add clang builtin for __builtin_amdgcn_ds_atomic_fadd_v2f16

2023-03-24 16:27:44 +01:00

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

[AMDGPU] Add clang builtin for __builtin_amdgcn_ds_atomic_fadd_v2f16

2023-03-24 16:27:44 +01:00

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

[AMDGPU] Replace target feature for global fadd32

2023-03-28 15:58:30 -06:00

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

[test] %clang_cc1 -emit-llvm: remove redundant -S

2024-05-04 17:00:29 -07:00

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

[test] %clang_cc1 -emit-llvm: remove redundant -S

2024-05-04 17:00:29 -07:00

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

[test] %clang_cc1 -emit-llvm: remove redundant -S

2024-05-04 17:00:29 -07:00

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

[test] %clang_cc1 -emit-llvm: remove redundant -S

2024-05-04 17:00:29 -07:00

builtins-amdgcn-gfx12-wmma-w64.cl

[test] %clang_cc1 -emit-llvm: remove redundant -S

2024-05-04 17:00:29 -07:00

builtins-amdgcn-gfx12.cl

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

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

builtins-amdgcn-gfx942.cl

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

2025-02-19 10:11:48 +01: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: Change bitop3 intrinsic operand to i32 (#118647)

2024-12-04 15:44:04 -05: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

[test] %clang_cc1 -emit-llvm: remove redundant -S

2024-05-04 17:00:29 -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.cl

[Clang][AMDGPU] Add builtins for instrinsic llvm.amdgcn.raw.ptr.buffer.load (#99258)

2024-07-18 15:33:03 -04: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

[test] %clang_cc1: remove redundant actions

2024-05-05 10:46:06 -07: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

[AMDGPU] Relax __builtin_amdgcn_update_dpp sema check (#113341)

2024-10-22 12:32:08 -07:00

builtins-amdgcn-wave32.cl

[test] %clang_cc1 -emit-llvm: remove redundant -S

2024-05-04 17:00:29 -07:00

builtins-amdgcn-wave64.cl

[test] %clang_cc1 -emit-llvm: remove redundant -S

2024-05-04 17:00:29 -07:00

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

[AMDGPU] Check wavefrontsize for GFX11 WMMA builtins (#79980)

2024-02-01 10:49:42 +00:00

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

[AMDGPU] - Add clang builtins for tied WMMA intrinsics (#70669)

2023-11-13 13:23:26 +01:00

builtins-amdgcn-wmma-w32.cl

[test] %clang_cc1 -emit-llvm: remove redundant -S

2024-05-04 17:00:29 -07:00

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

[AMDGPU] Check wavefrontsize for GFX11 WMMA builtins (#79980)

2024-02-01 10:49:42 +00:00

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

[AMDGPU] - Add clang builtins for tied WMMA intrinsics (#70669)

2023-11-13 13:23:26 +01:00

builtins-amdgcn-wmma-w64.cl

[test] %clang_cc1 -emit-llvm: remove redundant -S

2024-05-04 17:00:29 -07:00

builtins-amdgcn.cl

[InstCombine] Infer nusw + nneg -> nuw for getelementptr (#111144)

2024-12-05 14:36:40 +01: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

[test] %clang_cc1 -emit-llvm: remove redundant -S

2024-05-04 17:00:29 -07:00

builtins-r600.cl

[Clang] Swap range and noundef metadata to attribute for intrinsics. (#94851)

2024-06-19 17:23:53 +02:00

builtins.cl

[clang] Convert a few tests to opaque pointers

2023-05-14 21:00:15 +03:00

byval.cl

[clang][AMDGPU]: Don't use byval for struct arguments in function ABI

2023-08-11 16:37:42 -07:00

cast_image.cl

[clang] Convert a few tests to opaque pointers

2023-05-14 21:00:15 +03:00

cl20-device-side-enqueue-attributes.cl

[Clang] Avoid Using byval for ndrange_t when emitting __enqueue_kernel_basic (#116435)

2024-11-15 16:54:29 -05:00

cl20-device-side-enqueue.cl

[Clang] Avoid Using byval for ndrange_t when emitting __enqueue_kernel_basic (#116435)

2024-11-15 16:54:29 -05:00

cl-strict-aliasing.cl

…

cl-uniform-wg-size.cl

[Driver] Add -f[no-]offload-uniform-block

2023-07-27 16:36:02 -04:00

const-str-array-decay.cl

[test] Fix const-str-array-decay.cl failure on PowerPC

2023-05-15 19:56:28 +03:00

constant-addr-space-globals.cl

…

convergent.cl

…

debug-info-programming-language.cl

…

enqueue-kernel-non-entry-block.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

event_t.cl

…

ext-int-shift.cl

…

ext-vector-shuffle.cl

…

fdeclare-opencl-builtins.cl

…

fpmath.cl

clang: Add __builtin_elementwise_sqrt

2023-08-11 19:32:39 -04:00

func-call-dbg-loc.cl

…

half.cl

…

images.cl

…

implicit-addrspacecast-function-parameter.cl

[clang][CodeGen] sret args should always point to the alloca AS, so use that (#114062)

2025-02-14 11:20:45 +00: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

[Clang][SPIR-V] Emit target extension types for OpenCL types on SPIR-V.

2023-03-13 14:20:24 -04:00

kernel-arg-info-single-as.cl

…

kernel-arg-info.cl

…

kernel-attributes.cl

…

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

[clang] Convert a few tests to opaque pointers

2023-05-14 21:00:15 +03:00

no-signed-zeros.cl

…

norecurse.cl

…

null_queue.cl

…

numbered-address-space.cl

[test] %clang_cc1 -emit-llvm: remove redundant -S

2024-05-04 17:00:29 -07:00

opencl_types.cl

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

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

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

[clang] Convert a few tests to opaque pointers

2023-05-14 21:00:15 +03:00

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

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

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

relaxed-fpmath.cl

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

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

sampler.cl

[Clang][SPIR-V] Emit target extension types for OpenCL types on SPIR-V.

2023-03-13 14:20:24 -04:00

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

…

spir64_target.cl

…

spir_version.cl

…

spir-calling-conv.cl

…

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

…

spirv_target.cl

…

sqrt-fpmath.cl

[test] %clang_cc1 -emit-llvm: remove redundant -S

2024-05-04 17:00:29 -07:00

str_literals.cl

…

to_addr_builtin.cl

[clang] Convert a few tests to opaque pointers

2023-05-14 21:00:15 +03:00

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

[Clang] Use poison as base for vector literals

2023-12-19 11:53:18 +01:00

vectorLoadStore.cl

…

visibility.cl

[test] %clang_cc1 -emit-llvm: remove redundant -S

2024-05-04 17:00:29 -07:00

vla.cl

…