Page MenuHomePhabricator

yaxunl (Yaxun Liu)
User

Projects

User does not belong to any projects.

User Details

User Since
May 13 2015, 10:16 AM (209 w, 4 d)

Recent Activity

Wed, May 1

yaxunl added a comment to D61396: [hip] Fix ambiguity from `>>>` of CUDA..

LGTM too. Thanks Michael for fixing this.

Wed, May 1, 2:14 PM · Restricted Project, Restricted Project

Tue, Apr 30

yaxunl added inline comments to D61274: [Sema][AST] Explicit visibility for OpenCL/CUDA kernels/variables.
Tue, Apr 30, 3:27 PM · Restricted Project
yaxunl committed rG9e67d129f079: Add requires amdgpu-registered-target for amdgpu-float16.cpp (authored by yaxunl).
Add requires amdgpu-registered-target for amdgpu-float16.cpp
Tue, Apr 30, 12:06 PM
yaxunl committed rC359598: Add requires amdgpu-registered-target for amdgpu-float16.cpp.
Add requires amdgpu-registered-target for amdgpu-float16.cpp
Tue, Apr 30, 12:05 PM
yaxunl committed rL359598: Add requires amdgpu-registered-target for amdgpu-float16.cpp.
Add requires amdgpu-registered-target for amdgpu-float16.cpp
Tue, Apr 30, 12:05 PM
yaxunl committed rG44697012070c: AMDGPU: Enable _Float16 (authored by yaxunl).
AMDGPU: Enable _Float16
Tue, Apr 30, 11:35 AM
yaxunl committed rL359594: AMDGPU: Enable _Float16.
AMDGPU: Enable _Float16
Tue, Apr 30, 11:34 AM
yaxunl committed rC359594: AMDGPU: Enable _Float16.
AMDGPU: Enable _Float16
Tue, Apr 30, 11:34 AM

Fri, Apr 26

yaxunl accepted D61194: [HIP] Fix visibility of `__constant__` variables..

LGTM. Thanks!

Fri, Apr 26, 12:26 PM · Restricted Project
yaxunl added inline comments to D61194: [HIP] Fix visibility of `__constant__` variables..
Fri, Apr 26, 11:37 AM · Restricted Project

Wed, Apr 24

yaxunl created D61112: AMDGPU: Enable _Float16.
Wed, Apr 24, 8:08 PM

Apr 12 2019

yaxunl committed rG7bd8c37b1773: [HIP] Use -mlink-builtin-bitcode to link device library (authored by yaxunl).
[HIP] Use -mlink-builtin-bitcode to link device library
Apr 12 2019, 9:23 AM
yaxunl committed rC358290: [HIP] Use -mlink-builtin-bitcode to link device library.
[HIP] Use -mlink-builtin-bitcode to link device library
Apr 12 2019, 9:22 AM
yaxunl committed rL358290: [HIP] Use -mlink-builtin-bitcode to link device library.
[HIP] Use -mlink-builtin-bitcode to link device library
Apr 12 2019, 9:22 AM
yaxunl closed D60513: [HIP] Use -mlink-builtin-bitcode to link device library.
Apr 12 2019, 9:22 AM · Restricted Project
yaxunl updated the summary of D60620: [HIP] Support -offloading-target-id.
Apr 12 2019, 8:44 AM
yaxunl created D60620: [HIP] Support -offloading-target-id.
Apr 12 2019, 8:38 AM

Apr 10 2019

yaxunl created D60513: [HIP] Use -mlink-builtin-bitcode to link device library.
Apr 10 2019, 8:04 AM · Restricted Project

Apr 3 2019

yaxunl added inline comments to D59321: WIP: AMDGPU: Teach toolchain to link rocm device libs.
Apr 3 2019, 11:34 AM
yaxunl added reviewers for D59321: WIP: AMDGPU: Teach toolchain to link rocm device libs: b-sumner, t-tye.
Apr 3 2019, 11:23 AM

Apr 2 2019

yaxunl added inline comments to D60141: [HIP-Clang] Fat binary should not be produced for non GPU code.
Apr 2 2019, 1:32 PM · Restricted Project, Restricted Project
yaxunl accepted D60141: [HIP-Clang] Fat binary should not be produced for non GPU code.

LGTM. Thanks!

Apr 2 2019, 12:57 PM · Restricted Project, Restricted Project
yaxunl added a reviewer for D60141: [HIP-Clang] Fat binary should not be produced for non GPU code: tra.
Apr 2 2019, 11:04 AM · Restricted Project, Restricted Project

