Assertion fail "A key type larger than 8B is supported for only sm_70 and up." on Ampere GPUs

[nicknytko]

I'm getting the assertion failure:

static_map.cuh(157): error: static assertion failed with "A key type larger than 8B is supported for only sm_70 and up."
                detected during instantiation of class "cuco::static_map<Key, Value, Scope, Allocator> [with Key=coordinate_pair_t, Value=int64_t, Scope=cuda::__3::thread_scope_device, Allocator=cuco::cuda_allocator<char>]"

when trying to compile on a node with A100 GPUs and the flag -gencode=arch=compute_80,code=sm_80 passed to nvcc. I'm wondering if there's some other obvious thing I'm missing that might be preventing my code from compiling? I didn't want to open an issue because this is probably a user error, not a library error 🙂

For reference, this is the key struct I'm trying to use, which indeed is larger than 8 bytes:

struct coordinate_pair_t {
    int64_t row;
    int64_t col;

    __host__ __device__ coordinate_pair_t(int64_t _row, int64_t _col): row(_row), col(_col) {}

    __device__ bool operator==(const coordinate_pair_t& other) const {
        return (row == other.row && col == other.col);
    }
};

[PointKernel]

@nicknytko Thanks for creating cuCo's first discussion! 🎉

It seems you were building for multiple architectures which include pre-volta ones.

when trying to compile on a node with A100 GPUs and the flag -gencode=arch=compute_80,code=sm_80 passed to nvcc.

It is worth double-checking if you are passing the right build options to nvcc. Can you show the build log right before the error with make VERBOSE=1?

If you build with cmake. Specifying CUDA architectures should be as simple as cmake ../ -DCMAKE_CUDA_ARCHITECTURES=80.

[nicknytko]

These are the options I'm using (formatted for slightly easier reading, and ... inserted to omit some sensitive details)

/cvmfs/soft.computecanada.ca/easybuild/software/2020/Core/cudacore/11.4.2/bin/nvcc  
-I/.../ext/cuCollections/include
-I/localscratch/.../env/lib/python3.10/site-packages/torch/include 
-I/localscratch/.../env/lib/python3.10/site-packages/torch/include/torch/csrc/api/include
-I/localscratch/.../env/lib/python3.10/site-packages/torch/include/TH 
-I/localscratch/.../env/lib/python3.10/site-packages/torch/include/THC
-I/cvmfs/soft.computecanada.ca/easybuild/software/2020/Core/cudacore/11.4.2/include 
-I/.../env/include 
-I/cvmfs/soft.computecanada.ca/easybuild/software/2020/avx2/Core/python/3.10.2/include/python3.10 
-c -c /.../sparse.cu 
-o /.../build/temp.linux-x86_64-cpython-310/cpp/sparse.o 
-D__CUDA_NO_HALF_OPERATORS__ 
-D__CUDA_NO_HALF_CONVERSIONS__ 
-D__CUDA_NO_BFLOAT16_CONVERSIONS__ 
-D__CUDA_NO_HALF2_OPERATORS__ 
--expt-relaxed-constexpr --compiler-options ''"'"'-fPIC'"'"'' -std=c++17 
-DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' 
'-DPYBIND11_BUILD_ABI="_cxxabi1013"' -DTORCH_EXTENSION_NAME=... -D_GLIBCXX_USE_CXX11_ABI=1 
-gencode=arch=compute_80,code=sm_80

The make command is being run through Torch's extension setup code, so the above is auto generated. I would think that the last line would compile for compute 8.0 only, right..?

[sleeepyjack]

There seems to be an issue when building with CUDA 11.4.0-11.4.4.
Here is a reproducer: https://godbolt.org/z/aPrc6v93b

[sleeepyjack]

Tracked in #245

[nicknytko]

Good catch! I pulled in CUDA 11.7 and am able to get this compiling now. Going to mark this as "answered" because it's a simple enough workaround.

[mfbalin]

How can we compile our code if we also compile for older GPUs along with newer GPUs? I want to use cuco::static_map in my code but currently don't know how to do it. I want the cuco::static_map code path to be selected only for sm_70 and up in PR: dmlc/dgl#7239