[OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC)This adds -no-opaque-pointers to clang tests whose output willchange when opaque pointers are enabled by default. This isintended to be p
[OpaquePtrs][Clang] Add -no-opaque-pointers to tests (NFC)This adds -no-opaque-pointers to clang tests whose output willchange when opaque pointers are enabled by default. This isintended to be part of the migration approach described inhttps://discourse.llvm.org/t/enabling-opaque-pointers-by-default/61322/9.The patch has been produced by replacing %clang_cc1 with%clang_cc1 -no-opaque-pointers for tests that fail with opaquepointers enabled. Worth noting that this doesn't cover all tests,there's a remaining ~40 tests not using %clang_cc1 that will needa followup change.Differential Revision: https://reviews.llvm.org/D123115
show more ...
[HIPSPV] Fix literals are mapped to Generic address spaceThis issue is an oversight in D108621.Literals in HIP are emitted as global constant variables with defaultaddress space which maps to Ge
[HIPSPV] Fix literals are mapped to Generic address spaceThis issue is an oversight in D108621.Literals in HIP are emitted as global constant variables with defaultaddress space which maps to Generic address space for HIPSPV. InSPIR-V such variables translate to OpVariable instructions withGeneric storage class which are not legal. Fix by mapping literalsto CrossWorkGroup address space.The literals are not mapped to UniformConstant because the “flat”pointers in HIP may reference them and “flat” pointers are modeledas Generic pointers in SPIR-V. In SPIR-V/OpenCL UniformConstantpointers may not be casted to Generic.Patch by: Henry LinjamäkiReviewed by: Yaxun LiuDifferential Revision: https://reviews.llvm.org/D118876
[Clang/Test]: Rename enable_noundef_analysis to disable-noundef-analysis and turn it off by defaultTurning on `enable_noundef_analysis` flag allows better codegen by removing freeze instructions.I
[Clang/Test]: Rename enable_noundef_analysis to disable-noundef-analysis and turn it off by defaultTurning on `enable_noundef_analysis` flag allows better codegen by removing freeze instructions.I modified clang by renaming `enable_noundef_analysis` flag to `disable-noundef-analysis` and turning it off by default.Test updates are made as a separate patch: D108453Reviewed By: eugenisDifferential Revision: https://reviews.llvm.org/D105169
[HIPSPV] Convert HIP kernels to SPIR-V kernelsThis patch translates HIP kernels to SPIR-V kernels when the HIPcompilation mode is targeting SPIR-S. This involves:* Setting Cuda calling conventio
[HIPSPV] Convert HIP kernels to SPIR-V kernelsThis patch translates HIP kernels to SPIR-V kernels when the HIPcompilation mode is targeting SPIR-S. This involves:* Setting Cuda calling convention to CC_OpenCLKernel (which maps to SPIR_KERNEL in LLVM IR later on).* Coercing pointer arguments with default address space (AS) qualifier to CrossWorkGroup AS (__global in OpenCL). HIPSPV's device code is ultimately SPIR-V for OpenCL execution environment (as starter/default) where Generic or Function (OpenCL's private) is not supported as storage class for kernel pointer types. This leaves the CrossWorkGroup to be the only reasonable choice for HIP buffers.Reviewed By: yaxunlDifferential Revision: https://reviews.llvm.org/D109818
[HIPSPV] Add CUDA->SPIR-V address space mappingAdd mapping for CUDA address spaces for HIP to SPIR-Vtranslation. This change allows HIP device code to beemitted as valid SPIR-V by mapping unquali
[HIPSPV] Add CUDA->SPIR-V address space mappingAdd mapping for CUDA address spaces for HIP to SPIR-Vtranslation. This change allows HIP device code to beemitted as valid SPIR-V by mapping unqualified pointersto generic address space and by mapping __device__ and__shared__ AS to their equivalent AS in SPIR-V(CrossWorkgroup and Workgroup, respectively).Cuda's __constant__ AS is handled specially. In HIPunqualified pointers (aka "flat" pointers) can point to__constant__ objects. Mapping this AS to ConstantMemorywould produce to illegal address space casts togeneric AS. Therefore, __constant__ AS is mapped toCrossWorkgroup.Patch by linjamaki (Henry Linjamäki)!Differential Revision: https://reviews.llvm.org/D108621
[NFC] Fixed ignored .hip test.Reviewers: hliaoReviewed By: hliaoSubscribers: yaxunl, cfe-commits, llvm-commitsTags: #clangDifferential Revision: https://reviews.llvm.org/D82764
[clang][amdgpu] Prefer not using `fp16` conversion intrinsics.Reviewers: yaxunl, arsenmSubscribers: kzhuravl, jvesely, wdng, nhaehnle, dstuttard, tpr, t-tye, kerbowa, cfe-commitsTags: #clangD
[clang][amdgpu] Prefer not using `fp16` conversion intrinsics.Reviewers: yaxunl, arsenmSubscribers: kzhuravl, jvesely, wdng, nhaehnle, dstuttard, tpr, t-tye, kerbowa, cfe-commitsTags: #clangDifferential Revision: https://reviews.llvm.org/D81849
Don't dump IR output from this test to stdout.
[HIP][AMDGPU] expand printf when compiling HIP to AMDGPUSummary:This change implements the expansion in two parts:- Add a utility function emitAMDGPUPrintfCall() in LLVM.- Invoke the above funct
[HIP][AMDGPU] expand printf when compiling HIP to AMDGPUSummary:This change implements the expansion in two parts:- Add a utility function emitAMDGPUPrintfCall() in LLVM.- Invoke the above function from Clang CodeGen, when processing a HIP program for the AMDGPU target.The printf expansion has undefined behaviour if the format string isnot a compile-time constant. As a sufficient condition, the HIPToolChain now emits -Werror=format-nonliteral.Reviewed By: arsenmDifferential Revision: https://reviews.llvm.org/D71365