Page MenuHomePhabricator

Anastasia (Anastasia Stulova)
Compiler Engineer at ARM / Code Owner of OpenCL in Clang

Projects

User does not belong to any projects.

User Details

User Since
Mar 5 2015, 8:05 AM (219 w, 3 d)

Recent Activity

Fri, May 10

Anastasia added a comment to D60763: Prototype OpenCL BIFs using Tablegen.

The overall structure looks very clear now! I just have a couple of nitpicks! Thank you!

Fri, May 10, 4:23 AM · Restricted Project

Thu, May 9

Anastasia committed rGe6cf6c78f88f: [OpenCL] Make global ctor init function a kernel (authored by Anastasia).
[OpenCL] Make global ctor init function a kernel
Thu, May 9, 6:54 AM
Anastasia committed rG933e305ed909: [OpenCL] Switched CXX mode to be derived from C++17 (authored by Anastasia).
[OpenCL] Switched CXX mode to be derived from C++17
Thu, May 9, 4:55 AM
Anastasia added a comment to D61488: [OpenCL] Make global ctor init function a kernel.

The comments could use a bit more polishing but nothing that justifies making another iteration IMHO. LGTM!

Thu, May 9, 3:58 AM · Restricted Project
Anastasia committed rGeba9a6e08fcb: [SPIR] Simplified target checking. (authored by Anastasia).
[SPIR] Simplified target checking.
Thu, May 9, 3:24 AM
Anastasia committed rG3cdf8981054b: [SPIR] Simplified target checking. (authored by Anastasia).
[SPIR] Simplified target checking.
Thu, May 9, 3:16 AM

Wed, May 8

Anastasia updated the diff for D61488: [OpenCL] Make global ctor init function a kernel.
  • Improved comments
  • Switched to SPIR kernel
Wed, May 8, 8:56 AM · Restricted Project
Anastasia committed rG5b6dda33d122: [Sema][OpenCL] Make address space conversions a bit stricter. (authored by Anastasia).
[Sema][OpenCL] Make address space conversions a bit stricter.
Wed, May 8, 7:24 AM
Anastasia updated the diff for D61506: [OpenCL] Switch to C++17.

Added HexFloat!

Wed, May 8, 5:00 AM · Restricted Project
Anastasia added a comment to D58236: Make address space conversions a bit stricter..

This was accepted a while ago, but never landed. I don't have commit access; could someone commit it?

Wed, May 8, 4:08 AM · Restricted Project, Restricted Project
Anastasia added a comment to D61506: [OpenCL] Switch to C++17.

Per the OpenCL C++ 1.0 specification, section 2:

The OpenCL C++ programming language is based on the ISO/IEC JTC1 SC22 WG21 N 3690 language specification (a.k.a. C++14 specification).

I think it would be reasonable to permit changing the base C++ standard in OpenCL C++ mode

Indeed! There should be an option to pick the version of C++ the user wants for OpenCL.

Wed, May 8, 2:59 AM · Restricted Project

Tue, May 7

Anastasia added inline comments to D59988: [PR41276] Generate address space cast of 'this' for objects attributed by an address space in C++.
Tue, May 7, 7:44 AM · Restricted Project
Anastasia accepted D61639: Add Triple::isSPIR() to simplify code.

LGTM! Thanks!

Tue, May 7, 7:26 AM · Restricted Project, Restricted Project
Anastasia committed rGd6865b7d71bc: [OpenCL] Prevent mangling kernel functions. (authored by Anastasia).
[OpenCL] Prevent mangling kernel functions.
Tue, May 7, 7:21 AM
Anastasia added a comment to D61506: [OpenCL] Switch to C++17.

Per the OpenCL C++ 1.0 specification, section 2:

The OpenCL C++ programming language is based on the ISO/IEC JTC1 SC22 WG21 N 3690 language specification (a.k.a. C++14 specification).

I think it would be reasonable to permit changing the base C++ standard in OpenCL C++ mode, but we'd need a good reason to deviate from the behavior specified in the OpenCL C++ specification by default.

Tue, May 7, 2:57 AM · Restricted Project

