15 Commits

Author SHA1 Message Date
Nick Sarnie
52a96491e1
[clang][SPIR-V] Addrspace of opencl_global should always be 1 (#136753)
This fixes a CUDA SPIR-V regression introduced in
https://github.com/llvm/llvm-project/pull/134399.

---------

Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>
2025-04-24 14:20:13 +00:00
Alexander Shaposhnikov
297f0b3f4c
[CudaSPIRV] Allow using integral non-type template parameters as attribute args (#131546)
Allow using integral non-type template parameters as attribute arguments
of
reqd_work_group_size and work_group_size_hint.

Test plan:
ninja check-all
2025-03-19 10:11:18 -07:00
Pedro Lobo
ccf2109471
[Metadata] Change placeholder from undef to poison (#131469)
Replace `undef` constant metadata uses with `poison`.
2025-03-17 22:16:18 +00:00
Alexander Shaposhnikov
2dc123b33d
[clang][opencl] Allow passing all zeros to reqd_work_group_size (#131543)
Allow passing all zeros to reqd_work_group_size.

Test plan: ninja check-all
2025-03-16 16:21:46 -07:00
Alexander Shaposhnikov
df13acf344
[CudaSPIRV] Add support for optional spir-v attributes (#116589)
Add support for optional spir-v attributes.

Test plan:
ninja check-all
2024-11-19 13:14:45 -08:00
ShangwuYao
8bd13ad6c5 [CUDA][SPIRV] Match builtin types and __GCC_ATOMIC_XXX_LOCK_FREE macros on host/device
This change matches the CUDA/SPIRV behavior with CUDA/NVPTX, and makes some builtin types
and __GCC_ATOMIC_XXX_LOCK_FREE macros the same between the host and device. This is only
done when host triple is provided and known, otherwise the behavior is unchanged.

Reviewed By: yaxunl

Differential Revision: https://reviews.llvm.org/D144047
2023-02-22 15:44:46 +00:00
Nikita Popov
d8f91f2c1d [CodeGenCUDASPIRV] Convert tests to opaque pointers (NFC) 2022-12-14 13:36:07 +01:00
Shangwu Yao
31d8dbd1e5 [CUDA/SPIR-V] Force passing aggregate type byval
This patch forces copying aggregate type in kernel arguments by value when
compiling CUDA targeting SPIR-V. The original behavior is not passing by value
when there is any of destructor, copy constructor and move constructor defined
by user. This patch makes the behavior of SPIR-V generated from CUDA follow
the CUDA spec
(https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#global-function-argument-processing),
and matches the NVPTX
implementation (
41958f76d8/clang/lib/CodeGen/TargetInfo.cpp (L7241)).

Differential Revision: https://reviews.llvm.org/D130387
2022-07-22 20:30:15 +00:00
Fangrui Song
63fbc77121 [Driver][test] Remove unused/obsoleted REQUIRES: clang-driver
It (introduced by 556d713c70bfaf58ac18d089883f9c34c581633a) appears to be
related to the removed dragonegg project. In addition, the feature was a bit
misnamed and may lur users to unnecessarily use it.
2022-04-12 13:29:46 -07:00
Nikita Popov
b16a3b4f3b [Clang] Add -no-opaque-pointers to more tests (NFC)
This adds the flag to more tests that were not caught by the
mass-migration in 532dc62b907554b3f07f17205674aa71e76fc863.
2022-04-07 12:53:29 +02:00
Shangwu Yao
15a1769631 Emit OpenCL metadata when targeting SPIR-V
This is required for converting function calls such as get_global_id()
into SPIR-V builtins.

Differential Revision: https://reviews.llvm.org/D123049
2022-04-05 20:58:32 +00:00
Shangwu Yao
c2f501f395
[CUDA][SPIRV] Assign global address space to CUDA kernel arguments
(resubmit https://reviews.llvm.org/D119207 after fixing the test for
some build settings)

This patch converts CUDA pointer kernel arguments with default address
space to CrossWorkGroup address space (__global in OpenCL). This is
because Generic or Function (OpenCL's private) is not supported as
storage class for kernel pointer types.

Differential revision: https://reviews.llvm.org/D120366
2022-02-24 20:51:43 -08:00
Matthew Voss
9ce09099bb Revert "[CUDA][SPIRV] Assign global address space to CUDA kernel arguments"
This reverts commit 9de4fc0f2d3b60542956f7e5254951d049edeb1f.

Reverting due to test failure: https://lab.llvm.org/buildbot/#/builders/139/builds/17199
2022-02-17 14:32:10 -08:00
Shangwu Yao
9de4fc0f2d
[CUDA][SPIRV] Assign global address space to CUDA kernel arguments
This patch converts CUDA pointer kernel arguments with default address space to
CrossWorkGroup address space (__global in OpenCL). This is because Generic or
Function (OpenCL's private) is not supported as storage class for kernel pointer types.

Differential Revision: https://reviews.llvm.org/D119207
2022-02-17 09:38:06 -08:00
Daniele Castagna
abbdc13e68 [CUDA][SPIRV] Use OpenCLKernel CC for CUDA -> SPIRV
Select the OpenCLKernel calling convention for kernels when compiling
CUDA targeting SPIR-V.

In this way the generated LLVM IR will have a spir_kernel calling
convention that will be translated to an OpEntryPoint when converting
to SPIRV.

Reviewed By: jlebar, tra

Differential Revision: https://reviews.llvm.org/D114407
2021-12-06 15:06:57 -08:00