Mar 27 2019

yaxunl added inline comments to D59863: [HIP] Support gpu arch gfx906+sram-ecc.
Mar 27 2019, 11:05 AM

Mar 26 2019

yaxunl created D59863: [HIP] Support gpu arch gfx906+sram-ecc.
Mar 26 2019, 7:51 PM

Mar 21 2019

yaxunl created D59647: [CUDA][HIP] Warn shared var initialization.
Mar 21 2019, 8:39 AM · Restricted Project

Mar 15 2019

yaxunl added a comment to D59316: [HIP-Clang] propagate -mllvm options to opt and llc.
Mar 15 2019, 2:43 PM · Restricted Project
yaxunl added a comment to D59316: [HIP-Clang] propagate -mllvm options to opt and llc.

Here we are looking at the code which emulates a "linker" for HIP toolchain. The offloading action builder requests the offloading toolchain have a linker, but amdgpu does not have a real linker (ISA level linker), so we have to emulate that. If we have an ISA level linker we can get rid of all these stuff, but I don't think that will happen in short time.

This isn't really true. We do run lld to link the final executable. It also doesn't change that opt and llc should never be involved in the process

Can lld do ISA level linking? That is, one device function in one object file calls another device function in a different object file, and we let lld link them together?

We can't link multiple objects, but we do need to link the single object with lld. The relocations even for functions in the same module are 0 until lld fixes them up. Do we have execution tests for function calls using HIP? Since it looks like lld isn't getting used here, I suspect they aren't workingh

Mar 15 2019, 12:35 PM · Restricted Project
yaxunl added a comment to D59316: [HIP-Clang] propagate -mllvm options to opt and llc.

Here we are looking at the code which emulates a "linker" for HIP toolchain. The offloading action builder requests the offloading toolchain have a linker, but amdgpu does not have a real linker (ISA level linker), so we have to emulate that. If we have an ISA level linker we can get rid of all these stuff, but I don't think that will happen in short time.

This isn't really true. We do run lld to link the final executable. It also doesn't change that opt and llc should never be involved in the process

Mar 15 2019, 12:14 PM · Restricted Project

Mar 14 2019

yaxunl accepted D59316: [HIP-Clang] propagate -mllvm options to opt and llc.

LGTM. Thanks!

Mar 14 2019, 11:02 AM · Restricted Project
yaxunl added a comment to D59316: [HIP-Clang] propagate -mllvm options to opt and llc.

Here we are looking at the code which emulates a "linker" for HIP toolchain. The offloading action builder requests the offloading toolchain have a linker, but amdgpu does not have a real linker (ISA level linker), so we have to emulate that. If we have an ISA level linker we can get rid of all these stuff, but I don't think that will happen in short time.

Mar 14 2019, 11:02 AM · Restricted Project

Mar 5 2019

yaxunl committed rGc5be267003ef: [CUDA][HIP][Sema] Fix template kernel with function as template parameter (authored by yaxunl).
[CUDA][HIP][Sema] Fix template kernel with function as template parameter
Mar 5 2019, 10:20 AM
yaxunl committed rL355421: [CUDA][HIP][Sema] Fix template kernel with function as template parameter.
[CUDA][HIP][Sema] Fix template kernel with function as template parameter
Mar 5 2019, 10:20 AM
yaxunl committed rC355421: [CUDA][HIP][Sema] Fix template kernel with function as template parameter.
[CUDA][HIP][Sema] Fix template kernel with function as template parameter
Mar 5 2019, 10:20 AM
yaxunl closed D56411: [CUDA][HIP][Sema] Fix template kernel with function as template parameter.
Mar 5 2019, 10:19 AM · Restricted Project
yaxunl committed rG071949c3afae: Allow bundle size to be 0 in clang-offload-bundler (authored by yaxunl).
Allow bundle size to be 0 in clang-offload-bundler
Mar 5 2019, 9:52 AM
yaxunl committed rC355419: Allow bundle size to be 0 in clang-offload-bundler.
Allow bundle size to be 0 in clang-offload-bundler
Mar 5 2019, 9:52 AM
yaxunl committed rL355419: Allow bundle size to be 0 in clang-offload-bundler.
Allow bundle size to be 0 in clang-offload-bundler
Mar 5 2019, 9:52 AM
yaxunl closed D58057: Allow bundle size to be 0 in clang-offload-bundler.
Mar 5 2019, 9:52 AM · Restricted Project
yaxunl committed rGab851939fc6b: [HIP] Do not unbundle object files for -fno-gpu-rdc (authored by yaxunl).
[HIP] Do not unbundle object files for -fno-gpu-rdc
Mar 5 2019, 8:09 AM
yaxunl committed rC355410: [HIP] Do not unbundle object files for -fno-gpu-rdc.
[HIP] Do not unbundle object files for -fno-gpu-rdc
Mar 5 2019, 8:09 AM
yaxunl committed rL355410: [HIP] Do not unbundle object files for -fno-gpu-rdc.
[HIP] Do not unbundle object files for -fno-gpu-rdc
Mar 5 2019, 8:09 AM
yaxunl closed D58917: [HIP] Do not unbundle object files for -fno-gpu-rdc.
Mar 5 2019, 8:09 AM · Restricted Project, Restricted Project
yaxunl added a reviewer for D58057: Allow bundle size to be 0 in clang-offload-bundler: ABataev.
Mar 5 2019, 7:02 AM · Restricted Project
yaxunl added a comment to D58057: Allow bundle size to be 0 in clang-offload-bundler.