Fri, May 3

Anastasia added a comment to D60379: Make precompiled headers reproducible.

Can this be closed now due to commit r358674?

Fri, May 3, 7:58 AM · Restricted Project
Anastasia added inline comments to D59988: [PR41276] Generate address space cast of 'this' for objects attributed by an address space in C++.
Fri, May 3, 7:46 AM · Restricted Project
Anastasia created D61506: [OpenCL] Switch to C++17.
Fri, May 3, 7:05 AM · Restricted Project
Anastasia added inline comments to D61488: [OpenCL] Make global ctor init function a kernel.
Fri, May 3, 5:02 AM · Restricted Project
Anastasia created D61488: [OpenCL] Make global ctor init function a kernel.
Fri, May 3, 5:01 AM · Restricted Project
Anastasia updated the diff for D61318: [Sema] Prevent binding references with mismatching address spaces to temporaries.

Updated comments!

Fri, May 3, 3:02 AM
Anastasia added a comment to D57464: Generalize method overloading on addr spaces to C++.

I think I would lean towards the latter since it means less fudging around with a whole bunch of unrelated methods. Do @rjmccall or @rsmith have any further opinions on this?

Ok, I can change the patch to prototype this approach. I might need some example test cases though.

Alright!

Just to make sure of something here - are you waiting for me to provide some example cases? There hasn't been activity here in a while and I was just wondering if it was because you were waiting for this.

Fri, May 3, 2:47 AM
Anastasia added a comment to D60454: [OpenCL] Prevent mangling kernel functions.

I think it would be more reasonable to just change getDeclLanguageLinkage to check for a kernel function.

Fri, May 3, 2:39 AM · Restricted Project
Anastasia updated the diff for D60454: [OpenCL] Prevent mangling kernel functions.

Removed redundant C linkage checks!

Fri, May 3, 2:31 AM · Restricted Project

Thu, May 2

Anastasia committed rG26e095e84f0d: [OpenCL] Fix initialisation of this via pointer. (authored by Anastasia).
[OpenCL] Fix initialisation of this via pointer.
Thu, May 2, 9:11 AM
Anastasia accepted D61319: [PR41674] [OpenCL] Fix initialisation of this via pointer.

LGTM! Thanks!

Thu, May 2, 8:18 AM · Restricted Project
Anastasia updated the diff for D60454: [OpenCL] Prevent mangling kernel functions.

Extended FunctionDecl exten C logic to take kernel function into account.

Thu, May 2, 8:15 AM · Restricted Project
Anastasia committed rG44bb0aa99414: [OpenCL] Deduce static data members to __global addr space. (authored by Anastasia).
[OpenCL] Deduce static data members to __global addr space.
Thu, May 2, 7:39 AM
Anastasia updated the diff for D61318: [Sema] Prevent binding references with mismatching address spaces to temporaries.

Added extra overload for isAddressSpaceSupersetOf.

Thu, May 2, 6:42 AM
Anastasia added inline comments to D61274: [Sema][AST] Explicit visibility for OpenCL/CUDA kernels/variables.
Thu, May 2, 6:11 AM · Restricted Project
Anastasia added inline comments to D61319: [PR41674] [OpenCL] Fix initialisation of this via pointer.
Thu, May 2, 3:27 AM · Restricted Project

Wed, May 1

Anastasia added a comment to D61304: [OpenCL][PR41609] Deduce static data members to __global addr space.

Presumably a similar rule would apply to thread-locals if you supported them.

Wed, May 1, 2:08 AM · Restricted Project
Anastasia added inline comments to D61318: [Sema] Prevent binding references with mismatching address spaces to temporaries.
Wed, May 1, 2:03 AM
Anastasia added a comment to D61319: [PR41674] [OpenCL] Fix initialisation of this via pointer.

I would suggest an IR test, especially since what you want to achieve is addrspacecast in IR output. We use AST dump only if there is no other good way to test something. In fact we already have a test with similar logic that tests various OpenCL address spaces: test/CodeGenOpenCLCXX/addrspace-of-this.cl. I think it should be reasonable enough to extend it.

