[SYCL][COMPAT][cuda] Add "ptr_to_integer" syclcompat functions.

[JackAKirk]

Add "ptr_to_integer" (generic address space to .shared) syclcompat functions.

These functions are commonly required in optimized libraries that use inline ptx. The standard naming convention of removing "__" from corresponding cuda builtins has been applied. See the readme and accompanying test-e2e for example usage.

[joeatodd]

Since these functions do the same thing aside from casting to int/size_t, can we not implement them as a single templated function?

[JackAKirk]

Since these functions do the same thing aside from casting to int/size_t, can we not implement them as a single templated function?

Uncertainty around this is the reason I put them in experimental. It's a bit messy since the cuda versions of these api require different cuda toolkit versions (10.1 for the uint32_t and 11 for size_t, I think), but this does not affect these syclcompat translated versions. I was just told to translate them in this way so that cutlass sycl path can have corresponding apis to cuda runtime path. I don't think I really have the context to make a decision beyond this. It is probably best to ask @aacostadiaz what is best.

[JackAKirk]

Since these functions do the same thing aside from casting to int/size_t, can we not implement them as a single templated function?

@aacostadiaz wants them to be two separate functions, so I'll leave it as it is.

[JackAKirk]

@Alcpz @joeatodd Any more reviews for this?

Thanks

[joeatodd]

As discussed offline, these should be a single function with a template parameter describing the return type.

[npmiller]

Closing this after further discussions offline

[joeatodd]

Thanks for this @JackAKirk. Just a couple of formatting requests. Cheers!

[joeatodd]

Could you fix the formatting of this code section? Thanks

[JackAKirk]

I did clang-format it already using dpc++ format.

[joeatodd]

Possibly it's not running on code sections in markdown? I'd expect uint32_t fragment to align with T addr on the line above? And the line split on lines 975-976 looks pretty wacky? If I dump this code into a cpp file and autoformat this, I get:

  half *data = syclcompat::local_mem<half[NUM_ELEMENTS]>();
  // ...
  // ...
  T addr =
      syclcompat::ptr_to_int<T>(reinterpret_cast<char *>(data) + (id % 8) * 16);
  uint32_t fragment;
#if defined(__NVPTX__)
  asm volatile("ldmatrix.sync.aligned.m8n8.x1.shared.b16 {%0}, [%1];\n"
               : "=r"(fragment)
               : "r"(addr));
#endif

[JackAKirk]

I can see if that passes clang-format (in the test where it is used). The existing version passes the clang-format on the clang-format CI.

[joeatodd]

I don't think clang-format runs on the README tbh.

[JackAKirk]

Yeah, I use the same code in the test-e2e

[JackAKirk]

Possibly it's not running on code sections in markdown? I'd expect uint32_t fragment to align with T addr on the line above? And the line split on lines 975-976 looks pretty wacky? If I dump this code into a cpp file and autoformat this, I get:

  half *data = syclcompat::local_mem<half[NUM_ELEMENTS]>();
  // ...
  // ...
  T addr =
      syclcompat::ptr_to_int<T>(reinterpret_cast<char *>(data) + (id % 8) * 16);
  uint32_t fragment;
#if defined(__NVPTX__)
  asm volatile("ldmatrix.sync.aligned.m8n8.x1.shared.b16 {%0}, [%1];\n"
               : "=r"(fragment)
               : "r"(addr));
#endif

I've updated the README with this suggestion now

[JackAKirk]

@Alcpz is this OK now?
Thanks

[Alcpz]

@Alcpz is this OK now? Thanks

Yes. I agree with @joeatodd review.
Accepting your changes, assuming that you will finalize addressing his suggestions. Sorry for missing this.

[JackAKirk]

@Alcpz is this OK now? Thanks

Yes. I agree with @joeatodd review. Accepting your changes, assuming that you will finalize addressing his suggestions. Sorry for missing this.

Yes, I've updated the formatting now, thanks.

[joeatodd]

LGTM

[JackAKirk]

@intel/llvm-gatekeepers Please merge this.

Thanks