[all-commits] [llvm/llvm-project] 2abcdd: [CUDA] Add support for CUDA surfaces (#132883)

[Austin Schuh via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 2abcdd8cf08b9a170e6e5ad1b9facbf71135522f
      https://github.com/llvm/llvm-project/commit/2abcdd8cf08b9a170e6e5ad1b9facbf71135522f
  Author: Austin Schuh <AustinSchuh at users.noreply.github.com>
  Date:   2025-04-03 (Thu, 03 Apr 2025)

  Changed paths:
    M clang/lib/Headers/__clang_cuda_runtime_wrapper.h
    M clang/lib/Headers/__clang_cuda_texture_intrinsics.h
    A clang/test/CodeGen/nvptx-surface.cu
    M clang/test/Headers/Inputs/include/cuda.h
    A clang/test/Headers/Inputs/include/surface_indirect_functions.h

  Log Message:
  -----------
  [CUDA] Add support for CUDA surfaces (#132883)

This adds support for all the surface read and write calls to clang. It
extends the pattern used for textures to surfaces too.

I tested this by generating all the various permutations of the calls
and argument types in a python script, compiling them with both clang
and nvcc, and comparing the generated ptx for equivilence. They all
agree, ignoring register allocation, and some places where Clang picks
different memory write instructions. An example kernel is:

```
__global__ void testKernel(cudaSurfaceObject_t surfObj, int x, float2* result) {
    *result = surf1Dread<float2>(surfObj, x, cudaBoundaryModeZero);
}
```

---------

Signed-off-by: Austin Schuh <austin.linux at gmail.com>



To unsubscribe from these emails, change your notification settings at https://github.com/llvm/llvm-project/settings/notifications