Wed, May 1, 1:32 AM · Restricted Project

Tue, Apr 30

Anastasia added a comment to D58060: Fix diagnostic for addr spaces in reference binding.

@rjmccall or @ebevhan do you have any more feedback for this patch?

Tue, Apr 30, 9:33 AM
Anastasia created D61318: [Sema] Prevent binding references with mismatching address spaces to temporaries.
Tue, Apr 30, 9:14 AM
Anastasia added a comment to D61274: [Sema][AST] Explicit visibility for OpenCL/CUDA kernels/variables.

Seems reasonable for OpenCL kernels. You might want to add an AST dump test to check that the visibility is being set correctly in case it's being printed in AST.

Tue, Apr 30, 5:00 AM · Restricted Project
Anastasia created D61304: [OpenCL][PR41609] Deduce static data members to __global addr space.
Tue, Apr 30, 4:40 AM · Restricted Project

Sat, Apr 27

Anastasia added a comment to D60455: [SYCL] Add support for SYCL device attributes.

I think potentially reusing OpenCL features is desirable since the device code of SYCL is largely OpenCL. However I don't think we are clear enough about the overall device compilation flow of SYCL and I can easily suggest a number of different approaches including those that don't modify compiler at all. :) I am afraid until we have the big picture clear it will be hard to make any sensible decisions.

Sat, Apr 27, 12:19 AM · Restricted Project

Fri, Apr 26

Anastasia added a comment to D60763: Prototype OpenCL BIFs using Tablegen.

I also think we could reduce the size of the tables.
To sum up how this is working right now:

  1. The isOpenCLBuiltin(char* functionName) funcion is called to determine if a function is part of OpenCL builtin functions. If so, it returns its associated pair (index, number of signatures) in the OpenCLBuiltinDecl, otherwise it returns 0.
  2. The OpenCLBuiltinDecl table is storing, for each OpenCL builtin function, the list of the possible signature associated for this function (e.g. cos can be "float cos(float)", "double cos(double)", ...). The available signatures are stored in the OpenCLSignature table. In the OpenCLBuiltinDecl are stored the indexes corresponding to the possible signature for each function.
  3. The OpenCLSignature is storing the possible signatures.

E.g.: For the prototype float cos(float):

  1. isOpenCLBuiltin("cos") is returning the pair (345, 18)
  2. OpenCLBuiltinDecl[345] to OpenCLBuiltinDecl[345+18] are the available signatures of the builtin function "cos". Let say OpenCLBuiltinDecl[346] is containing our "float cos(float)" prototype. OpenCLBuiltinDecl[346] is containing the pair (123, 2), indexing the OpenCLSignature table and how many entries we have to read.
  3. OpenCLSignature[123] is storing the return type of the function, so the "float" type. OpenCLSignature[124] is containing the type of the first argument, so the float type again. We are not looking further in the table because we are only looking for 2 types.

    ---------------------------------------------------------------------------------------

    In the "float cos(float)" prototype, the information about the "float" type is duplicated. Plus, this "float" type is also the same as in the "float sin(float)" function. A third table, storing the different types, would discard duplicated definitions of types. The OpenCLSignature would store indexes of the required types, and the third table the type itself. This third table would be called OpenCLTypes, and would be as:

    ` struct OpenCLTypes { A type (e.g.: float, int, ...) OpenCLTypeID ID; Size of the vector (if applicable) unsigned VectorWidth; 0 if the type is not a pointer unsigned isPointer; Address space of the pointer (if applicable) clang::LangAS AS; } ` and OpenCLSignature:

    ` struct OpenCLSignature { unsigned OpenCLTypesIndex } ` --------------------------------------------------------------------------------------- Another way to save space would be to group functions. The "sin" and "cos" functions are similar (identical, we could say), regarding their use/ signature. However, they have two distinct lists of signatures in the OpenCLBuiltinDecl table. The consequence would be that isOpenCLBuiltin("cos") and isOpenCLBuiltin("sin") would give the same output. Such group of functions could be created manually by adding an attribute in the OpenCLBuiltins.td file, or automatically generated in the ClangOpenCLBuiltinEmitter.cpp file. The first solution would however make the feature potentially less understandable/ more complex to add new functions. The second solution is complex to implement/ could require a lot of time to process.
