Alex Voicu e13cbaca69 [clang][CodeGen][SPIR-V] Fix incorrect SYCL usage, implement missing interface (#109415) ...

This is primarily meant to address the issue identified in #109182,
around incorrect usage of `-fsycl-is-device`; we now have AMDGCN
flavoured SPIR-V which retains the desired behaviour around the default
AS and does not depend on the SYCL language being enabled to do so.
Overall, there are three changes:

1. We unconditionally use the `SPIRDefIsGen` AS map for AMDGCNSPIRV
target, as there is no case where the hack of setting default to private
would be desirable, and it can be used for languages other than OCL/HIP;
2. We implement `SPIRVTargetCodeGenInfo::getGlobalVarAddressSpace` for
SPIR-V in general, because otherwise using it from languages other than
HIP or OpenCL would yield 0, incorrectly;
3. We remove the incorrect usage of `-fsycl-is-device`.

2024-09-26 14:06:14 +01:00

..

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

…

addr-space-struct-arg.cl

[IRBuilder] Generate nuw GEPs for struct member accesses (#99538)

2024-08-09 13:25:04 +01:00

address-space-constant-initializers.cl

…

address-spaces-conversions.cl

…

address-spaces-mangling.cl

…

address-spaces.cl

…

amdgcn-automatic-variable.cl

…

amdgcn-buffer-rsrc-type.cl

[IRBuilder] Generate nuw GEPs for struct member accesses (#99538)

2024-08-09 13:25:04 +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

[IRBuilder] Generate nuw GEPs for struct member accesses (#99538)

2024-08-09 13:25:04 +01:00

amdgpu-abi-struct-coerce.cl

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

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

amdgpu-alignment.cl

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

2024-05-04 17:15:51 -07: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

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

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

[IRBuilder] Generate nuw GEPs for struct member accesses (#99538)

2024-08-09 13:25:04 +01: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-features.cl

[AMDGPU] Add a new target gfx1152 (#94534)

2024-06-06 12:16:11 -04:00

amdgpu-ieee.cl

…

amdgpu-nullptr.cl

…

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

[FuncAttrs] Deduce noundef attributes for return values (#76553)

2023-12-31 20:44:48 +08:00

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

[SPIRV][RFC] Rework / extend support for memory scopes (#106429)

2024-09-25 00:44:57 +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

…

atomics-cas-remarks-gfx90a.cl

clang/AMDGPU: Update test message order

2024-09-06 21:18:41 +04:00

atomics-unsafe-hw-remarks-gfx90a.cl

clang/AMDGPU: Update test message order

2024-09-06 21:18:41 +04:00

backend-unsupported-warning.ll

…

blocks.cl

[IRBuilder] Generate nuw GEPs for struct member accesses (#99538)

2024-08-09 13:25:04 +01:00

bool_cast.cl

…

builtins-alloca.cl

[Clang] Use private address space for builtin_alloca return type for OpenCL (#95750)

2024-07-26 15:24:06 +05:30

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

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

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

builtins-amdgcn-dl-insts-gfx11.cl

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

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

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

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

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

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

builtins-amdgcn-gfx10.cl

[AMDGPU] Extend permlane16, permlanex16 and permlane64 intrinsic lowering for generic types (#92725)

2024-06-26 09:24:09 +05:30

builtins-amdgcn-gfx11-err.cl

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

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

[clang][CodeGen][SPIR-V] Fix incorrect SYCL usage, implement missing interface (#109415)

2024-09-26 14:06:14 +01: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] Add target intrinsic for s_buffer_prefetch_data (#107293)

2024-09-06 11:41:21 -07: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] Add target intrinsic for s_buffer_prefetch_data (#107293)

2024-09-06 11:41:21 -07:00

builtins-amdgcn-gfx940.cl

[AMDGPU][Clang] Builtin for GLOBAL_LOAD_LDS on GFX940 (#92962)

2024-05-22 00:03:59 -04: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

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

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

[Clang][AMDGPU] Add a builtin for llvm.amdgcn.make.buffer.rsrc intrinsic (#95276)

2024-06-20 11:01:54 -04:00

builtins-amdgcn-mfma.cl

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

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

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

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

builtins-amdgcn-swmmac-w64.cl

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

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

builtins-amdgcn-vi.cl

clang/AMDGPU: Use atomicrmw for ds fmin/fmax builtins (#96738)

2024-06-27 15:32:08 +02: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

[clang][CodeGen][SPIR-V] Fix incorrect SYCL usage, implement missing interface (#109415)

2024-09-26 14:06:14 +01:00

builtins-f16.cl

[clang] Emit constraint intrinsics for arc and hyperbolic trig clang builtins (#98949)

2024-07-19 10:19:41 -04:00

builtins-fp-atomics-gfx8.cl

clang/AMDGPU: Emit atomicrmw for __builtin_amdgcn_global_atomic_fadd_{f32|f64} (#96872)

2024-08-15 22:59:24 +04:00

builtins-fp-atomics-gfx12.cl

clang/AMDGPU: Emit atomicrmw for global/flat fadd v2bf16 builtins (#96875)

2024-08-20 23:20:03 +04:00

builtins-fp-atomics-gfx90a.cl

clang/AMDGPU: Emit atomicrmw for flat/global atomic min/max f64 builtins (#96876)

2024-08-20 23:24:15 +04:00

builtins-fp-atomics-gfx940.cl

clang/AMDGPU: Emit atomicrmw for global/flat fadd v2bf16 builtins (#96875)

2024-08-20 23:20:03 +04: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

…

byval.cl

…

cast_image.cl

…

cl20-device-side-enqueue-attributes.cl

[IRBuilder] Generate nuw GEPs for struct member accesses (#99538)

2024-08-09 13:25:04 +01:00

cl20-device-side-enqueue.cl

…

cl-strict-aliasing.cl

…

cl-uniform-wg-size.cl

…

const-str-array-decay.cl

…

constant-addr-space-globals.cl

…

convergent.cl

…

debug-info-programming-language.cl

…

enqueue-kernel-non-entry-block.cl

…

event_t.cl

…

ext-int-shift.cl

…

ext-vector-shuffle.cl

…

fdeclare-opencl-builtins.cl

…

fpmath.cl

…

func-call-dbg-loc.cl

…

half.cl

…

images.cl

…

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

…

kernel-attributes.cl

…

kernel-metadata.cl

…

kernel-param-alignment.cl

…

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

…

lifetime.cl

…

local-initializer-undef.cl

…

local.cl

…

logical-ops.cl

…

memcpy.cl

…

no-half.cl

…

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

…

overload.cl

…

partial_initializer.cl

[Clang] Use poison as base for vector literals

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

pipe_builtin.cl

…

pipe_types_mangling.cl

…

pipe_types.cl

…

preserve_vec3.cl

…

printf.cl

…

private-array-initialization.cl

…

ptx-calls.cl

…

ptx-kernels.cl

…

reflect.cl

[NVPTX] Add clang builtin for __nvvm_reflect intrinsic (#81277)

2024-02-09 14:11:01 -06:00

relaxed-fpmath.cl

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

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

sampler.cl

…

shifts.cl

…

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

…

unroll-hint.cl

…

vector_literals.cl

[Clang] Use poison as base for vector literals

2023-12-19 11:53:18 +01: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

…