[SYCL][CUDA] read-only vectors (marked with 'const' and '__restrict__') translated to the instruction 'ld.global.nc'

[zjin-lcf]

Please see the descriptions of the issue below:

AdaptiveCpp/AdaptiveCpp#848

[AerialMantis]

@zjin-lcf thanks for raising this, we will look into it.

[zjin-lcf]

Thanks.

[JackAKirk]

The issue is that the compiler will only use the ld.global.nc instruction if it is sure that the return value is not being written to throughout the kernel execution. In some cases it appears that the compilers referenced by the linked issue (nvcc and clang) cannot work this out and it reverts to the standard global load. It is for this reason that it is recommended to explicitly call e.g. the __ldg clang builtins (https://github.com/intel/llvm/blob/6485ec461da8a7abfbd72b930204bb782a1ce3ca/clang/lib/Headers/__clang_cuda_intrinsics.h) when compiling cuda code, which guarantees that ld.global.nc instruction will be used provided that the cached variable is not written to throughout the kernel. This is the standard way of accessing the texture cache for gpgpu use cases: see the examples that I reference in #7397.
I've introduced an extension that exposes these builtins to sycl via a new function cache_read. See https://github.com/intel/llvm/blob/d99eda476ef2ac1b66b186014c9d2f25b9989696/sycl/test/check_device_code/cuda/cache_read.cpp for a simple usage example. Note that when using cache_read it is no longer necessary to also wrap the pointer in const __restrict__ which further simplifies the code referenced in this issue.

[zjin-lcf]

Thank you for the extension. Could you please explain why cache_read needs to be introduced when const __restrict__ are familiar to many GPU developers ? In other words, does const __restrict__ provide sufficient hint/information for a compiler to generate the "ld.global.nc" ? instructions ? Thanks.

[JackAKirk]

Thank you for the extension. Could you please explain why cache_read needs to be introduced when const __restrict__ are familiar to many GPU developers ? In other words, does const __restrict__ provide sufficient hint/information for a compiler to generate the "ld.global.nc" ? instructions ? Thanks.

cache_reads has to be introduced unless we can guarantee that the const __restrict__ hint will work in all possible use cases (that work when the __ldg functions is explicitly called) without also calling an equivalent to the __ldg function (cache_read). My understanding is that this is very hard to guarantee and in practice clearly not the case.

The CUDA programming guide also appears to always recommend using __ldg(aka cache_read) without just relying on const __restrict:

• "Data that is read-only for the entire lifetime of the kernel can also be cached in the read-only data cache described in the previous section by reading it using the __ldg() function (see Read-Only Data Cache Load Function). When the compiler detects that the read-only condition is satisfied for some data, it will use __ldg() to read it. The compiler might not always be able to detect that the read-only condition is satisfied for some data. Marking pointers used for loading such data with both the const and restrict qualifiers increases the likelihood that the compiler will detect the read-only condition."

As stated in the extension document and in the above excerpt, if the compiler believes that the read-only condition is broken the cache is not used; adding the qualifiers can also make it more likely that the compiler does not incorrectly believe that the variable is written to. Note that it is stated that using both the const and restrict qualifiers only increases the likelihood rather than guaranteeing that the read-only condition is detected.

We also have an internal issue to try to improve the const __restrict__ hints ability to generate ld.global.nc without calling the __ldg, since it sounds like sycl is performing less well than cuda in this regard from your linked issue. However based on the number of issues complaining about failure of the const __restrict__ hints to achieve this without calling __ldg on online forums, we do not want to rely solely on this.

I think that historically it has always been recommended to use an explicit API to access the cuda texture cache. Originally the cuda texture API dedicated to to image processing was the only API exposing the texture cache, and this API was widely used also in more general GPGPU programming. Eventually the __ldg instruction was provided in the CUDA API and later extended so that the texture ld.global.nc instruction could be called generally without using CUDA textures.
So also for people that know the hardware and have used texture memory previously it is common that they will expect to find an explicit API that they can call within DPC++.

[zjin-lcf]

Does the following example explain "Note that when using cache_read it is no longer necessary to also wrap the pointer in const restrict which further simplifies the code referenced in this issue." ?

void add(nd_item<1> &item, int *a, int *b, int *c) {
  int i = item.get_global_id(0);
  c[i] = cache_read(a+i) + cache_read(b+i);
}

q.submit([=](sycl::handler &h) {
   h.parallel_for(...) {
     add(item, a, b, c);
   });
});

[JackAKirk]

I've just tested this code you've put here and it also uses the tex cache without const __restrict__. I don't know yet if there will be cases that won't work with cache_read without the const __restrict__, and if so, whether there will also be cases where it doesn't work with both cache_read and const __restrict__ used (and cached variable not written to). I'll update the extension accordingly as it becomes clearer with more testing.
BTW it looks like there isn't interest from other backends so I think we have decided to rename cache_read as __ldg for the initial extension, so that it is recognizable by CUDA programmers.