Fri, Apr 26, 9:04 PM · Restricted Project

Wed, Apr 24

Anastasia accepted D60764: Add clang cc1 option to generate OpenCL builtin functions.

LGTM! Thanks!

Wed, Apr 24, 11:44 PM · Restricted Project

Apr 19 2019

Anastasia added a comment to D60455: [SYCL] Add support for SYCL device attributes.

SYCL is similar to OpenMP 5 for C++, where you use only C++ classes instead of #pragma. So it is quite C++-friendlier than OpenMP.

Apr 19 2019, 5:38 AM · Restricted Project

Apr 18 2019

Anastasia accepted D60835: [Serialization] Stable serialization order for OpenCLTypeExtMap and OpenCLDeclExtMap.

We could still go for something like the following but accept that some a random failure might happen not necessarily on a commit that introduces it?

I am not sure that this is needed. Non-deterministic tests are really annoying.

Yes. I think this might be a wider topic to address if we happen to have more issues of this kind.

Apr 18 2019, 6:58 AM · Restricted Project, Restricted Project
Anastasia added a comment to D60602: [InferAddressSpaces] Add AS parameter to the pass factory.

This doesn't fix anything broken in LLVM but I did check that it didn't break any of the existing tests either :).

I did think about adding a test but in this case testing the new parameter would require a full C++ app that makes use of the pass and grep'ing in the tests returned zero matches for similar tests.

Apr 18 2019, 6:38 AM · Restricted Project
Anastasia added a comment to D60602: [InferAddressSpaces] Add AS parameter to the pass factory.

The pass is used by the AMDGPU and NVPTX backends. All their tests are passing with this change.

Yes, but they were passing before? Or is this fixing any bug in the existing tests?

If the change is adding functionality it is a good practice to add a test demonstrating its use. Otherwise, someone can commit something later that break the functionality and we might not have any way to detect that it's broken.

Apr 18 2019, 6:18 AM · Restricted Project
Anastasia added a comment to D60602: [InferAddressSpaces] Add AS parameter to the pass factory.

The pass is used by the AMDGPU and NVPTX backends. All their tests are passing with this change.

Apr 18 2019, 6:15 AM · Restricted Project
Anastasia added a comment to D60602: [InferAddressSpaces] Add AS parameter to the pass factory.

Btw, how is this tested? Do we have an existing test that exercises this already?

Apr 18 2019, 4:43 AM · Restricted Project
Anastasia added a comment to D60835: [Serialization] Stable serialization order for OpenCLTypeExtMap and OpenCLDeclExtMap.

By the way, I am wondering about how much this is tested. I did look quickly in test/PCH and it appears that there are only 3 (short) tests : ocl_types.cl, opencl-extensions.cl and no-validate-pch.

Apr 18 2019, 4:12 AM · Restricted Project, Restricted Project

Apr 17 2019

Anastasia added a comment to D60455: [SYCL] Add support for SYCL device attributes.

I think what we are missing currently is a thorough analysis/comparison between SYCL device mode and OpenCL kernel language mode to understand what's the best implementation strategy. That would apply to many other features: kernel function restrictions, address spaces, vectors, special types, etc.

That would make definitely sense when we target OpenCL.

I still see no point in polluting our code base with extra code that just does the same thing. It will save us a lot of time to just work cooperatively on the same problem and even improve readability of the code. But of course this can only be done if there is no need to diverge the implementation significantly.

Yes. Probably that even when the target is not OpenCL, the general principles remain similar. Probably the same for CUDA & OpenMP 5 too...

Apr 17 2019, 9:05 AM · Restricted Project
Anastasia added a comment to D60778: Make precompiled headers reproducible by switching OpenCL extension to std::map.

Maybe, but is there an equivalent to getTypeID for declarations ?