Alexey, could you please also review this patch? Thanks.

Mar 5 2019, 7:02 AM · Restricted Project
yaxunl added inline comments to D58917: [HIP] Do not unbundle object files for -fno-gpu-rdc.
Mar 5 2019, 7:02 AM · Restricted Project, Restricted Project

Mar 4 2019

yaxunl created D58917: [HIP] Do not unbundle object files for -fno-gpu-rdc.
Mar 4 2019, 12:01 PM · Restricted Project, Restricted Project

Feb 28 2019

yaxunl committed rGfd2c5c05fc8f: Partial revert of r353952: [HIP] Handle compile -m options and propagate into… (authored by yaxunl).
Partial revert of r353952: [HIP] Handle compile -m options and propagate into…
Feb 28 2019, 9:08 AM
yaxunl committed rC355106: Partial revert of r353952: [HIP] Handle compile -m options and propagate into….
Partial revert of r353952: [HIP] Handle compile -m options and propagate into…
Feb 28 2019, 9:08 AM
yaxunl committed rL355106: Partial revert of r353952: [HIP] Handle compile -m options and propagate into….
Partial revert of r353952: [HIP] Handle compile -m options and propagate into…
Feb 28 2019, 9:08 AM

Feb 27 2019

yaxunl committed rG785cbd850b76: [NFC] minor revision of r354929 [CUDA][HIP] Check calling convention based on… (authored by yaxunl).
[NFC] minor revision of r354929 [CUDA][HIP] Check calling convention based on…
Feb 27 2019, 7:48 AM
yaxunl committed rL354990: [NFC] minor revision of r354929 [CUDA][HIP] Check calling convention based on….
[NFC] minor revision of r354929 [CUDA][HIP] Check calling convention based on…
Feb 27 2019, 7:45 AM
yaxunl committed rC354990: [NFC] minor revision of r354929 [CUDA][HIP] Check calling convention based on….
[NFC] minor revision of r354929 [CUDA][HIP] Check calling convention based on…
Feb 27 2019, 7:45 AM
yaxunl added inline comments to D57716: [CUDA][HIP] Check calling convention based on function target.
Feb 27 2019, 7:11 AM · Restricted Project
yaxunl added inline comments to D57716: [CUDA][HIP] Check calling convention based on function target.
Feb 27 2019, 6:57 AM · Restricted Project

Feb 26 2019

yaxunl committed rGe739ac0e2555: [HIP] change kernel stub name (authored by yaxunl).
[HIP] change kernel stub name
Feb 26 2019, 6:03 PM
yaxunl committed rC354948: [HIP] change kernel stub name.
[HIP] change kernel stub name
Feb 26 2019, 6:03 PM
yaxunl committed rL354948: [HIP] change kernel stub name.
[HIP] change kernel stub name
Feb 26 2019, 6:03 PM
yaxunl closed D58518: [HIP] change kernel stub name.
Feb 26 2019, 6:03 PM · Restricted Project, Restricted Project
yaxunl committed rGfa49c3a888e8: [CUDA][HIP] Check calling convention based on function target (authored by yaxunl).
[CUDA][HIP] Check calling convention based on function target
Feb 26 2019, 2:27 PM
yaxunl committed rL354929: [CUDA][HIP] Check calling convention based on function target.
[CUDA][HIP] Check calling convention based on function target
Feb 26 2019, 2:27 PM
yaxunl committed rC354929: [CUDA][HIP] Check calling convention based on function target.
[CUDA][HIP] Check calling convention based on function target
Feb 26 2019, 2:27 PM
yaxunl closed D57716: [CUDA][HIP] Check calling convention based on function target.
Feb 26 2019, 2:27 PM · Restricted Project
yaxunl accepted D58623: [AMDGPU] Allow using integral non-type template parameters.

