error in sampling ， error: identifier "cudaCGGetIntrinsicHandle" is undefined return (cudaCGGetIntrinsicHandle(cudaCGScopeMultiGrid))

When I tested under windows server 2019, get the following error.

it is due the cuda toolkit version , or cuDNN issue(version cudnn-windows-x86_64-8.9.6.50_cuda12-archive) or the display card is to low?

my envirement:
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2023 NVIDIA Corporation
Built on Tue_Aug_15_22:09:35_Pacific_Daylight_Time_2023
Cuda compilation tools, release 12.2, V12.2.140
Build cuda_12.2.r12.2/compiler.33191640_0

[SPARSE] Backend: spconv, Attention: xformers
Warp 1.5.0 initialized:
CUDA Toolkit 12.6, Driver 12.2
Devices:
"cpu" : "Intel64 Family 6 Model 85 Stepping 4, GenuineIntel"
"cuda:0" : "Quadro P600" (2 GiB, sm_61, mempool enabled)
Kernel cache:
C:\Users\Administrator\AppData\Local\NVIDIA\warp\Cache\1.5.0
D:\miniconda3\envs\trellis1\lib\site-packages\gradio_client\utils.py:1097: UserWarning: file() is deprecated and will be removed in a future version. Use handle_file() instead.
warnings.warn(
[SPARSE][CONV] spconv algo: auto
[ATTENTION] Using backend: xformers

error info:
Sampling: 0%| | 0/12 [00:00<?, ?it/s]D:\miniconda3\envs\trellis1\Lib\site-packages\cumm\include\tensorview/core/nvrtc/core.h(39): warning #867-D: declaration of "size_t" does not match the expected type "unsigned long long"
typedef unsigned long size_t;
^

Remark: The warnings can be suppressed with "-diag-suppress "

D:\miniconda3\envs\trellis1\Lib\site-packages\cumm\include\tensorview/core/nvrtc/limits.h(689): warning #61-D: integer operation result is out of range
return -INT_MIN - 1;
^

D:\miniconda3\envs\trellis1\Lib\site-packages\cumm\include\tensorview/core/nvrtc/limits.h(689): warning #61-D: integer operation result is out of range
return -INT_MIN - 1;
^

D:\miniconda3\envs\trellis1\Lib\site-packages\cumm\include\tensorview/core/array.h(821): warning #161-D: unrecognized #pragma
#pragma GCC diagnostic pop
^

**C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.2\include\cooperative_groups/details/helpers.h(453): error: identifier "cudaCGGetIntrinsicHandle" is undefined
return (cudaCGGetIntrinsicHandle(cudaCGScopeMultiGrid));
^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.2\include\cooperative_groups/details/helpers.h(458): error: identifier "cudaCGSynchronize" is undefined
cudaError_t err = cudaCGSynchronize(handle, 0);
^

_**C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.2\include\cooperative_groups/details/helpers.h(464): error: identifier "cudaCGGetSize" is undefined
cudaCGGetSize(&numThreads, NULL, handle);
^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.2\include\cooperative_groups/details/helpers.h(471): error: identifier "cudaCGGetRank" is undefined
cudaCGGetRank(&threadRank, NULL, handle);
^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.2\include\cooperative_groups/details/helpers.h(478): error: identifier "cudaCGGetRank" is undefined
cudaCGGetRank(NULL, &gridRank, handle);
^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.2\include\cooperative_groups/details/helpers.h(485): error: identifier "cudaCGGetSize" is undefined
cudaCGGetSize(NULL, &numGrids, handle);_
^

spconv/conv_params/ConvParams.h(52): warning #427-D: qualified name is not allowed in member declaration
host device ConvParams::ConvParams(ConvProblem problem, const tv::half_t* A, const tv::half_t* B, tv::half_t* C, const tv::half_t* D, const uint32_t* mask_ptr, const int* mask_argsort_ptr, const int* indice_ptr, uint32_t* mask_out_ptr, uint32_t mask_filter, bool reverse_mask, float alpha, float beta, float act_alpha, float act_beta, tv::gemm::Activation act_type, int split_k_slices, bool d_is_bias) : problem(problem), itera_params_(problem, indice_ptr, mask_argsort_ptr), mask_out_ptr(mask_out_ptr), iterb_params_(problem, LayoutB::from_shape(problem.get_weight_shape())), ptr_A(A), ptr_B(B), ptr_C(C), ptr_D(D), mask_ptr(mask_ptr), mask_filter(mask_filter), reverse_mask(reverse_mask), alpha(alpha), beta(beta), act_alpha(act_alpha), act_beta(act_beta), act_type(act_type) {
^

6 errors detected in the compilation of "kernel.cu.cu".

Build Error. Kernel Code:
#pragma once
#include <cumm/common/TensorViewNVRTCKernel.h>
#include <cumm/gemm/layout/RowMajor.h>
#include <cumm/gemm/layout/ColumnMajor.h>
#include <cumm/common/GemmBasicKernel.h>
#include <cumm/common/GemmBasic.h>
#include <spconv/inpitera/ForwardDgradSparseIOIterator.h>
#include <spconv/inpiterb/WeightIteratorDP4A.h>
#include <spconv/layouta/TensorGeneric.h>
#include <spconv/layoutb/TensorGeneric.h>
#include <spconv/layoutc/TensorGeneric.h>
#include <spconv/gemm_smem_storage/BlockMmaStorage.h>
#include <spconv/out_smem_storage/OutputSmemStorage.h>
#include <spconv/conv_params/ConvParams.h>
#include <spconv/conv_params/ConvProblem.h>
#include <spconv/out_iter/OutIterator.h>
#include <spconv/out_iter_const/OutIterator.h>
#include <spconv/out_op/LinearCombination.h>
#include <spconv/mma/Mma.h>
#include <spconv/mma_miterd/MaskIGemmIteratorMaskLoaderDynamic.h>
#include <spconv/output/Output.h>
namespace spconv {
using TensorViewNVRTCKernel = cumm::common::TensorViewNVRTCKernel;
using RowMajor = cumm::gemm::layout::RowMajor;
using ColumnMajor = cumm::gemm::layout::ColumnMajor;
using GemmBasicKernel = cumm::common::GemmBasicKernel;
using GemmBasic = cumm::common::GemmBasic;
using InputIteratorA = spconv::inpitera::ForwardDgradSparseIOIterator;
using InputIteratorB = spconv::inpiterb::WeightIteratorDP4A;
using LayoutA = spconv::layouta::TensorGeneric;
using LayoutB = spconv::layoutb::TensorGeneric;
using LayoutC = spconv::layoutc::TensorGeneric;
using BlockMmaStorage = spconv::gemm_smem_storage::BlockMmaStorage;
using OutputSmemStorage = spconv::out_smem_storage::OutputSmemStorage;
using ConvParams = spconv::conv_params::ConvParams;
using ConvProblem = spconv::conv_params::ConvProblem;
using OutIter = spconv::out_iter::OutIterator;
using ConstOutIter = spconv::out_iter_const::OutIterator;
using OutputOp = spconv::out_op::LinearCombination;
using Mma = spconv::mma::Mma;
using MaskIGemmIteratorDynamic = spconv::mma_miterd::MaskIGemmIteratorMaskLoaderDynamic;
using Output = spconv::output::Output;
constant int kSizeOfParams = sizeof(ConvParams);
constant int kNumThreads = 256;
constant int kSmemSize = 10752;
constant tv::array<uint8_t, sizeof(ConvParams)> params_raw;
global void conv_kernel() {

ConvParams params = (reinterpret_cast<ConvParams>(
params_raw.data()));
#if (defined(CUDA_ARCH) && (CUDA_ARCH >= 0))
constexpr bool kSplitKSerial = false;
extern shared uint8_t SharedStorage[];
auto gemm_shared_mem =
reinterpret_cast<BlockMmaStorage *>(SharedStorage);
auto out_shared_mem =
reinterpret_cast<OutputSmemStorage >(SharedStorage);
int tile_offset_m = blockIdx.x;
int tile_offset_n = blockIdx.y;
int tile_offset_k = blockIdx.z;
if (tile_offset_m >= params.grid_dims.x ||
tile_offset_n >= params.grid_dims.y) {
return;
}
tv::array<int, 2> block_offset_A{tile_offset_m * 32, tile_offset_k * 16};
tv::array<int, 2> block_offset_B{tile_offset_n * 128, tile_offset_k * 16};
int thread_idx = threadIdx.x;
InputIteratorA input_iter_A(
params.itera_params_, params.problem, params.ptr_A,
thread_idx,
block_offset_A);
InputIteratorB input_iter_B(
params.iterb_params_, params.problem, params.ptr_B,
thread_idx,
block_offset_B);
int warp_idx = __shfl_sync(0xffffffff, threadIdx.x / 32, 0);
int lane_idx = threadIdx.x % 32;
int warp_mn =
warp_idx % (1 * 4);
int warp_idx_k =
warp_idx / (1 * 4);
int warp_m = warp_mn % 1;
int warp_n = warp_mn / 1;
uint32_t kmask = 0;
tv::array<uint32_t, 1> masks;
masks.clear();
TV_PRAGMA_UNROLL
for (int i = 0; i < 1; ++i){
if (tile_offset_m * 32 + i * 32 + lane_idx < params.m){
masks[i] = params.mask_ptr[tile_offset_m * 32 + i * 32 + lane_idx];
}
}
TV_PRAGMA_UNROLL
for (int i = 0; i < 1; ++i){
kmask |= masks[i];
}
// perform a warp reduce to get block mask
TV_PRAGMA_UNROLL
for (int mask = 16; mask > 0; mask /= 2) {
kmask |= shfl_xor_sync(0xffffffff, kmask, mask, 32);
}
kmask &= params.mask_filter;
if (params.mask_out_ptr != nullptr){
params.mask_out_ptr[tile_offset_m] = kmask;
}
Mma mma(gemm_shared_mem, thread_idx, warp_idx_k, warp_m, warp_n, lane_idx);
tv::array<float, 32, 0> accumulators;
accumulators.clear();
if (!kSplitKSerial || params.gemm_k_iterations > 0){
if (kmask != 0){
mma(params.gemm_k_iterations, accumulators, input_iter_A, input_iter_B, accumulators, kmask, params.problem.kernel_volume);
}
}
// // C = alpha * A@B + beta * D, D can be C
OutputOp output_op(params.alpha, params.beta, params.act_alpha, params.act_beta, params.act_type);
tv::array<int, 2> block_offset_C{tile_offset_m * 32,
tile_offset_n * 128};
tv::array<int, 2> block_extent_C{params.m, params.n};
OutIter out_iter_C(params.out_params, params.ptr_C, block_extent_C,
block_offset_C,
thread_idx);
ConstOutIter out_iter_source(params.out_params_source, params.ptr_D, block_extent_C,
block_offset_C,
thread_idx);
Output out(out_shared_mem, thread_idx, warp_idx_k, warp_m, warp_n, lane_idx);
out.run(output_op, accumulators, out_iter_C, out_iter_source);
#else
tv::printf2_once("this arch isn't supported!");
assert(0);
#endif
}
global void nvrtc_kernel_cpu_out(tv::gemm::SparseConvNVRTCParams p, ConvParams out) {

ConvProblem problem(p.N, p.C, p.K, p.kernel_volume, static_casttv::gemm::ConvMode(p.mode),
p.split_k_slices, p.groups);
ConvParams params(problem, reinterpret_cast<const tv::half_t*>(p.ptr_A), reinterpret_cast<const tv::half_t*>(p.ptr_B), reinterpret_casttv::half_t*(p.ptr_C), reinterpret_cast<const tv::half_t*>(p.ptr_D), p.mask_ptr, p.mask_argsort_ptr, p.indice_ptr, p.mask_out_ptr, p.mask_filter, p.reverse_mask, float(p.alpha), float(p.beta), float(p.act_alpha), float(p.act_beta), static_casttv::gemm::Activation(p.act_type), p.split_k_slices, p.d_is_bias);
auto out_ptr = reinterpret_cast<uint8_t*>(out);
auto param_ptr = reinterpret_cast<const uint8_t*>(&params);
for (auto i : tv::KernelLoopX(sizeof(ConvParams))){
out_ptr[i] = param_ptr[i];
}
// out[0] = params;
}
struct ConvKernel {
};
} // namespace spconv
Sampling: 0%| | 0/12 [00:02<?, ?it/s]
Traceback (most recent call last):
File "D:\miniconda3\envs\trellis1\lib\site-packages\gradio\queueing.py", line 536, in process_events
response = await route_utils.call_process_api(
File "D:\miniconda3\envs\trellis1\lib\site-packages\gradio\route_utils.py", line 322, in call_process_api
output = await app.get_blocks().process_api(
File "D:\miniconda3\envs\trellis1\lib\site-packages\gradio\blocks.py", line 1935, in process_api
result = await self.call_function(
File "D:\miniconda3\envs\trellis1\lib\site-packages\gradio\blocks.py", line 1520, in call_function
prediction = await anyio.to_thread.run_sync( # type: ignore
File "D:\miniconda3\envs\trellis1\lib\site-packages\anyio\to_thread.py", line 56, in run_sync
return await get_async_backend().run_sync_in_worker_thread(
File "D:\miniconda3\envs\trellis1\lib\site-packages\anyio_backends_asyncio.py", line 2505, in run_sync_in_worker_thread
return await future
File "D:\miniconda3\envs\trellis1\lib\site-packages\anyio_backends_asyncio.py", line 1005, in run
result = context.run(func, *args)
File "D:\miniconda3\envs\trellis1\lib\site-packages\gradio\utils.py", line 826, in wrapper
response = f(*args, **kwargs)
File "D:\0AI\TRELLIS\app.py", line 140, in image_to_3d
outputs = pipeline.run(
File "D:\miniconda3\envs\trellis1\lib\site-packages\torch\utils_contextlib.py", line 116, in decorate_context
return func(*args, **kwargs)
File "D:\0AI\TRELLIS\trellis\pipelines\trellis_image_to_3d.py", line 283, in run
slat = self.sample_slat(cond, coords, slat_sampler_params)
File "D:\0AI\TRELLIS\trellis\pipelines\trellis_image_to_3d.py", line 243, in sample_slat
slat = self.slat_sampler.sample(
File "D:\miniconda3\envs\trellis1\lib\site-packages\torch\utils_contextlib.py", line 116, in decorate_context
return func(*args, **kwargs)
File "D:\0AI\TRELLIS\trellis\pipelines\samplers\flow_euler.py", line 199, in sample
return super().sample(model, noise, cond, steps, rescale_t, verbose, neg_cond=neg_cond, cfg_strength=cfg_strength, cfg_interval=cfg_interval, **kwargs)
File "D:\miniconda3\envs\trellis1\lib\site-packages\torch\utils_contextlib.py", line 116, in decorate_context
return func(*args, **kwargs)
File "D:\0AI\TRELLIS\trellis\pipelines\samplers\flow_euler.py", line 112, in sample
out = self.sample_once(model, sample, t, t_prev, cond, **kwargs)
File "D:\miniconda3\envs\trellis1\lib\site-packages\torch\utils_contextlib.py", line 116, in decorate_context
return func(*args, **kwargs)
File "D:\0AI\TRELLIS\trellis\pipelines\samplers\flow_euler.py", line 73, in sample_once
pred_x_0, pred_eps, pred_v = self._get_model_prediction(model, x_t, t, cond, **kwargs)
File "D:\0AI\TRELLIS\trellis\pipelines\samplers\flow_euler.py", line 43, in _get_model_prediction
pred_v = self._inference_model(model, x_t, t, cond, **kwargs)
File "D:\0AI\TRELLIS\trellis\pipelines\samplers\guidance_interval_mixin.py", line 11, in _inference_model
pred = super()._inference_model(model, x_t, t, cond, **kwargs)
File "D:\0AI\TRELLIS\trellis\pipelines\samplers\flow_euler.py", line 40, in _inference_model
return model(x_t, t, cond, **kwargs)
File "D:\miniconda3\envs\trellis1\lib\site-packages\torch\nn\modules\module.py", line 1736, in _wrapped_call_impl
return self._call_impl(*args, **kwargs)
File "D:\miniconda3\envs\trellis1\lib\site-packages\torch\nn\modules\module.py", line 1747, in _call_impl
return forward_call(*args, **kwargs)
File "D:\0AI\TRELLIS\trellis\models\structured_latent_flow.py", line 245, in forward
h = block(h, t_emb)
File "D:\miniconda3\envs\trellis1\lib\site-packages\torch\nn\modules\module.py", line 1736, in _wrapped_call_impl
return self._call_impl(*args, **kwargs)
File "D:\miniconda3\envs\trellis1\lib\site-packages\torch\nn\modules\module.py", line 1747, in _call_impl
return forward_call(*args, **kwargs)
File "D:\0AI\TRELLIS\trellis\models\structured_latent_flow.py", line 59, in forward
h = self.conv1(h)
File "D:\miniconda3\envs\trellis1\lib\site-packages\torch\nn\modules\module.py", line 1736, in _wrapped_call_impl
return self._call_impl(*args, **kwargs)
File "D:\miniconda3\envs\trellis1\lib\site-packages\torch\nn\modules\module.py", line 1747, in _call_impl
return forward_call(*args, **kwargs)
File "D:\0AI\TRELLIS\trellis\modules\sparse\conv\conv_spconv.py", line 26, in forward
new_data = self.conv(x.data)
File "D:\miniconda3\envs\trellis1\lib\site-packages\torch\nn\modules\module.py", line 1736, in _wrapped_call_impl
return self._call_impl(*args, **kwargs)
File "D:\miniconda3\envs\trellis1\lib\site-packages\torch\nn\modules\module.py", line 1747, in _call_impl
return forward_call(*args, kwargs)
File "D:\miniconda3\envs\trellis1\lib\site-packages\spconv\pytorch\conv.py", line 755, in forward
return self._conv_forward(self.training,
File "D:\miniconda3\envs\trellis1\lib\site-packages\spconv\pytorch\conv.py", line 467, in _conv_forward
out_features, , _ = ops.implicit_gemm(
File "D:\miniconda3\envs\trellis1\lib\site-packages\spconv\pytorch\ops.py", line 1513, in implicit_gemm
mask_width, tune_res_cpp = ConvGemmOps.implicit_gemm(
File "D:\miniconda3\envs\trellis1\lib\site-packages\spconv\algo.py", line 208, in cached_get_nvrtc_params
mod, ker = self.compile_nvrtc_module(desp)
File "D:\miniconda3\envs\trellis1\lib\site-packages\spconv\algo.py", line 196, in compile_nvrtc_module
mod = CummNVRTCModule([kernel],
File "D:\miniconda3\envs\trellis1\lib\site-packages\cumm\nvrtc_init.py", line 389, in init
super().init(mod_params.code,
File "D:\miniconda3\envs\trellis1\lib\site-packages\cumm\nvrtc_init.py", line 280, in init
super().init(code,
**File "D:\miniconda3\envs\trellis1\lib\site-packages\cumm\tensorview_init.py", line 191, in init
self._mod = _NVRTCModule(code, headers, opts, program_name,
RuntimeError: D:\a\cumm\cumm\include\tensorview/cuda/nvrtc.h(96)
compileResult == NVRTC_SUCCESS assert faild. nvrtc compile failed.

Start with the reported NVRTC build failure in cumm/include/tensorview/core/nvrtc and the CUDA cooperative_groups/details/helpers.h header, then inspect the generated spconv ConvParams.h kernel. Reproduce the sampling failure on Windows Server 2019 with CUDA 12.2, CUDA 12.6, and the Quadro P600; done means the sparse-convolution sampling kernel compiles or the supported-version limitation is documented.