Apr 17 2019, 8:56 AM
Anastasia added a comment to D60778: Make precompiled headers reproducible by switching OpenCL extension to std::map.

Not sure how to test this though? I guess we can trust that std::map is always sorted just as it says in its description.

You could add a test that contains several entries in OpenCLTypeExtMap and several entries in OpenCLDeclExtMap.
In that test, write te PCH, and then apply llvm-bcanalyzer to that PCH. It should dump it the bitcode in textual dump.
And then simply apply FileCheck to that textual dump, with CHECK lines requiring the specific order of these entries.
If the order is not deterministic in PCH, then that test would (should!) fail sporadically.
At least that is my guess.

Ok I could do that, but I guess it can then fail on the commits that didn't actually trigger the issue. Would it not be a problem?

Why would it fail if this fixes the issue?

Apr 17 2019, 8:54 AM
Anastasia added a comment to D60778: Make precompiled headers reproducible by switching OpenCL extension to std::map.

Not sure how to test this though? I guess we can trust that std::map is always sorted just as it says in its description.

You could add a test that contains several entries in OpenCLTypeExtMap and several entries in OpenCLDeclExtMap.
In that test, write te PCH, and then apply llvm-bcanalyzer to that PCH. It should dump it the bitcode in textual dump.
And then simply apply FileCheck to that textual dump, with CHECK lines requiring the specific order of these entries.
If the order is not deterministic in PCH, then that test would (should!) fail sporadically.
At least that is my guess.

Apr 17 2019, 6:26 AM
Anastasia added a comment to D60455: [SYCL] Add support for SYCL device attributes.

Just to understand how this will work. I would imagine you can have a device function definition preceding its use. When the function is being parsed it's not known yet whether it will be called from the device or not, so it won't be possible to set the language mode correctly and hence provide the right diagnostics. So is the plan to launch a separate parsing phase then just to extract the call graph and annotate the device functions?

I'm not an expert in clang terminology but I will try briefly explain our current implementation approach.
In SYCL all kernel functions should be template functions so these functions have a deferred instantiation. If we found that we instantiated a sycl kernel function - we add it to a special array with sycl device functions (you can see the corresponding code here and here, actually AddSyclKernel adds declaration to the special array inside the Sema). After performing
pending instantiations we run a recursive AST visitor for each SYCL kernel to mark all device functions and add them to a special array with SYCL device functions (here we start traverse AST from MarkDevice function, code of MarkDevice is here).
To get a correct set of SYCL device functions in produced module we added a check for all declarations inside the CodeGen on sycl_device attribute existence - so it will ignore declarations without sycl_device attribute if we are compiling for SYCL device (code is here). But with this check it's possible situation when function was parsed and ignored by the CodeGen before we added sycl_device attribute to it so we added yet another parsing action inside the clang::ParseAST to generate code for all SYCL device functions from the special array (code is here).

Apr 17 2019, 4:36 AM · Restricted Project
Anastasia added a comment to D60455: [SYCL] Add support for SYCL device attributes.

Ok, my question is whether you are planning to duplicate the same logic as for OpenCL kernel which doesn't really seem like an ideal design choice. Is this the only difference then we can simply add an extra check for SYCL compilation mode in this template handling case. The overall interaction between OpenCL and SYCL implementation is still a very big unknown to me so it's not very easy to judge about the implementations details...

Of course, if nothing prevents us to re-use OpenCL kernel attribute for SYCL I assume it would be good idea.
But I'm thinking about the situation with https://reviews.llvm.org/D60454 . If we re-use OpenCL kernel attributes - we affected by OpenCL-related changes and OpenCL-related changes shouldn't violate SYCL semantics. Will it be usable for SYCL/OpenCL clang developers? @bader , what do you think about it?

I also think it's worth trying. We should be able to cover "SYCL semantics" with LIT test to avoid regressions by OpenCL related changes. E.g. add a test case checking that -fsycl-is-device option disables restriction on applying __kernel to template functions.
I want to confirm that everyone is okay to enable __kernel keyword for SYCL extension and cover SYCL use cases with additional regression tests. IIRC, on yesterday call, @keryell, said that having SYCL specific attributes useful for separation of concerns.

