[SYCL] Support missing types with ldg extension #8748

jchlanda · 2023-03-23T10:32:07Z

No description provided.

jchlanda · 2023-03-23T10:34:14Z

Marking as a draft till the pull-down is done (#8738), as this patch relies on the half ldg builtins that were added in to the mainline here: llvm/llvm-project@71b0658

jchlanda · 2023-03-23T10:34:34Z

cc: @JackAKirk

JackAKirk · 2023-03-23T11:07:38Z

LGTM

btw you can add half cases to the e2e test: intel/llvm-test-suite#1417 (although initial test cases aren't merged yet either)

gmlueck · 2023-03-27T22:18:12Z

sycl/doc/extensions/experimental/sycl_ext_oneapi_cuda_tex_cache_read.asciidoc

@@ -98,7 +98,7 @@ T ldg(const T* ptr);
 `ldg` returns the data of type `T` located at address `ptr`. When called from the `ext_oneapi_cuda` backend the data is cached in the read-only texture cache.
 When called from any other backend a copy of the data stored at address `ptr` is returned without using any special cache.

-The template parameter `T` can be one of `char`, `signed char`, `short`, `int`, `long`, `long long`, `unsigned char`, `unsigned short`, `unsigned int`, `unsigned long`, `unsigned long long`, `vec<char, 2>`, `vec<char, 4>`, `vec<short, 2>`, `vec<short, 4>`, `vec<int, 2>`, `vec<int, 4>`, `vec<long long, 2>`, `vec<uchar, 2>`, `vec<uchar, 4>`, `vec<ushort, 2>`, `vec<ushort, 4>`, `vec<uint, 2>`, `vec<uint, 4>`, `vec<unsigned long long, 2>`, `float`, `vec<float, 2>`, `vec<float, 4>`, `double`, or `vec<double, 2>`.
+The template parameter `T` can be one of `char`, `signed char`, `short`, `int`, `long`, `long long`, `unsigned char`, `unsigned short`, `unsigned int`, `unsigned long`, `unsigned long long`, `half`, `float`, `double`, `vec<char, 2>`, `vec<char, 4>`, `vec<short, 2>`, `vec<short, 4>`, `vec<int, 2>`, `vec<int, 4>`, `vec<long long, 2>`, `vec<uchar, 2>`, `vec<uchar, 4>`, `vec<ushort, 2>`, `vec<ushort, 4>`, `vec<uint, 2>`, `vec<uint, 4>`, `vec<unsigned long long, 2>`, `vec<half, 2>`, `vec<float, 2>`, `vec<float, 4>`, or `vec<double, 2>`.


Did you intentionally omit vec<half, 4>? If so, why?

~~vec<half, 4> isn't natively supported with a single __ldg instruction.~~ But it could be supported by just calling the vec2 instruction twice. There is at least one case I saw where an application implements this themselves for double4, so this could be something useful to add in the future.

Actually maybe half4/double4 are natively implementable, at least this is what is suggested by https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-ld-global-nc
We could investigate this in the future.
The types that we added so far match those exposed through CUDA: https://docs.nvidia.com/cuda/cuda-c-programming-guide/#read-only-data-cache-load-function

If this is what the CUDA docs say are natively implemented, that is fine. I only pointed it out because it doesn't seem to match your existing pattern. For the other types, it seems like there are 2- and 4-element versions for all the types that are 32-bits or less. (E.g. float2, float4, int2, int4, short2, short4, etc.)

Yes this matches the pattern here: https://docs.nvidia.com/cuda/cuda-c-programming-guide/#read-only-data-cache-load-function

However I don't see why we can't also implement it for additional types like double4 in the future, especially if the ptx allows this to be done via a single instruction. It is just that the equivalent of double4 etc is not exposed via CUDA Runtime API.

That is a good point I followed the types we exposed as builtins, but you are right, I should have been more consistent with types already supported. I'll add the vec<half, 4> as calls to the existing builtins and will revisit it if/when we implement support for instructions mentioned by Jack.

If we are adding vec<half, 4> can you also add vec<double, 4>? MD uses the 4D version for double (basically they need 3 dimensions for positions/velocities): glotzerlab/hoomd-blue@730a809#diff-c20b0af2267267f69ae103ae2c7e4aecea76b0e5feb3d3f22f0e48f4b652c08dR28

This will save them having to implement this themselves.

Thanks

Sure, if we are sticking to vectors of 2 and 4 elements, I'll add 4 elements variants of longlong and ulonglong.

gmlueck · 2023-03-29T12:32:40Z

sycl/doc/extensions/experimental/sycl_ext_oneapi_cuda_tex_cache_read.asciidoc

@@ -98,7 +98,7 @@ T ldg(const T* ptr);
 `ldg` returns the data of type `T` located at address `ptr`. When called from the `ext_oneapi_cuda` backend the data is cached in the read-only texture cache.
 When called from any other backend a copy of the data stored at address `ptr` is returned without using any special cache.

-The template parameter `T` can be one of `char`, `signed char`, `short`, `int`, `long`, `long long`, `unsigned char`, `unsigned short`, `unsigned int`, `unsigned long`, `unsigned long long`, `vec<char, 2>`, `vec<char, 4>`, `vec<short, 2>`, `vec<short, 4>`, `vec<int, 2>`, `vec<int, 4>`, `vec<long long, 2>`, `vec<uchar, 2>`, `vec<uchar, 4>`, `vec<ushort, 2>`, `vec<ushort, 4>`, `vec<uint, 2>`, `vec<uint, 4>`, `vec<unsigned long long, 2>`, `float`, `vec<float, 2>`, `vec<float, 4>`, `double`, or `vec<double, 2>`.
+The template parameter `T` can be one of `char`, `signed char`, `short`, `int`, `long`, `long long`, `unsigned char`, `unsigned short`, `unsigned int`, `unsigned long`, `unsigned long long`, `half`, `float`, `double`, `vec<char, 2>`, `vec<char, 4>`, `vec<short, 2>`, `vec<short, 4>`, `vec<int, 2>`, `vec<int, 4>`, `vec<long long, 2>`, `vec<long long, 4>`, `vec<uchar, 2>`, `vec<uchar, 4>`, `vec<ushort, 2>`, `vec<ushort, 4>`, `vec<uint, 2>`, `vec<uint, 4>`, `vec<unsigned long long, 2>`, `vec<unsigned long long, 4>`, `vec<half, 2>`, `vec<half, 4>`, `vec<float, 2>`, `vec<float, 4>`, `vec<double, 2>`, or `vec<double, 4>`.


At this point you are missing these combinations:

vec<signed char, 2>

vec<signed char, 4>

vec<long, 2>

vec<long, 4>

vec<unsigned long, 2>

vec<unsigned long, 4>

Does it make sense to add them to make a complete set?

I also noticed that you use aliases uchar, ushort, and uint which are being deprecated in the core SYCL spec because they don't mean what you would expect. I think you probably want to change these to unsigned char, unsigned short, and unsigned int.

Yes to both.

I've added missing types and tidied up the syntax on the vector usage. This PR will need to stay in a draft mode for a wee bit longer, as we didn't have all the relevant builtins added in upstream llvm.

gmlueck

spec changes LGTM.

zjin-lcf · 2023-04-14T15:41:09Z

Are uchar3 and char3 equivalent to uchar4 and char4 ?

zjin-lcf · 2023-04-14T16:27:57Z

#ifdef ENABLE_BF16
template<>
inline __device__ __nv_bfloat162 ldg(const __nv_bfloat162* val) {
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 800
    return val[0];
#else
    return __ldg(val);
#endif
}

template<>
inline __device__ __nv_bfloat16 ldg(const __nv_bfloat16* val) {
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 800
    return val[0];
#else
    return __ldg(val);
#endif
}

zjin-lcf · 2023-04-20T15:22:46Z

Could the reviewers (cperkinsintel and intel llvm-reviewers) comment on the PR ? Quite a few programs need the extension for half-precision and char3 data type. Thanks.

jchlanda · 2023-06-30T10:12:23Z

@gmlueck are you still happy with the patch, with this PR in ldg should cover the following types:

char,
signed char,
short,
int,
long,
long long,
unsigned char,
unsigned short,
unsigned int,
unsigned long,
unsigned long long,
half,
float,
double

alongside their vectors of 2, 3 and 4 elements.

gmlueck

Yes, spec still looks good!

jchlanda · 2023-07-04T09:15:54Z

@intel/llvm-reviewers-runtime gentle ping, is there anything else that this PR needs?

steffenlarsen

I am a little sad that we do so many seemingly repeated yet different branches, but I don't have a good alternative.

LGTM!

jchlanda · 2023-07-06T06:12:59Z

I am a little sad that we do so many seemingly repeated yet different branches, but I don't have a good alternative.

LGTM!

Thanks @steffenlarsen, had a go at macros, but it got counterproductive really quick.

jinz2014 · 2023-07-10T14:49:49Z

@jchlanda
I have some question. Some CUDA kernels contain an argument of user-defined struct/class types. Could we still apply "ldg" ? Thanks.

jchlanda · 2023-07-11T05:26:34Z

@jchlanda I have some question. Some CUDA kernels contain an argument of user-defined struct/class types. Could we still apply "ldg" ? Thanks.

Hi @jinz2014
This will not work with the current implementation, sorry. And I am not sure if it could ever be implemented for arbitrary user types. While ISA specifies the load.global.nc instruction can operate on the untyped b64 you would still be limited by the max number of vector elements for the load, which is set to 4. See: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-ld-global-nc for details.

jchlanda · 2023-07-24T08:54:20Z

@intel/llvm-gatekeepers is this PR OK to merge in?

jchlanda requested review from a team as code owners March 23, 2023 10:32

jchlanda requested a review from cperkinsintel March 23, 2023 10:32

jchlanda marked this pull request as draft March 23, 2023 10:32

jchlanda temporarily deployed to aws March 23, 2023 10:51 — with GitHub Actions Inactive

jchlanda force-pushed the jakub/ldg_half branch from 9148b06 to 93ea36f Compare March 23, 2023 11:00

jchlanda temporarily deployed to aws March 23, 2023 11:19 — with GitHub Actions Inactive

gmlueck reviewed Mar 27, 2023

View reviewed changes

jchlanda temporarily deployed to aws March 29, 2023 11:28 — with GitHub Actions Inactive

gmlueck reviewed Mar 29, 2023

View reviewed changes

jchlanda temporarily deployed to aws March 29, 2023 12:35 — with GitHub Actions Inactive

jchlanda temporarily deployed to aws March 29, 2023 13:21 — with GitHub Actions Inactive

jchlanda force-pushed the jakub/ldg_half branch from abc174b to f0ef805 Compare March 31, 2023 06:24

jchlanda temporarily deployed to aws March 31, 2023 06:52 — with GitHub Actions Inactive

jchlanda temporarily deployed to aws March 31, 2023 07:36 — with GitHub Actions Inactive

jchlanda temporarily deployed to aws March 31, 2023 11:28 — with GitHub Actions Inactive

gmlueck approved these changes Mar 31, 2023

View reviewed changes

jchlanda changed the title ~~[SYCL] Support half and vec2 half types with ldg extension~~ [SYCL] Support missing types with ldg extension Mar 31, 2023

cperkinsintel requested a review from steffenlarsen April 20, 2023 16:57

jinz2014 mentioned this pull request Apr 20, 2023

[SYCL] half-precision tanh function #9100

Closed

jchlanda force-pushed the jakub/ldg_half branch from 880e699 to 16fa357 Compare May 29, 2023 13:30

jchlanda temporarily deployed to aws May 29, 2023 13:49 — with GitHub Actions Inactive

JackAKirk mentioned this pull request May 31, 2023

[SYCL] Added tests for ldg builtins. #9374

Closed

jchlanda force-pushed the jakub/ldg_half branch from 16fa357 to 756ff5f Compare June 2, 2023 08:22

jchlanda temporarily deployed to aws June 15, 2023 17:35 — with GitHub Actions Inactive

jchlanda force-pushed the jakub/ldg_half branch from 25008bc to 176b817 Compare June 16, 2023 08:46

jchlanda temporarily deployed to aws June 16, 2023 09:01 — with GitHub Actions Inactive

jchlanda temporarily deployed to aws June 16, 2023 09:39 — with GitHub Actions Inactive

jchlanda force-pushed the jakub/ldg_half branch from 176b817 to eae74fd Compare June 27, 2023 11:27

jchlanda temporarily deployed to aws June 27, 2023 11:43 — with GitHub Actions Inactive

jchlanda temporarily deployed to aws June 27, 2023 12:35 — with GitHub Actions Inactive

jchlanda force-pushed the jakub/ldg_half branch from eae74fd to 0a0811f Compare June 30, 2023 10:00

jchlanda requested a review from gmlueck June 30, 2023 10:12

jchlanda temporarily deployed to aws June 30, 2023 11:09 — with GitHub Actions Inactive

jchlanda temporarily deployed to aws June 30, 2023 13:24 — with GitHub Actions Inactive

gmlueck approved these changes Jun 30, 2023

View reviewed changes

steffenlarsen approved these changes Jul 5, 2023

View reviewed changes

jchlanda added 6 commits July 24, 2023 08:55

[SYCL] Support missing types with ldg extension

83b646d

Add <4 x half> support

e841037

Typo fix

207129a

d4, ll4 and ull4

f55ea82

sign char, vec<sign char, 2>, vec<unsigned|long, 2|4>

897981d

Vec 3

d8c4f82

jchlanda force-pushed the jakub/ldg_half branch from 0a0811f to d8c4f82 Compare July 24, 2023 07:58

jchlanda temporarily deployed to aws July 24, 2023 08:13 — with GitHub Actions Inactive

jchlanda temporarily deployed to aws July 24, 2023 08:51 — with GitHub Actions Inactive

steffenlarsen merged commit aec8a35 into intel:sycl Jul 24, 2023

mdtoguchi pushed a commit to mdtoguchi/llvm that referenced this pull request Oct 18, 2023

[SYCL] Support missing types with ldg extension (intel#8748)

3c6ea0d

[SYCL] Support missing types with ldg extension #8748

[SYCL] Support missing types with ldg extension #8748

Uh oh!

Conversation

jchlanda commented Mar 23, 2023

Uh oh!

jchlanda commented Mar 23, 2023

Uh oh!

jchlanda commented Mar 23, 2023

Uh oh!

JackAKirk commented Mar 23, 2023

Uh oh!

Choose a reason for hiding this comment

Uh oh!

JackAKirk Mar 28, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

JackAKirk Mar 29, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

gmlueck left a comment

Choose a reason for hiding this comment

Uh oh!

zjin-lcf commented Apr 14, 2023

Uh oh!

zjin-lcf commented Apr 14, 2023

Uh oh!

zjin-lcf commented Apr 20, 2023

Uh oh!

jchlanda commented Jun 30, 2023

Uh oh!

gmlueck left a comment

Choose a reason for hiding this comment

Uh oh!

jchlanda commented Jul 4, 2023

Uh oh!

steffenlarsen left a comment • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

jchlanda commented Jul 6, 2023

Uh oh!

jinz2014 commented Jul 10, 2023

Uh oh!

jchlanda commented Jul 11, 2023

Uh oh!

jchlanda commented Jul 24, 2023

Uh oh!

Uh oh!

JackAKirk Mar 28, 2023 •

edited

Loading

JackAKirk Mar 29, 2023 •

edited

Loading

steffenlarsen left a comment •

edited

Loading