Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this request and that I agree to the Code of Conduct

Area

libcu++

Is your feature request related to a problem? Please describe.

As a CUDA developer using libcu++, I want to be able to use architecture dependent features of libcudacxx in my CUDA application. For any given libcudacxx header and feature, I need to be able to do the following:

#include <cuda/header>

__global__ void kernel(...){
  NV_DISPATCH_TARGET(
    NV_IS_EXACTLY_SM_60,   ( do_sm60_thing(); ),
    NV_PROVIDES_SM_70, ( do_sm70_thing(); ),
    NV_PROVIDES_SM_90, ( do_sm90_thing(); )
  )
}

I need to be able to compile this file with any set of architectures (-gencode arch=compute_XX,code=sm_XX) and for it to be able to compile and link successfully so long as I am always careful to use an architecture dependent feature in an appropriately guarded code path, whether using NV_IF_TARGET or __CUDA_ARCH__.

However, this does not work universally today. For example, the following fails to compile when compiled with -gencode arch=compute_52,code=sm_52 -gencode arch=compute_70,code=sm_70

#include <cuda/atomic>
#include <nv/target>

__global__ void kernel(){
  NV_IF_TARGET(
    NV_PROVIDES_SM_70,
      cuda::atomic<int> i;
  )
}

https://godbolt.org/z/ddMaW65Ej

This is because the cuda/atomic header will unconditionally error any time it is included in a TU that compiles for an architecture less than sm60, even if the feature is never used in code paths for the unsupported architecture.

A similar problem exists with cuda/barrier: https://godbolt.org/z/aEjsMT5YK

Describe the solution you'd like

I should be able to do the following with all libcu++ headers and features:

#include <cuda/header>

__global__ void kernel(...){
  NV_DISPATCH_TARGET(
    NV_IS_EXACTLY_SM_60,   ( do_sm60_thing(); ),
    NV_PROVIDES_SM_70, ( do_sm70_thing(); ),
    NV_PROVIDES_SM_90, ( do_sm90_thing(); )
  )
}

Describe alternatives you've considered

If libcu++ doesn't do this, then I am forced to use lower level things like atomicAdd() or inline PTX.

Additional context

Related issues:
#997
#1082
#624