I'm not comfortable with that decision unless the attribute semantics are sufficiently related to justify it. If we're just going to have a lot of KernelAttr->isSYCL() vs KernelAttr->isOpenCL() accessor calls, it may make more sense to use separate semantic attributes (even if they share spellings), though then I'd be curious how a user combines OpenCL and SYCL attributes.

Apr 17 2019, 4:18 AM · Restricted Project

Apr 16 2019

Anastasia added a comment to D60379: Make precompiled headers reproducible.

FYI, I switched to std::map in this review:
https://reviews.llvm.org/D60778

Apr 16 2019, 8:55 AM · Restricted Project
Anastasia created D60778: Make precompiled headers reproducible by switching OpenCL extension to std::map.
Apr 16 2019, 8:51 AM
Anastasia accepted D60775: [libclang] Expose ext_vector_type.

LGTM! Thanks!

Apr 16 2019, 7:47 AM · Restricted Project, Restricted Project
Anastasia added inline comments to D60764: Add clang cc1 option to generate OpenCL builtin functions.
Apr 16 2019, 7:29 AM · Restricted Project
Anastasia updated subscribers of D60763: Prototype OpenCL BIFs using Tablegen.

Not related to this patch but it might be good to start thinking about testing this functionality properly. In the past we haven't tested the header because it would take a lot of testing time. So I would suggest we keep a light minimal basic testing in regular clang tests as is, but then we either add special build target to tests extra header functionality or add special cmake option (i.e. LLVM_INCLUDE_OPENCL_BIFS_TESTS) that would enable such testing with regular check-clang. I assume the latter can even be used to avoid generating the extra tables from TableGen minimizing the impact on the clang binary size for the situations OpenCL isn't required in the installation.

Apr 16 2019, 7:23 AM · Restricted Project
Anastasia added a reviewer for D60763: Prototype OpenCL BIFs using Tablegen: asavonic.
Apr 16 2019, 4:51 AM · Restricted Project
Anastasia added inline comments to D60455: [SYCL] Add support for SYCL device attributes.
Apr 16 2019, 2:48 AM · Restricted Project

Apr 15 2019

Anastasia added a comment to D58236: Make address space conversions a bit stricter..

Well, it doesn't seem to me like there is consensus on prohibiting nested address space conversion like this.

I can simply redo the patch to only include the bugfix on implicit conversions and drop the nesting level checks.

Apr 15 2019, 10:11 AM · Restricted Project, Restricted Project
Anastasia added a comment to D60455: [SYCL] Add support for SYCL device attributes.

Our current approach is to add an attribute, which SYCL runtime will use to mark code passed to cl::sycl::handler::parallel_for as "kernel functions". Obviously runtime library can't mark foo as "device" code - this is a compiler job: to traverse all symbols accessible from kernel functions and add them to the "device part" of the code.

Here is a link to the code in the SYCL runtime using sycl_kernel attribute: https://github.com/intel/llvm/blob/sycl/sycl/include/CL/sycl/handler.hpp#L267

I'm quite sure something similar should happen for other "single source" programming models like OpenMP/CUDA, except these attributes are exposed to the user and there is a specific requirement on attributes/pragma/keyword names.

Apr 15 2019, 9:12 AM · Restricted Project

Apr 9 2019

Anastasia requested changes to D60455: [SYCL] Add support for SYCL device attributes.

I was just wondering since SYCL is intended to be a single source C++ for OpenCL and this attribute is for internal use is it possible to just reuse existing OpenCL kernel attribute?

Apr 9 2019, 8:47 AM · Restricted Project
Anastasia added inline comments to D60379: Make precompiled headers reproducible.
Apr 9 2019, 4:58 AM · Restricted Project
Anastasia created D60454: [OpenCL] Prevent mangling kernel functions.
Apr 9 2019, 4:47 AM · Restricted Project

Apr 8 2019

Anastasia added inline comments to D60379: Make precompiled headers reproducible.
Apr 8 2019, 10:27 AM · Restricted Project
Anastasia accepted D59985: [OpenCL] Re-fix invalid address space generation for clk_event_t arguments of enqueue_kernel builtin function.

