Page MenuHomePhabricator

richardmembarth (Richard Membarth)
User

Projects

User does not belong to any projects.

User Details

User Since
Nov 8 2018, 5:39 AM (49 w, 3 d)

Recent Activity

Jun 5 2019

richardmembarth added a comment to D54258: [Clang] Fix pretty printing of CUDA address spaces.

Thanks for clarification and merging!

Jun 5 2019, 11:52 AM · Restricted Project, Restricted Project

Jun 4 2019

richardmembarth added a comment to D54258: [Clang] Fix pretty printing of CUDA address spaces.

Merging in two weeks is fine for me.

Jun 4 2019, 2:37 AM · Restricted Project, Restricted Project

May 31 2019

Herald added a project to D54258: [Clang] Fix pretty printing of CUDA address spaces: Restricted Project.

Do you know when this will be merged?

May 31 2019, 1:15 AM · Restricted Project, Restricted Project

Nov 13 2018

richardmembarth added a comment to D54258: [Clang] Fix pretty printing of CUDA address spaces.

In D54258#1297191, @Anastasia wrote:

I agree with the change itself... but it's quite annoying that such things can't be tested. :(

Nov 13 2018, 2:23 PM · Restricted Project, Restricted Project
richardmembarth added a comment to D54258: [Clang] Fix pretty printing of CUDA address spaces.

CUDA maps __shared__ internally also to __attribute__((shared)):

#define __annotate__(a) \
        __attribute__((a))
#define __location__(a) \
        __annotate__(a)
...
#define __shared__ \
        __location__(shared)

My guess is that Clang does it just the same way and only converts to LangAS::cuda_shared for code generation in GetGlobalVarAddressSpace:
https://clang.llvm.org/doxygen/CodeGenModule_8cpp_source.html#l03305
In contrast, OpenCL uses keywords that are mapped directly to LangAS::opencl_local etc.

Nov 13 2018, 2:25 AM · Restricted Project, Restricted Project

Nov 12 2018

richardmembarth added a comment to D54258: [Clang] Fix pretty printing of CUDA address spaces.

There are external tools (e.g. hipacc) that generate Clang AST. This AST uses LangAS annotations and emits incorrect memory space specifiers for CUDA when pretty-printed.

Nov 12 2018, 1:02 AM · Restricted Project, Restricted Project

Nov 9 2018

richardmembarth added a comment to D54258: [Clang] Fix pretty printing of CUDA address spaces.

Same problem here: The CUDA memory space specifiers are represented via attributes, e.g. CUDASharedAttr and only converted in CodeGen to LangAS::cuda_shared.
We would need a different frontend that annotates LangAS::cuda_shared.

Nov 9 2018, 6:32 AM · Restricted Project, Restricted Project

Nov 8 2018

richardmembarth added a comment to D54258: [Clang] Fix pretty printing of CUDA address spaces.

I think it's not so easy to provide such tests for CUDA.
CUDA memory space specifiers are implemented via attributes, e.g. #define __shared__ __attribute__((shared)).
As a result of this, they are pretty-printed via a different code path.
In my example, I call Ctx.getAddrSpaceQualType(QT, LangAS::cuda_shared), which is then pretty-printed via the code above.
Any hints how to provide tests for this one?

Nov 8 2018, 2:28 PM · Restricted Project, Restricted Project
richardmembarth updated the summary of D54258: [Clang] Fix pretty printing of CUDA address spaces.
Nov 8 2018, 6:45 AM · Restricted Project, Restricted Project
richardmembarth created D54258: [Clang] Fix pretty printing of CUDA address spaces.
Nov 8 2018, 6:36 AM · Restricted Project, Restricted Project