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 (381 w, 3 d)

Recent Activity

Fri, Jun 24

Anastasia accepted D128436: [OpenCL] Remove fast_ half geometric builtins.

LGTM! Thanks

Fri, Jun 24, 11:15 AM · Restricted Project, Restricted Project
Anastasia accepted D128434: [OpenCL] Remove half scalar vload/vstore builtins.

LGTM! Thanks

Fri, Jun 24, 11:13 AM · Restricted Project, Restricted Project
Anastasia added a comment to D127579: [clang][WIP] add option to keep types of ptr args for non-kernel functions in metadata.

@Anastasia Thanks, that does sound like a legitimate reason to include the information. I want to double check though, does linking the modules actually fail if the functions have signatures that differ only by pointer types? At least for normal LLVM IR this would work fine, and would just result in the insertion of a bitcast during linking (and then typically the bitcast would get shifted from the called function to the call arguments later).

@nikic If I use spirv-link with two modules that have mismatching pointee type in a function parameter I get an error:

error: 0: Type mismatch on symbol "foo" between imported variable/function %6 and exported variable/function %17.

The way I understand a bitcast instruction in SPIR-V (OpBitcast in https://www.khronos.org/registry/SPIR-V/specs/unified1/SPIRV.html#_conversion_instructions) is that it can only apply to pointer types which are distinct from function types. Note that I believe that function pointers are illegal, at least we disallow them in OpenCL.

Okay ... can we maybe turn this around then? Always emit function parameters as i8* and bitcast them as needed, even if it is possible to guess a better type based on the definition? (Let's ignore the image type case here, which seems to have different requirements from the rest.)

So where would bitcasts be emitted to reconstruct the function prototypes correctly?

The bitcasts would be in the definition only, when the pointer arguments actually get used in a typed manner. The prototype would always include i8* arguments only.

Fri, Jun 24, 11:10 AM · Restricted Project, Restricted Project

Tue, Jun 21

Anastasia accepted D127961: [OpenCL] Reduce emitting candidate notes for builtins.

LGTM! Thanks

Tue, Jun 21, 2:17 AM · Restricted Project, Restricted Project
Anastasia added a comment to D128012: [HLSL] Add ExternalSemaSource & vector alias.

aha, are you having a lot of these types then? Ok that sounds similar problem to OpenCL builtin functions as we had to go away from a simple header include due to long parsing time.

Very similar problem. We have a two-part problem where we both have enough of them that parsing a header would be too slow _and_ the HLSL language actually can't represent them. The later we're working to fix. My hope is that eventually we can write the HLSL libraries in HLSL and use a pre-compilation step (PCH or Module-style) to get past the performance issues.

Tue, Jun 21, 2:15 AM · Restricted Project, Restricted Project
Anastasia added a comment to D127579: [clang][WIP] add option to keep types of ptr args for non-kernel functions in metadata.

@Anastasia Thanks, that does sound like a legitimate reason to include the information. I want to double check though, does linking the modules actually fail if the functions have signatures that differ only by pointer types? At least for normal LLVM IR this would work fine, and would just result in the insertion of a bitcast during linking (and then typically the bitcast would get shifted from the called function to the call arguments later).

@nikic If I use spirv-link with two modules that have mismatching pointee type in a function parameter I get an error:

error: 0: Type mismatch on symbol "foo" between imported variable/function %6 and exported variable/function %17.

The way I understand a bitcast instruction in SPIR-V (OpBitcast in https://www.khronos.org/registry/SPIR-V/specs/unified1/SPIRV.html#_conversion_instructions) is that it can only apply to pointer types which are distinct from function types. Note that I believe that function pointers are illegal, at least we disallow them in OpenCL.

Okay ... can we maybe turn this around then? Always emit function parameters as i8* and bitcast them as needed, even if it is possible to guess a better type based on the definition? (Let's ignore the image type case here, which seems to have different requirements from the rest.)

Tue, Jun 21, 2:13 AM · Restricted Project, Restricted Project

Fri, Jun 17

Anastasia added a comment to D127579: [clang][WIP] add option to keep types of ptr args for non-kernel functions in metadata.

Btw do we have any issue tracking this failure, I am having a feeling that the translation of special types might be related to https://github.com/llvm/llvm-project/issues/55121? I am not aware at which point libclc is expected to be linked for SPIR-V but if linking happens at SPIR-V binary level it is likely going to face the other problem I described here.

Fri, Jun 17, 7:58 AM · Restricted Project, Restricted Project
Anastasia added a comment to D128012: [HLSL] Add ExternalSemaSource & vector alias.

For just the type alias, adding it during parsing would probably work. Where things get more complicated is we have a whole mess of other data types that we’ll need to add too eventually. One of the advantages of using an external sema source is that we can create incomplete records during initialization, and complete them when the sema source gets the callback from lookup. That gives us cheap and efficient lazy-initialization of our buffer types which (in DXC) has a huge impact on compile time.

Fri, Jun 17, 7:45 AM · Restricted Project, Restricted Project
Anastasia added a comment to D128012: [HLSL] Add ExternalSemaSource & vector alias.

I wonder if accepting type alias with identifier vector while parsing would be simpler to implement? Then you could just add those type aliases into the internal header. I assume user code is not allowed to redefine those i.e. it would result in undefined behavior or something like for most of C++ standard library features...

Fri, Jun 17, 4:04 AM · Restricted Project, Restricted Project
Anastasia added inline comments to D127961: [OpenCL] Reduce emitting candidate notes for builtins.
Fri, Jun 17, 3:32 AM · Restricted Project, Restricted Project
Anastasia added a comment to D127579: [clang][WIP] add option to keep types of ptr args for non-kernel functions in metadata.

Clang is only permitted to encode pointer type information if that pointer type has some kind of direct semantic meaning -- say, if the pointer element type affects the call ABI in some way. If it does not have direct semantic meaning, then deriving the pointer type based on usage (as DXIL does) and using that for emission would be right approach.

Fri, Jun 17, 1:32 AM · Restricted Project, Restricted Project
Anastasia added a comment to D127579: [clang][WIP] add option to keep types of ptr args for non-kernel functions in metadata.

If I'm understanding correctly, previously the LLVM pointee type was used to represent a frontend OpenCL type. An alternative to marking everything with elementtype or adding metadata, can something be done in the frontend to mangle the function name with OpenCL pointee types, then the backend translates that to the proper SPIRV types?

Fri, Jun 17, 1:29 AM · Restricted Project, Restricted Project

Thu, Jun 16

Anastasia added a comment to D127579: [clang][WIP] add option to keep types of ptr args for non-kernel functions in metadata.

The way I understand a bitcast instruction in SPIR-V (OpBitcast in https://www.khronos.org/registry/SPIR-V/specs/unified1/SPIRV.html#_conversion_instructions) is that it can only apply to pointer types which are distinct from function types. Note that I believe that function pointers are illegal, at least we disallow them in OpenCL.

FYI: we are experimenting with function pointers on Intel HW programmed via SPIR-V. Extension draft - https://github.com/intel/llvm/blob/sycl/sycl/doc/design/spirv-extensions/SPV_INTEL_function_pointers.asciidoc.

Thu, Jun 16, 5:50 AM · Restricted Project, Restricted Project
Anastasia added a comment to D127579: [clang][WIP] add option to keep types of ptr args for non-kernel functions in metadata.

@Anastasia Thanks, that does sound like a legitimate reason to include the information. I want to double check though, does linking the modules actually fail if the functions have signatures that differ only by pointer types? At least for normal LLVM IR this would work fine, and would just result in the insertion of a bitcast during linking (and then typically the bitcast would get shifted from the called function to the call arguments later).

Thu, Jun 16, 5:25 AM · Restricted Project, Restricted Project

Wed, Jun 15

Anastasia added a comment to D127579: [clang][WIP] add option to keep types of ptr args for non-kernel functions in metadata.

@nikic the most important thing you need to know about SPIR-V is that it is a virtual ISA based on LLVM IR. The ISA itself encodes types for pointers just like LLVM IR would.

Okay ... I guess my next question would be how the SPIR-V situation differs from the DXIL situation. For DXIL you have already implemented code to "guess" pointer types insofar as they are relevant -- does SPIR-V need something more than that?

Clang is only permitted to encode pointer type information if that pointer type has some kind of direct semantic meaning -- say, if the pointer element type affects the call ABI in some way. If it does not have direct semantic meaning, then deriving the pointer type based on usage (as DXIL does) and using that for emission would be right approach.

Wed, Jun 15, 9:09 AM · Restricted Project, Restricted Project
Anastasia added a comment to D127579: [clang][WIP] add option to keep types of ptr args for non-kernel functions in metadata.

The patch seems very straight forward. I wonder though if we should always enable this for SPIR-V w/o adding any flag until the untyped pointers are added to SPIR-V so it will become a target setting controlled instead of a flag?

Wed, Jun 15, 7:42 AM · Restricted Project, Restricted Project

Thu, Jun 2

Anastasia accepted D126857: [HLSL] Add WaveActiveCountBits as Langugage builtin function for HLSL.

LGTM! Thanks

Thu, Jun 2, 12:56 PM · Restricted Project, Restricted Project
Anastasia added inline comments to D126857: [HLSL] Add WaveActiveCountBits as Langugage builtin function for HLSL.
Thu, Jun 2, 12:26 PM · Restricted Project, Restricted Project
Anastasia added inline comments to D126857: [HLSL] Add WaveActiveCountBits as Langugage builtin function for HLSL.
Thu, Jun 2, 11:56 AM · Restricted Project, Restricted Project
Anastasia added a comment to D126857: [HLSL] Add WaveActiveCountBits as Langugage builtin function for HLSL.

ok, so the reason you are adding this to clang is that it needs to be mapped into a target intrinsic?

Thu, Jun 2, 11:47 AM · Restricted Project, Restricted Project
Anastasia added a comment to D124753: [HLSL] Set main as default entry..

From the current change it seems to me that what you need to be testing is a just that the frontend options are being passed correctly? This should then be a driver test with -### checking for the options to be set for the frontend invocation...

There's already a driver test with '-###' in https://reviews.llvm.org/D124751#change-af6Z62NjlfGb

This test doesn't seem to correspond to the change being added as you are changing the command-line flags. You don't actually add/generate any attributes in this patch.

Sorry to make things confusing. I should not split default value for -E option as a separate PR :(

There's dxc_E.hlsl (https://reviews.llvm.org/D124751#change-af6Z62NjlfGb) in https://reviews.llvm.org/D124751 where the -E option is added.
dxc_E.hlsl will test -E option translated into -hlsl-entry for cc1.

There was a test for codeGen of -E option in https://reviews.llvm.org/D124752, but I removed it because it is to almost the same as https://reviews.llvm.org/D124752#change-w4NWvaT68Dhk which test codeGen for ShaderAttr.

And the entry_default.hlsl in current PR test codeGen for the default main and -E option.

I can add a separate codeGen test for -E option if that's what you thought is missing.

Thu, Jun 2, 11:45 AM · Restricted Project, Restricted Project

Wed, Jun 1

Anastasia accepted D126660: [OpenCL] Reword unknown extension pragma diagnostic.

Ok, makes sense! Thanks!

Wed, Jun 1, 8:32 AM · Restricted Project, Restricted Project
Anastasia accepted D124752: [HLSL] clang codeGen for HLSLShaderAttr..

LGTM! Thanks

Wed, Jun 1, 8:25 AM · Restricted Project, Restricted Project
Anastasia added a comment to D124753: [HLSL] Set main as default entry..

From the current change it seems to me that what you need to be testing is a just that the frontend options are being passed correctly? This should then be a driver test with -### checking for the options to be set for the frontend invocation...

There's already a driver test with '-###' in https://reviews.llvm.org/D124751#change-af6Z62NjlfGb

Wed, Jun 1, 8:23 AM · Restricted Project, Restricted Project

Mon, May 30

Anastasia accepted D125052: [HLSL] Enable vector types for hlsl..

Please make sure to reflect in your commit message that you are adding the clang internal headers too.

Mon, May 30, 7:50 AM · Restricted Project, Restricted Project
Anastasia added a comment to D124753: [HLSL] Set main as default entry..

From the current change it seems to me that what you need to be testing is a just that the frontend options are being passed correctly? This should then be a driver test with -### checking for the options to be set for the frontend invocation...

Mon, May 30, 7:46 AM · Restricted Project, Restricted Project
Anastasia added inline comments to D124752: [HLSL] clang codeGen for HLSLShaderAttr..
Mon, May 30, 7:39 AM · Restricted Project, Restricted Project

May 27 2022

Anastasia committed rG7df25978ef78: [Doc][OpenCL] Misc wording improvements for SPIR-V (authored by Anastasia).
[Doc][OpenCL] Misc wording improvements for SPIR-V
May 27 2022, 3:14 AM · Restricted Project, Restricted Project, Restricted Project

May 26 2022

Anastasia added inline comments to D124752: [HLSL] clang codeGen for HLSLShaderAttr..
May 26 2022, 11:12 AM · Restricted Project, Restricted Project
Anastasia added inline comments to D124752: [HLSL] clang codeGen for HLSLShaderAttr..
May 26 2022, 11:09 AM · Restricted Project, Restricted Project
Anastasia added inline comments to D124753: [HLSL] Set main as default entry..
May 26 2022, 10:48 AM · Restricted Project, Restricted Project
Anastasia added inline comments to D125052: [HLSL] Enable vector types for hlsl..
May 26 2022, 10:45 AM · Restricted Project, Restricted Project
Anastasia committed rG3087afb421bf: [OpenCL][Doc] Misc improvements related to SPIR-V support. (authored by Anastasia).
[OpenCL][Doc] Misc improvements related to SPIR-V support.
May 26 2022, 7:55 AM · Restricted Project, Restricted Project

May 25 2022

Anastasia committed rG730dc4e9bce8: [Clang] Added options for integrated backend. (authored by Anastasia).
[Clang] Added options for integrated backend.
May 25 2022, 4:13 AM · Restricted Project, Restricted Project
Anastasia closed D125679: [Clang] Added options for integrated backend only used for SPIR-V for now.
May 25 2022, 4:13 AM · Restricted Project, Restricted Project

May 24 2022

Anastasia added inline comments to D125052: [HLSL] Enable vector types for hlsl..
May 24 2022, 7:22 AM · Restricted Project, Restricted Project
Anastasia added inline comments to D124753: [HLSL] Set main as default entry..
May 24 2022, 5:54 AM · Restricted Project, Restricted Project
Anastasia added inline comments to D124752: [HLSL] clang codeGen for HLSLShaderAttr..
May 24 2022, 5:51 AM · Restricted Project, Restricted Project
Anastasia committed rGd61ded1034bb: [OpenCL] Make -cl-ext a driver option. (authored by Anastasia).
[OpenCL] Make -cl-ext a driver option.
May 24 2022, 3:36 AM · Restricted Project, Restricted Project
Anastasia closed D125243: [OpenCL] Make -cl-ext a driver option.
May 24 2022, 3:36 AM · Restricted Project, Restricted Project

May 23 2022

Anastasia committed rG72832efc941a: [SPIR-V] Allow setting SPIR-V version via target triple. (authored by Anastasia).
[SPIR-V] Allow setting SPIR-V version via target triple.
May 23 2022, 6:24 AM · Restricted Project, Restricted Project
Anastasia closed D124776: [SPIR-V] Allow setting SPIR-V version via target triple.
May 23 2022, 6:24 AM · Restricted Project, Restricted Project

May 18 2022

Anastasia requested changes to D124382: [Clang] Recognize target address space in superset calculation.

I feel that to progress further on this change, it would be good to get details about the use cases and the limitations first.

May 18 2022, 4:09 PM · Restricted Project, Restricted Project
Anastasia added a comment to D124382: [Clang] Recognize target address space in superset calculation.

And I think we could add this feature in a very light way for example by reserving the numbers from the clang LangAS enum to be used with the language address spaces in the prototypes of builtins.

I'm not sure I understand how that would look, could you please elaborate?

May 18 2022, 10:40 AM · Restricted Project, Restricted Project

May 16 2022

Anastasia added a comment to D110685: [HIPSPV][4/4] Add option to use llc to emit SPIR-V.

Actually after some more thinking I have decided to go for a generic flag and I have created the review for it: https://reviews.llvm.org/D125679. Perhaps you can built your functionality on top of it when you get to it.

May 16 2022, 4:11 AM · Restricted Project, Restricted Project
Anastasia updated the diff for D124776: [SPIR-V] Allow setting SPIR-V version via target triple.

Fixed typo in docs

May 16 2022, 4:05 AM · Restricted Project, Restricted Project
Anastasia requested review of D125679: [Clang] Added options for integrated backend only used for SPIR-V for now.
May 16 2022, 4:03 AM · Restricted Project, Restricted Project
Anastasia accepted D125208: [OpenCL] Fix __remove_address_space documentation code example.

LGTM! Thanks!

May 16 2022, 3:20 AM · Restricted Project, Restricted Project
Anastasia accepted D125401: [OpenCL] Do not guard vload/store_half builtins.

LGTM! Thanks

May 16 2022, 3:18 AM · Restricted Project, Restricted Project
Anastasia added inline comments to D125243: [OpenCL] Make -cl-ext a driver option.
May 16 2022, 3:09 AM · Restricted Project, Restricted Project

May 9 2022

Anastasia updated the diff for D124776: [SPIR-V] Allow setting SPIR-V version via target triple.
  • Added v1.5.
  • Changed wording in documentation.
May 9 2022, 9:58 AM · Restricted Project, Restricted Project
Anastasia requested review of D125243: [OpenCL] Make -cl-ext a driver option.
May 9 2022, 9:54 AM · Restricted Project, Restricted Project
Anastasia committed rGcd99227c7876: [Docs] Added my office hours. (authored by Anastasia).
[Docs] Added my office hours.
May 9 2022, 9:41 AM · Restricted Project, Restricted Project
Anastasia added a comment to D124776: [SPIR-V] Allow setting SPIR-V version via target triple.

@Anastasia, when are you going to add support for v1.5? We would use it in the SPIR-V backend.

May 9 2022, 8:41 AM · Restricted Project, Restricted Project

May 6 2022

Herald added a project to D110685: [HIPSPV][4/4] Add option to use llc to emit SPIR-V: Restricted Project.

I think it might make more sense to use the option -fno-llvm-spirv as we don't advertise llc as a tool to the developers but llvm-spirv we do.

May 6 2022, 1:59 AM · Restricted Project, Restricted Project
Anastasia updated the summary of D124776: [SPIR-V] Allow setting SPIR-V version via target triple.
May 6 2022, 1:51 AM · Restricted Project, Restricted Project

May 2 2022

Anastasia requested review of D124776: [SPIR-V] Allow setting SPIR-V version via target triple.
May 2 2022, 8:56 AM · Restricted Project, Restricted Project

Apr 28 2022

Anastasia added a comment to D124382: [Clang] Recognize target address space in superset calculation.

Can you provide an example of where it could be useful? Note that I feel that
such functionality could be implemented on top of full implementation of
target specific address space proposed in https://reviews.llvm.org/D62574.

The use case we had was when calling target builtin (that specifies address
space) from within OpenCL C. Currently this errors out, similarly, the explicit
type cast to address space yields an error
((__attribute__((address_space(3))) int *)&l_woof in the example below). This
is important for libclc, which is implemented in OpenCL C and deals directly
with target builtins.

__kernel void woof() {
  __local int l_woof;
  __nvvm_cp_async_mbarrier_arrive_shared(&l_woof);
}

I wasn't aware of that patch, sorry, I've not had a close look yet, but it
seems worryingly dated.

Apr 28 2022, 10:55 AM · Restricted Project, Restricted Project

Apr 26 2022

Anastasia added inline comments to D123498: [clang] Adding Platform/Architecture Specific Resource Header Installation Targets.
Apr 26 2022, 2:08 AM · Restricted Project, Restricted Project
Anastasia accepted D124256: [OpenCL] Add cl_khr_subgroup_rotate builtins.

LGTM! I imagine tablegen side is being tested automatically?

Apr 26 2022, 1:59 AM · Restricted Project, Restricted Project

Apr 25 2022

Anastasia added a comment to D124382: [Clang] Recognize target address space in superset calculation.

You should be able to provide an address space of the pointer using the number, see details in:
https://github.com/llvm/llvm-project/blob/main/clang/include/clang/Basic/Builtins.def#L65

Apr 25 2022, 1:36 PM · Restricted Project, Restricted Project

Apr 11 2022

Anastasia accepted D122728: [OpenCL] opencl-c.h: Add const to get_image_num_samples.

LGTM! Thanks

Apr 11 2022, 3:53 AM · Restricted Project, Restricted Project

Mar 8 2022

Anastasia committed rG481f6818670a: [AST] Fix typo in assert messages (authored by kuzkry).
[AST] Fix typo in assert messages
Mar 8 2022, 3:09 AM · Restricted Project
Anastasia closed D120221: [AST] Fix typos.
Mar 8 2022, 3:09 AM · Restricted Project, Restricted Project
Anastasia accepted D120221: [AST] Fix typos.

Can I ask you to deliver this one for me? I don't have permissions. My name and email in git format is "Krystian Kuzniarek <krystian.kuzniarek@gmail.com>"

Mar 8 2022, 3:01 AM · Restricted Project, Restricted Project

Feb 17 2022

Anastasia accepted D120032: [OpenCL] opencl-c.h: use uint/ulong consistently.

LGTM! Thanks

Feb 17 2022, 4:19 AM · Restricted Project

Feb 16 2022

Anastasia closed D119713: [Docs] Release 14 notes for SPIR-V in clang.

Committed https://github.com/llvm/llvm-project/commit/b54c95790b8a410b8eab1cbaae306cdc56a8e82a

Feb 16 2022, 4:34 AM
Anastasia closed D119710: [Docs][OpenCL] Release 14 notes .

Committed https://github.com/llvm/llvm-project/commit/e8712accba1150b94168d8ace6c7dd7b9612e174

Feb 16 2022, 4:33 AM
Anastasia committed rGfdd615d4f91f: [Docs][OpenCL] Update OpenCL 3.0 status in docs. (authored by Anastasia).
[Docs][OpenCL] Update OpenCL 3.0 status in docs.
Feb 16 2022, 4:33 AM
Anastasia closed D119719: [Docs][OpenCL] Update OpenCL 3.0 status.
Feb 16 2022, 4:32 AM · Restricted Project

Feb 15 2022

Anastasia accepted D119858: [OpenCL] Guard 64-bit atomic types.

I imagine this is another change to align with opencl-c.h?

Feb 15 2022, 10:02 AM · Restricted Project
Anastasia added a comment to D119398: [OpenCL] Guard atomic_double with cl_khr_int64_base_atomics and cl_khr_int64_extended_atomics.

Hi, @Anastasia. Please help to land it, thanks very much.

Feb 15 2022, 6:20 AM · Restricted Project
Anastasia added a comment to D119560: [OpenCL] opencl-c.h: remove arg names from atomics.

also makes the header no longer "claim" the identifiers "success",
"failure", "desired", "value" (such that you can compile with -Dvalue=...
when including the header for example, which currently breaks parsing
of the header).

I don't get what you mean by this. :)

Compiling a CL source file with e.g. clang -cl-std=CL2.0 -Xclang -finclude-default-header -cl-no-stdinc -Dvalue=1 clang/test/CodeGenOpenCL/as_type.cl gives lots of errors such as the following, because defining value as a macro (which is not a reserved identifier as far as I'm aware) collides with the argument names in the header:

In file included from <built-in>:1:
lib/clang/15.0.0/include/opencl-c.h:13277:58: error: expected ')'
void __ovld atomic_init(volatile atomic_int *object, int value);
                                                         ^
<command line>:1:15: note: expanded from here
#define value 1
Feb 15 2022, 6:08 AM · Restricted Project
Anastasia added a comment to D119398: [OpenCL] Guard atomic_double with cl_khr_int64_base_atomics and cl_khr_int64_extended_atomics.

Can this be committed now?

Feb 15 2022, 4:27 AM · Restricted Project
Anastasia accepted D119560: [OpenCL] opencl-c.h: remove arg names from atomics.

LGTM! Thanks

Feb 15 2022, 4:26 AM · Restricted Project
Anastasia added a comment to D119560: [OpenCL] opencl-c.h: remove arg names from atomics.

also makes the header no longer "claim" the identifiers "success",
"failure", "desired", "value" (such that you can compile with -Dvalue=...
when including the header for example, which currently breaks parsing
of the header).

Feb 15 2022, 4:26 AM · Restricted Project
Anastasia updated the diff for D119710: [Docs][OpenCL] Release 14 notes .
  • Fixed review comments
  • Updated OpenCL 3 state after backports
Feb 15 2022, 4:06 AM
Anastasia updated the diff for D119719: [Docs][OpenCL] Update OpenCL 3.0 status.
Feb 15 2022, 3:59 AM · Restricted Project

Feb 14 2022

Anastasia requested review of D119719: [Docs][OpenCL] Update OpenCL 3.0 status.
Feb 14 2022, 7:15 AM · Restricted Project
Anastasia added a comment to D119713: [Docs] Release 14 notes for SPIR-V in clang.

@linjamaki feel free to suggest other topics to document for the release.

Feb 14 2022, 6:25 AM
Anastasia requested review of D119713: [Docs] Release 14 notes for SPIR-V in clang.
Feb 14 2022, 6:24 AM
Anastasia updated subscribers of D119710: [Docs][OpenCL] Release 14 notes .
Feb 14 2022, 6:19 AM
Anastasia added inline comments to D119710: [Docs][OpenCL] Release 14 notes .
Feb 14 2022, 6:18 AM
Anastasia requested review of D119710: [Docs][OpenCL] Release 14 notes .
Feb 14 2022, 6:16 AM

Feb 11 2022

Anastasia accepted D118999: [OpenCL] Adjust diagnostic for subgroup support..

LGTM! Thanks

Feb 11 2022, 6:38 AM · Restricted Project
Anastasia accepted D118605: [OpenCL] Add support of language builtins for OpenCL C 3.0.

LGTM! Thanks

Feb 11 2022, 4:49 AM · Restricted Project
Anastasia added a comment to D118605: [OpenCL] Add support of language builtins for OpenCL C 3.0.

There are tests checking for this (e.g. clang/test/Frontend/opencl.cl), so we need this check to preserve the existing behavior indeed.

Thanks. The other test is SemaOpenCL/clang-builtin-version.cl.

But it might be worth asking someone outside of the OpenCL community whether it's desirable to use the LanguageID enum in this way.

I personally think this looks good now, for OpenCL in particularly, as it became version-agnostic (except for DSE). But we still are querying language options only, and we expect language options for generic AS, pipes and DSE to be immutable at this point.

Feb 11 2022, 4:48 AM · Restricted Project

Feb 10 2022

Anastasia accepted D119420: [OpenCL] Add OpenCL 3.0 atomics to -fdeclare-opencl-builtins.

LGTM! The rename can be made on commit.

Feb 10 2022, 4:58 AM · Restricted Project
Anastasia accepted D119398: [OpenCL] Guard atomic_double with cl_khr_int64_base_atomics and cl_khr_int64_extended_atomics.

This might interfere with https://reviews.llvm.org/D119420

Feb 10 2022, 4:45 AM · Restricted Project
Anastasia added inline comments to D119420: [OpenCL] Add OpenCL 3.0 atomics to -fdeclare-opencl-builtins.
Feb 10 2022, 3:49 AM · Restricted Project
Anastasia added inline comments to D118605: [OpenCL] Add support of language builtins for OpenCL C 3.0.
Feb 10 2022, 2:39 AM · Restricted Project

Feb 4 2022

Anastasia added a comment to D118999: [OpenCL] Adjust diagnostic for subgroup support..

It seems hard to find a direct mention in the spec, but in the API spec:

CL_DEVICE_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS: Is CL_TRUE if this device supports independent forward progress of sub-groups, CL_FALSE otherwise. This query must return CL_TRUE for devices that support the cl_khr_subgroups extension, and must return CL_FALSE for devices that do not support subgroups.

Feb 4 2022, 9:54 AM · Restricted Project
Anastasia added a comment to D119011: [clang] Cache OpenCL types.

Just to understand the intent - is this a performance optimization or functionality fix?

Feb 4 2022, 9:51 AM · Restricted Project
Anastasia added a comment to D118999: [OpenCL] Adjust diagnostic for subgroup support..

Btw can you point us to the spec reference please, I can't seem to find anything related to this: https://www.khronos.org/registry/OpenCL/specs/3.0-unified/html/OpenCL_C.html#subgroup-functions.

Feb 4 2022, 7:42 AM · Restricted Project
Anastasia added inline comments to D118605: [OpenCL] Add support of language builtins for OpenCL C 3.0.
Feb 4 2022, 7:33 AM · Restricted Project

Feb 3 2022

Anastasia added inline comments to D118894: [OpenCL] Mark kernel arguments as ABI aligned.
Feb 3 2022, 3:38 PM · Restricted Project
Anastasia added inline comments to D118894: [OpenCL] Mark kernel arguments as ABI aligned.
Feb 3 2022, 1:11 PM · Restricted Project
Anastasia added inline comments to D118894: [OpenCL] Mark kernel arguments as ABI aligned.
Feb 3 2022, 1:07 PM · Restricted Project
Anastasia added inline comments to D118894: [OpenCL] Mark kernel arguments as ABI aligned.
Feb 3 2022, 12:47 PM · Restricted Project
Anastasia added a comment to D115523: [OpenCL] Set external linkage for block enqueue kernels.

It is possible that block kernels are defined and invoked in static functions, therefore two block kernels in different TU's may have the same name. Making such kernels external may cause duplicate symbols.

Potentially we should append the name of the translation unit to all kernel wrapper names for the enqueued blocks to resolve this? For example, global constructors stubs are using such a similar naming scheme taken from the translation unit.

But the kernel function in OpenCL has to be globally visible and many tools have been built with this assumption. Additionally, some toolchains might require the enqueued kernels to be globally visible as well in order to access them as an execution entry point.

If we have to externalize such block kernels whose names are derived from variables with internal linkage, we may need a way to make the block kernel names unique.

One approach would be use MD5 hash of the file path and compile options, similar to https://github.com/llvm/llvm-project/blob/main/clang/lib/Driver/Driver.cpp#L2623

Ok, however it might be enough to do what C++ handling does with global ctors. Should we file a bug for this for now?

Do you mean the global ctors like this? https://github.com/llvm/llvm-project/blob/main/clang/test/CodeGenCXX/global-init.cpp#L201

They do not have unique names across TU's. They work because they have internal linkage. However OpenCL kernels have external linkage.

I suggest opening a bug to track this. Thanks.

Feb 3 2022, 8:50 AM · Restricted Project