[CUDA/SPIR-V] Force passing aggregate type byvalThis patch forces copying aggregate type in kernel arguments by value whencompiling CUDA targeting SPIR-V. The original behavior is not passing by v
[CUDA/SPIR-V] Force passing aggregate type byvalThis patch forces copying aggregate type in kernel arguments by value whencompiling CUDA targeting SPIR-V. The original behavior is not passing by valuewhen there is any of destructor, copy constructor and move constructor definedby user. This patch makes the behavior of SPIR-V generated from CUDA followthe CUDA spec(https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#global-function-argument-processing),and matches the NVPTXimplementation (https://github.com/llvm/llvm-project/blob/41958f76d8a2c47484fa176cba1de565cfe84de7/clang/lib/CodeGen/TargetInfo.cpp#L7241).Differential Revision: https://reviews.llvm.org/D130387
show more ...
[Driver][test] Remove unused/obsoleted REQUIRES: clang-driverIt (introduced by 556d713c70bfaf58ac18d089883f9c34c581633a) appears to berelated to the removed dragonegg project. In addition, the fea
[Driver][test] Remove unused/obsoleted REQUIRES: clang-driverIt (introduced by 556d713c70bfaf58ac18d089883f9c34c581633a) appears to berelated to the removed dragonegg project. In addition, the feature was a bitmisnamed and may lur users to unnecessarily use it.
[Clang] Add -no-opaque-pointers to more tests (NFC)This adds the flag to more tests that were not caught by themass-migration in 532dc62b907554b3f07f17205674aa71e76fc863.
Emit OpenCL metadata when targeting SPIR-VThis is required for converting function calls such as get_global_id()into SPIR-V builtins.Differential Revision: https://reviews.llvm.org/D123049
[CUDA][SPIRV] Assign global address space to CUDA kernel arguments(resubmit https://reviews.llvm.org/D119207 after fixing the test forsome build settings)This patch converts CUDA pointer kernel
[CUDA][SPIRV] Assign global address space to CUDA kernel arguments(resubmit https://reviews.llvm.org/D119207 after fixing the test forsome build settings)This patch converts CUDA pointer kernel arguments with default addressspace to CrossWorkGroup address space (__global in OpenCL). This isbecause Generic or Function (OpenCL's private) is not supported asstorage class for kernel pointer types.Differential revision: https://reviews.llvm.org/D120366
Revert "[CUDA][SPIRV] Assign global address space to CUDA kernel arguments"This reverts commit 9de4fc0f2d3b60542956f7e5254951d049edeb1f.Reverting due to test failure: https://lab.llvm.org/buildbo
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
[CUDA][SPIRV] Assign global address space to CUDA kernel argumentsThis patch converts CUDA pointer kernel arguments with default address space toCrossWorkGroup address space (__global in OpenCL).
[CUDA][SPIRV] Assign global address space to CUDA kernel argumentsThis patch converts CUDA pointer kernel arguments with default address space toCrossWorkGroup address space (__global in OpenCL). This is because Generic orFunction (OpenCL's private) is not supported as storage class for kernel pointer types.Differential Revision: https://reviews.llvm.org/D119207
[CUDA][SPIRV] Use OpenCLKernel CC for CUDA -> SPIRVSelect the OpenCLKernel calling convention for kernels when compilingCUDA targeting SPIR-V.In this way the generated LLVM IR will have a spir_k
[CUDA][SPIRV] Use OpenCLKernel CC for CUDA -> SPIRVSelect the OpenCLKernel calling convention for kernels when compilingCUDA targeting SPIR-V.In this way the generated LLVM IR will have a spir_kernel callingconvention that will be translated to an OpEntryPoint when convertingto SPIRV.Reviewed By: jlebar, traDifferential Revision: https://reviews.llvm.org/D114407