LGTM. Thanks!

Feb 26 2019, 9:38 AM · Restricted Project
yaxunl added a comment to D57716: [CUDA][HIP] Check calling convention based on function target.

ping

Feb 26 2019, 8:23 AM · Restricted Project
yaxunl added a reviewer for D58518: [HIP] change kernel stub name: rjmccall.
Feb 26 2019, 8:23 AM · Restricted Project, Restricted Project
yaxunl committed rGd83c74028db0: [OpenCL] Fix assertion due to blocks (authored by yaxunl).
[OpenCL] Fix assertion due to blocks
Feb 26 2019, 8:20 AM
yaxunl committed rL354893: [OpenCL] Fix assertion due to blocks.
[OpenCL] Fix assertion due to blocks
Feb 26 2019, 8:20 AM
yaxunl committed rC354893: [OpenCL] Fix assertion due to blocks.
[OpenCL] Fix assertion due to blocks
Feb 26 2019, 8:20 AM
yaxunl closed D58658: [OpenCL] Fix assertion due to blocks.
Feb 26 2019, 8:20 AM · Restricted Project
yaxunl added a comment to D56411: [CUDA][HIP][Sema] Fix template kernel with function as template parameter.

I would like to fix the validation issue only and leave the overload resolution issue for future.

As I understand it, the "validation issue" is just that you'd like a diagnostic to be emitted when resolving the template argument in order to force SFINAE to pick a different template. I think that's actually just the overload-resolution issue.

Feb 26 2019, 7:16 AM · Restricted Project

Feb 25 2019

yaxunl created D58658: [OpenCL] Fix assertion due to blocks.
Feb 25 2019, 5:39 PM · Restricted Project

Feb 22 2019

yaxunl added inline comments to D58518: [HIP] change kernel stub name.
Feb 22 2019, 3:47 PM · Restricted Project, Restricted Project
yaxunl reopened D58518: [HIP] change kernel stub name.
Feb 22 2019, 1:45 PM · Restricted Project, Restricted Project
yaxunl updated the diff for D58518: [HIP] change kernel stub name.

Fixed regressions.

Feb 22 2019, 1:45 PM · Restricted Project, Restricted Project
yaxunl updated the diff for D57716: [CUDA][HIP] Check calling convention based on function target.

modify test to use non-template functions.

Feb 22 2019, 8:44 AM · Restricted Project
yaxunl added inline comments to D57716: [CUDA][HIP] Check calling convention based on function target.
Feb 22 2019, 8:40 AM · Restricted Project

Feb 21 2019

yaxunl committed rG00ebc0cb92e9: revert r354615: [HIP] change kernel stub name (authored by yaxunl).
revert r354615: [HIP] change kernel stub name
Feb 21 2019, 8:20 PM
yaxunl committed rL354651: revert r354615: [HIP] change kernel stub name.
revert r354615: [HIP] change kernel stub name
Feb 21 2019, 8:19 PM
yaxunl committed rC354651: revert r354615: [HIP] change kernel stub name.
revert r354615: [HIP] change kernel stub name
Feb 21 2019, 8:19 PM
yaxunl updated the diff for D57716: [CUDA][HIP] Check calling convention based on function target.

Revised by Artem's comments.

Feb 21 2019, 2:04 PM · Restricted Project
yaxunl committed rG8d7cf0e2d4b5: [HIP] change kernel stub name (authored by yaxunl).
[HIP] change kernel stub name
Feb 21 2019, 12:13 PM
yaxunl committed rL354615: [HIP] change kernel stub name.
[HIP] change kernel stub name
Feb 21 2019, 12:12 PM
yaxunl committed rC354615: [HIP] change kernel stub name.
[HIP] change kernel stub name
Feb 21 2019, 12:11 PM
yaxunl closed D58518: [HIP] change kernel stub name.
Feb 21 2019, 12:11 PM · Restricted Project, Restricted Project
yaxunl updated the diff for D56411: [CUDA][HIP][Sema] Fix template kernel with function as template parameter.

I would like to fix the validation issue only and leave the overload resolution issue for future.