LGTM! Great! Thanks!

Apr 8 2019, 10:19 AM · Restricted Project, Restricted Project

Apr 4 2019

Anastasia committed rG9b4c6b8c7b18: [PR41157][OpenCL] Prevent implicit init of local addr space var in C++ mode. (authored by Anastasia).
[PR41157][OpenCL] Prevent implicit init of local addr space var in C++ mode.
Apr 4 2019, 4:08 AM
Anastasia added a comment to D59985: [OpenCL] Re-fix invalid address space generation for clk_event_t arguments of enqueue_kernel builtin function.

Alternative way to fix it is to use isNullPointerConstant like we do in SemaOpenCLBuiltinEnqueueKernel. So in case we have a zero literal value we can emit ConstantPointerNull directly, without EmitScalarExpr .

Ok and if it's not 0 the code gets rejected?

It must be rejected by SemaOpenCLBuiltinEnqueueKernel and it is already done https://godbolt.org/z/MFN3VU

Apr 4 2019, 3:50 AM · Restricted Project, Restricted Project
Anastasia committed rG094c72660a30: [PR41276] Fixed incorrect generation of addr space cast for 'this' in C++. (authored by Anastasia).
[PR41276] Fixed incorrect generation of addr space cast for 'this' in C++.
Apr 4 2019, 3:49 AM
Anastasia added a comment to D59985: [OpenCL] Re-fix invalid address space generation for clk_event_t arguments of enqueue_kernel builtin function.

Alternative way to fix it is to use isNullPointerConstant like we do in SemaOpenCLBuiltinEnqueueKernel. So in case we have a zero literal value we can emit ConstantPointerNull directly, without EmitScalarExpr .

Apr 4 2019, 2:47 AM · Restricted Project, Restricted Project

Apr 3 2019

Anastasia updated the diff for D59646: [PR40778][PR41157][OpenCL] Prevent implicit initialization of local address space objects.

Improved comment about initializers in __local addr space.

Apr 3 2019, 5:32 AM · Restricted Project
Anastasia added inline comments to D60193: [OpenCL] Added addrspace_cast operator.
Apr 3 2019, 5:05 AM
Anastasia added inline comments to D59985: [OpenCL] Re-fix invalid address space generation for clk_event_t arguments of enqueue_kernel builtin function.
Apr 3 2019, 4:52 AM · Restricted Project, Restricted Project
Anastasia created D60193: [OpenCL] Added addrspace_cast operator.
Apr 3 2019, 4:46 AM

Apr 2 2019

Anastasia updated the diff for D59988: [PR41276] Generate address space cast of 'this' for objects attributed by an address space in C++.
  • Use AggValueSlot in the constructor call generation to store/retrieve address space of 'this'.
  • Fixed detecting the address space conversion while performing qualification conversion.
Apr 2 2019, 7:09 AM · Restricted Project

Apr 1 2019

Anastasia added inline comments to D59985: [OpenCL] Re-fix invalid address space generation for clk_event_t arguments of enqueue_kernel builtin function.
Apr 1 2019, 3:53 AM · Restricted Project, Restricted Project
Anastasia removed a reviewer for D59988: [PR41276] Generate address space cast of 'this' for objects attributed by an address space in C++: brunodefraine.
Apr 1 2019, 3:22 AM · Restricted Project

Mar 29 2019

Anastasia added inline comments to D59988: [PR41276] Generate address space cast of 'this' for objects attributed by an address space in C++.
Mar 29 2019, 4:53 AM · Restricted Project
Anastasia created D59988: [PR41276] Generate address space cast of 'this' for objects attributed by an address space in C++.
Mar 29 2019, 4:49 AM · Restricted Project
Herald added a project to D53809: Fix invalid address space generation for clk_event_t: Restricted Project.
Mar 29 2019, 4:33 AM · Restricted Project

Mar 28 2019

