In amdgcn target, null pointers in global, constant, and generic address space take value 0 but null pointers in private and local address space take value -1. Currently LLVM assumes all null pointers take value 0, which results in incorrectly translated IR. To workaround this issue, instead of emit null pointers in local and private address space, a null pointer in generic address space is emitted and casted to local and private address space.
Tentative definition of global variables with non-zero initializer will have weak linkage instead of common linkage since common linkage requires zero initializer and does not have explicit section to hold the non-zero value.
Virtual member functions getNullPointer and performAddrSpaceCast are added to TargetCodeGenInfo which by default returns ConstantPointerNull and emitting addrspacecast instruction. A virtual member function getNullPointerValue is added to TargetInfo which by default returns 0. Each target can override these virtual functions to get target specific null pointer and the null pointer value for specific address space, and perform specific translations for addrspacecast.
Wrapper functions getNullPointer is added to CodegenModule and getTargetNullPointerValue is added to ASTContext to facilitate getting the target specific null pointers and their values.
This change has no effect on other targets except amdgcn target. Other targets can provide support of non-zero null pointer in a similar way.
This change only provides support for non-zero null pointer for C and OpenCL. Supporting for other languages will be added later incrementally.
Differential Revision: https://reviews.llvm.org/D26196
llvm-svn: 289252
Avoid using shortcut for const qualified non-constant address space
aggregate variables while generating them on the stack such that
the alloca object is used instead of a global variable containing
initializer.
Review: https://reviews.llvm.org/D27109
llvm-svn: 288163
The wave barrier represents the discardable barrier. Its main purpose is to
carry convergent attribute, thus preventing illegal CFG optimizations. All lanes
in a wave come to convergence point simultaneously with SIMT, thus no special
instruction is needed in the ISA. The barrier is discarded during code generation.
Differential Revision: https://reviews.llvm.org/D26584
llvm-svn: 287006
Make handling integer parameters more flexible:
- For the number of events argument allow to pass larger
integers than 32 bits as soon as compiler can prove that
the range fits in 32 bits. If not, the diagnostic will be given.
- Change type of the arguments specifying the sizes of
the corresponding block arguments to be size_t.
Review: https://reviews.llvm.org/D26509
llvm-svn: 286849
- Accept NULL pointer as a valid parameter value for clk_event.
- Generate clk_event_t arguments of internal
__enqueue_kernel_XXX function as pointers in generic address space.
Review: https://reviews.llvm.org/D26507
llvm-svn: 286836
It doesn't make sense to use the target's address space ids in this context as
this is metadata that should be referring to the "logical" OpenCL address spaces.
For flat AS machines like all "CPUs" in general, the logical AS info gets lost as
there's only one address space (0).
This commit changes the logic such that we always use the SPIR address space
ids for the argument metadata. It thus allows implementing the clGetKernelArgInfo()
and the other detection needs.
https://reviews.llvm.org/D26157
llvm-svn: 286819
This change makes sure single-precision floating point types are used if the
cl_fp64 extension is not supported by the target.
Also removed the check to see whether the OpenCL version is >= 1.2, as this has
been incorporated into the extension setting code.
Differential Revision: https://reviews.llvm.org/D24235
llvm-svn: 286815
Certain OpenCL builtin functions are supposed to be executed by all threads in a work group or sub group. Such functions should not be made divergent during transformation. It makes sense to mark them with convergent attribute.
The adding of convergent attribute is based on Ettore Speziale's work and the original proposal and patch can be found at https://www.mail-archive.com/cfe-commits@lists.llvm.org/msg22271.html.
Differential Revision: https://reviews.llvm.org/D25343
llvm-svn: 285725
Summary: Setting constant address space for global constants used for memcpy-initialization of arrays.
Patch by Alexey Sotkin.
Reviewers: bader, yaxunl, Anastasia
Subscribers: cfe-commits, AlexeySotkin
Differential Revision: https://reviews.llvm.org/D25305
llvm-svn: 285557
Currently Clang allows partial initializer for C99 but not for OpenCL, e.g.
float a[16][16] = {1.0f, 2.0f};
is allowed in C99 but not allowed in OpenCL.
This patch fixes that.
Differential Revision: https://reviews.llvm.org/D25335
llvm-svn: 283891
__builtin_astype is used to cast OpenCL opaque types to other types, as such, it needs to be able to handle casting from and to pointer types correctly.
Current it cannot handle 1) casting between pointers of different addr spaces 2) casting between pointer type and non-pointer types.
This patch fixes that.
Differential Revision: https://reviews.llvm.org/D25123
llvm-svn: 283114
__attribute__((amdgpu_flat_work_group_size(<min>, <max>))) - request minimum and maximum flat work group size
__attribute__((amdgpu_waves_per_eu(<min>[, <max>]))) - request minimum and/or maximum waves per execution unit
Differential Revision: https://reviews.llvm.org/D24513
llvm-svn: 282371
Fix target options for fp32/64-denormals so that
+fp64-denormals is set if fp64 is supported
-fp32-denormals if fp32 denormals is not supported, or -cl-denorms-are-zero is set
+fp32-denormals if fp32 denormals is supported and -cl-denorms-are-zero is not set
If target feature fp32/64-denormals is explicitly set, they will override default options and options deduced from -cl-denorms-are-zero.
Differential Revision: https://reviews.llvm.org/D24512
llvm-svn: 281357
Summary:
Remove access qualifiers on images in arg info metadata:
* kernel_arg_type
* kernel_arg_base_type
Image access qualifiers are inseparable from type in clang implementation,
but OpenCL spec provides a special query to get access qualifier
via clGetKernelArgInfo with CL_KERNEL_ARG_ACCESS_QUALIFIER.
Besides that OpenCL conformance test_api get_kernel_arg_info expects
image types without access qualifier.
Patch by Evgeniy Tyurin.
Reviewers: bader, yaxunl, Anastasia
Subscribers: cfe-commits
Differential Revision: https://reviews.llvm.org/D23915
llvm-svn: 280699
Structs are currently handled as pointer + byval, which makes AMDGPU
LLVM backend generate incorrect code when structs are used. This patch
changes struct argument to be handled directly and without flattening,
which Clover (Mesa 3D Gallium OpenCL state tracker) will be able to
handle. Flattening would expand the struct to individual elements and
pass each as a separate argument, which Clover can not
handle. Furthermore, such expansion does not fit the OpenCL
programming model which requires to explicitely specify each argument
index, size and memory location.
Patch by Vedran Miletić
llvm-svn: 279463
Summary:
int __builtin_amdgcn_ds_swizzle (int a, int imm);
while imm is a constant.
Differential Revision:
http://reviews.llvm.org/D23682
llvm-svn: 279165
Pointers of certain GPUs in AMDGCN target in private address space is 32 bit but pointers in other address spaces are 64 bit. size_t type should be defined as 64 bit for these GPUs so that it could hold pointers in all address spaces. Also fixed issues in pointer arithmetic codegen by using pointer specific intptr type.
Differential Revision: https://reviews.llvm.org/D23361
llvm-svn: 279121
Let the driver pass the option to frontend. Do not set precision metadata for division instructions when this option is set. Set function attribute "correctly-rounded-divide-sqrt-fp-math" based on this option.
Differential Revision: https://reviews.llvm.org/D22940
llvm-svn: 278155
Adjust target features for amdgcn target when -cl-denorms-are-zero is set.
Denormal support is controlled by feature strings fp32-denormals fp64-denormals in amdgcn target. If -cl-denorms-are-zero is not set and the command line does not set fp32/64-denormals feature string, +fp32-denormals +fp64-denormals will be on for GPU's supporting them.
A new virtual function virtual void TargetInfo::adjustTargetOptions(const CodeGenOptions &CGOpts, TargetOptions &TargetOpts) const is introduced to allow adjusting target option by codegen option.
Differential Revision: https://reviews.llvm.org/D22815
llvm-svn: 278151
Summary:
In order to re-define OpenCL built-in functions
'to_{private,local,global}' in OpenCL run-time library LLVM names must
be different from the clang built-in function names.
Reviewers: yaxunl, Anastasia
Subscribers: cfe-commits
Differential Revision: https://reviews.llvm.org/D23120
llvm-svn: 277743
Currently Clang use int32 to represent sampler_t, which have been a source of issue for some backends, because in some backends sampler_t cannot be represented by int32. They have to depend on kernel argument metadata and use IPA to find the sampler arguments and global variables and transform them to target specific sampler type.
This patch uses opaque pointer type opencl.sampler_t* for sampler_t. For each use of file-scope sampler variable, it generates a function call of __translate_sampler_initializer. For each initialization of function-scope sampler variable, it generates a function call of __translate_sampler_initializer.
Each builtin library can implement its own __translate_sampler_initializer(). Since the real sampler type tends to be architecture dependent, allowing it to be initialized by a library function simplifies backend design. A typical implementation of __translate_sampler_initializer could be a table lookup of real sampler literal values. Since its argument is always a literal, the returned pointer is known at compile time and easily optimized to finally become some literal values directly put into image read instructions.
This patch is partially based on Alexey Sotkin's work in Khronos Clang (3d4eec6162).
Differential Revision: https://reviews.llvm.org/D21567
llvm-svn: 277024
Allows AMDGCN target to generate images (such as %opencl.image2d_t) in constant address space.
Images will still be generated in global address space by default.
Added tests to existing opencl-types.cl in test\CodeGenOpenCL.
Patch by Aaron En Ye Shi.
Differential Revision: https://reviews.llvm.org/D22523
llvm-svn: 276161
Added the opencl.ocl.version metadata to be emitted with amdgcn. Created a static function emitOCLVerMD which is shared between triple spir and target amdgcn.
Also added new testcases to existing test file, spir_version.cl inside test/CodeGenOpenCL.
Patch by Aaron En Ye Shi.
Differential Revision: https://reviews.llvm.org/D22424
llvm-svn: 276010
This reverts also r275029, "Update Clang tests after adding inference for the returned argument attribute"
It broke LTO build. Seems miscompilation.
llvm-svn: 275756
Improved test with user define structure pipe type case.
Reviewers: Anastasia, pxli168
Subscribers: yaxunl, cfe-commits
Differential revision: http://reviews.llvm.org/D21744
llvm-svn: 275259
Yaxun Liu