Feb 21 2019, 11:43 AM · Restricted Project
yaxunl accepted D58509: [CodeGen] Fix string literal address space casting..

LGTM. Thanks!

Feb 21 2019, 11:28 AM · Restricted Project
yaxunl added a comment to D58518: [HIP] change kernel stub name.
In D58518#1406124, @tra wrote:

My guess is that this is needed because HIP debugger can see symbols from both host and device executables at the same time. Is that so?

If that's the case, I guess HIP may have similar naming problem for __host__ __device__ foo() if it's used on both host and device.

Feb 21 2019, 11:24 AM · Restricted Project, Restricted Project
yaxunl created D58518: [HIP] change kernel stub name.
Feb 21 2019, 10:05 AM · Restricted Project, Restricted Project

Feb 15 2019

yaxunl added a comment to D56411: [CUDA][HIP][Sema] Fix template kernel with function as template parameter.

It is totally unreasonable, at the time you are resolving a template argument, to investigate how the corresponding template parameter is used within the template and use that to shape how the template argument is resolved. That is simply not how the C++ template model works. Given that CODA doesn't distinguish between host and device functions in the type system, if you are going to have a rule here, it has to be based on, at most, (1) the current semantic context (which may not even be a function), (2) the template being specialized, and (3) the declarations in the template-argument set.

As I've said before on a previous patch, I think the *best* rule would be to recognize a hard difference between host and device function types, probably by making function types default to being host function types and requiring function pointers that can store device function pointers to be explicitly annotated. However, that would not be source-compatible with ordinary CUDA, which is presumably unacceptable.

The second-best rule would be to preserve compatibility by making an unannotated function type still be "unknown whether host or device", but to also allow the creation of explicitly host-only and device-only function types. For source compatibility, DREs to functions would formally have the unknown function type. Converting a pointer to an unknown function into a pointer to a host function would do some basic checking on the operand expression (basically to verify that it's not obviously a device function), and resolving an overload set in the context of a host-only function pointer type would do the obvious filtering.

Otherwise, you're going to be stuck where you are right now, which is that you're messing around with heuristics because somebody added a language extension that isn't actually very well thought out. But if that's what you have to do, it's what you have to do. For this specific question, where you are trying to resolve an overloaded template argument, I think there are basically two sensible options.

  • You can filter the overloads by the host-ness of the template. This makes some sense, because it's probably most likely that a function template that takes a function as a template argument is going to call it — but not necessarily, because it very well might decide instead to call over to the device to invoke the function. Also, not all templates have a "host-ness"; that's pretty much exclusive to function templates.
  • You can filter the overload by the host-ness of the current context. Again, this makes some sense because it's likely that a host function is trying to pass down a host function — but again, it's not hard to think of exceptions. And again, this has the problem that the context isn't always a function and so doesn't necessarily have a host-ness. Any sort of additional template-specific guidance seems doomed to gradually turn into the second design I mentioned above where you have the ability to be more specific about function types.

    For the time being, this is still a Clang extension, and while Artem mentioned that NVIDIA is investigating it, that's presumably still an investigation and we still have an opportunity to shape their thinking. So I would really recommend taking the second approach, or maybe even trying to convince them to take the first. (How common is higher-order programming on the device, anyway, that you can't break source compatibility for it?) For this specific line of inquiry, that would probably mean not trying to automatically use any particular filter on the overload set but instead just relying on the programmer to annotation what kind of function they want.
Feb 15 2019, 9:28 PM · Restricted Project
yaxunl added a comment to D56411: [CUDA][HIP][Sema] Fix template kernel with function as template parameter.

But what we've just been talking about is not a validity rule, it's an overload-resolution rule. It's not *invalid* to use a device function as a template argument to a host function template (or to a class template, which of course is neither host nor device). All you need to do is to resolve otherwise-intractable overload ambiguities by matching with the host-ness of the current context, which there's probably already code to do for when an overload set is used as e.g. a function argument.

Feb 15 2019, 4:39 PM · Restricted Project

Feb 14 2019

yaxunl added a comment to D56411: [CUDA][HIP][Sema] Fix template kernel with function as template parameter.
In D56411#1398291, @tra wrote:

That said, does CUDA have a general rule resolving __host__ vs. __device__ overloads based on context? And does it allow overloading based solely on __host__ vs. __device__?

NVCC does not. Clang does. See https://goo.gl/EXnymm for the details.

AFAICT, NVIDIA is starting to consider adopting Clang's approach:
http://lists.llvm.org/pipermail/cfe-dev/2018-November/060070.html (original message from Bryce apparently didn't make it to the cfe-dev archive)