Anastasia committed rG314fab6d7fa9: [PR41247] Fixed parsing of private keyword in C++. (authored by Anastasia).
[PR41247] Fixed parsing of private keyword in C++.
Mar 28 2019, 4:48 AM

Mar 27 2019

Anastasia updated the diff for D59874: [PR41247] Fixed parsing of private keyword in C++.

Extended test case and explained the intention of the test.

Mar 27 2019, 11:48 AM · Restricted Project
Anastasia added a comment to D59603: [PR40707][PR41011][OpenCL] Allow addr space spelling without double underscore in C++ mode.

Looks like this may cause some unexpected failures. See https://bugs.llvm.org/show_bug.cgi?id=41247 for more details.

Mar 27 2019, 6:02 AM · Restricted Project
Anastasia created D59874: [PR41247] Fixed parsing of private keyword in C++.
Mar 27 2019, 6:01 AM · Restricted Project

Mar 26 2019

Anastasia closed D59492: [OpenCL] Allow variadic macros as Clang feature.

Committed in r356987.

Mar 26 2019, 11:11 AM
Anastasia committed rG545652b96485: [OpenCL] Allow variadic macros as Clang feature. (authored by Anastasia).
[OpenCL] Allow variadic macros as Clang feature.
Mar 26 2019, 4:22 AM

Mar 25 2019

Anastasia added inline comments to D59646: [PR40778][PR41157][OpenCL] Prevent implicit initialization of local address space objects.
Mar 25 2019, 7:26 AM · Restricted Project
Anastasia added inline comments to D59646: [PR40778][PR41157][OpenCL] Prevent implicit initialization of local address space objects.
Mar 25 2019, 5:39 AM · Restricted Project
Anastasia updated the diff for D59646: [PR40778][PR41157][OpenCL] Prevent implicit initialization of local address space objects.

Fixed typo in the comment.

Mar 25 2019, 5:39 AM · Restricted Project
Anastasia added inline comments to D59492: [OpenCL] Allow variadic macros as Clang feature.
Mar 25 2019, 5:20 AM
Anastasia updated the diff for D59492: [OpenCL] Allow variadic macros as Clang feature.

Changed diagnostic wording to indicate that this feature is a Clang extension.

Mar 25 2019, 5:20 AM
Anastasia committed rG948e37c8ca8b: [OpenCL] Allow addr space spelling without __ prefix in C++. (authored by Anastasia).
[OpenCL] Allow addr space spelling without __ prefix in C++.
Mar 25 2019, 4:54 AM

Mar 21 2019

Anastasia added a comment to D57464: Generalize method overloading on addr spaces to C++.

I think I would lean towards the latter since it means less fudging around with a whole bunch of unrelated methods. Do @rjmccall or @rsmith have any further opinions on this?

Mar 21 2019, 10:33 AM
Anastasia added a comment to D58236: Make address space conversions a bit stricter..

Any more input on this?

I could redo the patch to simply fix the bug and not make the conversions stricter, if that's preferable.

Mar 21 2019, 10:30 AM · Restricted Project, Restricted Project
Anastasia created D59646: [PR40778][PR41157][OpenCL] Prevent implicit initialization of local address space objects.
Mar 21 2019, 8:20 AM · Restricted Project
Anastasia retitled D59646: [PR40778][PR41157][OpenCL] Prevent implicit initialization of local address space objects from [PR40778][PR41157] Prevent implicit initialization of local address space objects to [PR40778][PR41157][OpenCL] Prevent implicit initialization of local address space objects.
Mar 21 2019, 8:20 AM · Restricted Project
Anastasia added a comment to D59603: [PR40707][PR41011][OpenCL] Allow addr space spelling without double underscore in C++ mode.

I was a little bit worried about

struct top_level { int i; };

private ::top_level a;

but it should be fine because in that case we have a tok::coloncolon

Mar 21 2019, 4:43 AM · Restricted Project
Anastasia updated the diff for D59603: [PR40707][PR41011][OpenCL] Allow addr space spelling without double underscore in C++ mode.

Updated test to use root namespace (commented by Ronan).

Mar 21 2019, 4:39 AM · Restricted Project