[zjin-lcf]

Renaming is helpful for migrating CUDA/HIP codes. Thanks.

"
Textures and Cache Control
Compute programs sometimes use textures either to access dedicated texture caches or to use the texture-sampling hardware for interpolation and clamping. The former approach uses simple point samplers with linear interpolation, essentially only reading a single point. The latter approach uses the sampler hardware to interpolate and combine multiple samples. AMD hardware, as well as recent competing hardware, has a unified texture/L1 cache, so it no longer has a dedicated texture cache. But the nvcc path often caches global loads in the L2 cache, and some programs may benefit from explicit control of the L1 cache contents. We recommend the __ldg instruction for this purpose.

AMD compilers currently load all data into both the L1 and L2 caches, so __ldg is treated as a no-op.

We recommend the following for functional portability:

For programs that use textures only to benefit from improved caching, use the __ldg instruction

Programs that use texture object and reference APIs, work well on HIP
`

[FRYoussef]

I was trying to use @JackAKirk 's solution, but I'm unable to disassemble the resulting code.

#include <iostream>
#include <CL/sycl.hpp>
#include <sycl/ext/oneapi/experimental/cuda/builtins.hpp>

using namespace sycl::ext::oneapi::experimental::cuda;
using namespace sycl::ext::oneapi::experimental;

class CUDASelector : public cl::sycl::device_selector {
    public:
        int operator()(const cl::sycl::device &Device) const override {
            const std::string DriverVersion = Device.get_info<cl::sycl::info::device::driver_version>();

            if (Device.is_gpu() && (DriverVersion.find("CUDA") != std::string::npos))
                return 1;

            return 0;
        }
};

void sum(cl::sycl::nd_item<1> item, const double* __restrict__ A, const double* __restrict__ B, double* C){
        int idx = item.get_global_id(0);
        C[idx] = __ldg(&A[idx]) + __ldg(&B[idx]);  
}

int main()
{
    constexpr size_t m = 1024;
    constexpr size_t n = 128;

    cl::sycl::queue device_queue{CUDASelector{}};
    std::cout << "device: "
              << device_queue.get_device().get_info<cl::sycl::info::device::name>()
              << std::endl;

    double* A = cl::sycl::malloc_device<double>(m, device_queue);
    double* B = cl::sycl::malloc_device<double>(m, device_queue);
    double* C = cl::sycl::malloc_device<double>(m, device_queue);

    device_queue.memset(A, 1, sizeof(double)*m);
    device_queue.memset(B, 1, sizeof(double)*m);
    device_queue.wait();

   cl::sycl::range<1> num_groups{(size_t)m/n + (m % n == 0 ? 0 : 1)};
    cl::sycl::range<1> group_size{(size_t)n};

    device_queue.submit([&](cl::sycl::handler &h) {
        h.parallel_for<class sum_vectors>(cl::sycl::nd_range(num_groups*group_size, group_size), [=](cl::sycl::nd_item<1> item){
            sum(item, A, B, C);
        });
    });
    device_queue.wait();
    free(A, device_queue);
    free(B, device_queue);
    free(C, device_queue);

    exit(0);
}

I'm compiling the code with: clang++ -fsycl -O3 -w -fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend --cuda-gpu-arch=sm_86 -fsycl-device-only -fsycl-unnamed-lambda -fno-sycl-use-bitcode vec_sum2.cpp and disassembling it with cuobjdump a.out -sass -ptx. However, cuobjdump fails showing cuobjdump info : File 'a.out' does not contain device code

Is any way to disassemble it and check whether it is correct?

[JackAKirk]

I'm compiling the code with: clang++ -fsycl -O3 -w -fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend --cuda-gpu-arch=sm_86 -fsycl-device-only -fsycl-unnamed-lambda -fno-sycl-use-bitcode vec_sum2.cpp and disassembling it with cuobjdump a.out -sass -ptx. However, cuobjdump fails showing cuobjdump info : File 'a.out' does not contain device code

Is any way to disassemble it and check whether it is correct?

cuobjdump will only work on the fat binary. The easiest way to get the ptx for you is probably:

mkdir savetemps
cd savetemps
clang++ -fsycl -O3 -w -fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend --cuda-gpu-arch=sm_86 -fsycl-unnamed-lambda -save-temps ../vec_sum2.cpp

This should output a load of temporary headers/binaries/ptx in savetemps directory.

The ptx is located in the longest filename with .s extension.

[JackAKirk]

I just checked this and it worked. I can send you the ptx I got if you want. Do you also want the sass? If so, savetemps also outputs the fat binary, e.g.:

cuobjdump vec_sum2-sycl-nvptx64-nvidia-cuda-sm_86-3efds61b.o -sass

[FRYoussef]

I've tried what you posted and work well, it is using the ld.global.nc instruction.

Well done!!