Okay. Probably the template-argument rule ought to be the same as the address-of-function rule, which I assume means that there's a final pass that resolves ambiguities in favor of functions that can be used from the current context, to the extent that that's meaningful. It's hard to tell because that document does not appear to include a formal specification.

Regardless, that has no effect on this patch.

Feb 14 2019, 1:00 PM · Restricted Project
yaxunl added a comment to D56411: [CUDA][HIP][Sema] Fix template kernel with function as template parameter.
In D56411#1398291, @tra wrote:

That said, does CUDA have a general rule resolving __host__ vs. __device__ overloads based on context? And does it allow overloading based solely on __host__ vs. __device__?

NVCC does not. Clang does. See https://goo.gl/EXnymm for the details.

AFAICT, NVIDIA is starting to consider adopting Clang's approach:
http://lists.llvm.org/pipermail/cfe-dev/2018-November/060070.html (original message from Bryce apparently didn't make it to the cfe-dev archive)

Feb 14 2019, 11:28 AM · Restricted Project
yaxunl added a comment to D56411: [CUDA][HIP][Sema] Fix template kernel with function as template parameter.

I think the diagnostic should come during instantiation when you find an evaluated use of a host function within a device function.

It seems the body of function template is checked only during parsing of the definition of the template itself. When a function
template is instantiated, the body of the instantiated function is not checked again.

No, that's not correct. However, it's checked somewhat differently, and it's possible that the existing diagnostic is not set up to fire along all common paths. Try moving the diagnostic to MarkFunctionReferenced, and note that OdrUse will be false in all the unevaluated contexts.

You are right. After I disable current diagnostic, I saw PerformPendingInstantiations at the end of parsing the TU, where the AST of the instantiated function is iterated and MarkFunctionReferenced is called. I will try to fix my patch as suggested. Thanks.

I got one concern. If we want to do overload resolution of function type template argument based on host or device, we need to do that before template instantiation, right?

e.g. we have two functions having the same name f and type, but one is __host__ and the other is __device__, and we pass it as a template argument to a template function g. We want to choose __device__ f if g itself is __device__ and __host__ f if g itself is __host__. If we want to do this we have to do the check before template instantiation, right?

Yes, you would need to check that when resolving the overload to a single declaration. That would be separate from diagnosing uses.

That said, does CUDA have a general rule resolving __host__ vs. __device__ overloads based on context? And does it allow overloading based solely on __host__ vs. __device__?

Feb 14 2019, 9:03 AM · Restricted Project
yaxunl added a comment to D56411: [CUDA][HIP][Sema] Fix template kernel with function as template parameter.

I think the diagnostic should come during instantiation when you find an evaluated use of a host function within a device function.

It seems the body of function template is checked only during parsing of the definition of the template itself. When a function
template is instantiated, the body of the instantiated function is not checked again.

No, that's not correct. However, it's checked somewhat differently, and it's possible that the existing diagnostic is not set up to fire along all common paths. Try moving the diagnostic to MarkFunctionReferenced, and note that OdrUse will be false in all the unevaluated contexts.

You are right. After I disable current diagnostic, I saw PerformPendingInstantiations at the end of parsing the TU, where the AST of the instantiated function is iterated and MarkFunctionReferenced is called. I will try to fix my patch as suggested. Thanks.

Feb 14 2019, 8:35 AM · Restricted Project

Feb 13 2019

yaxunl committed rGc18e9ecd4fc1: [CUDA][HIP] Use device side kernel and variable names when registering them (authored by yaxunl).
[CUDA][HIP] Use device side kernel and variable names when registering them
Feb 13 2019, 6:00 PM
yaxunl committed rL354004: [CUDA][HIP] Use device side kernel and variable names when registering them.
[CUDA][HIP] Use device side kernel and variable names when registering them
Feb 13 2019, 6:00 PM
yaxunl committed rC354004: [CUDA][HIP] Use device side kernel and variable names when registering them.
[CUDA][HIP] Use device side kernel and variable names when registering them
Feb 13 2019, 6:00 PM
yaxunl closed D58163: [CUDA][HIP] Use device side kernel and variable names when registering them.
Feb 13 2019, 5:59 PM · Restricted Project
yaxunl updated the diff for D58163: [CUDA][HIP] Use device side kernel and variable names when registering them.

Revised by Artem's comments.

Feb 13 2019, 11:42 AM · Restricted Project