错误信息：调用hipLaunchKernelGGL时出错，错误日志如下：

KERNEL VMFault !!!! <<<<<<

PID: 22202 !!!! <<<<<<
=========> STREAM <0x2627540>: VMFault HSA QUEUE ANALYSIS <=========
STREAM <0x2627540>: >>>>>>>> DUMP KERNEL AQL PACKET <<<<<<<<<
STREAM <0x2627540>: header: 770
STREAM <0x2627540>: setup: 3
STREAM <0x2627540>: workgroup: x:256, y:1, z:1
STREAM <0x2627540>: grid: x:163840, y:1, z:1
STREAM <0x2627540>: group_segment_size: 512
STREAM <0x2627540>: private_segment_size: 16944
STREAM <0x2627540>: kernel_object: 47728628124224

SUCCESS: FIND SAME KERNEL OBJECT COMMAND IN USE LIST. useIdx: 32
STREAM <0x2627540>: >>>>>>>> FIND MATCH KERNEL COMMAND <<<<<<<<<
STREAM <0x2627540>: kernel name: ZN8spmv_csr20csr64DynamicWarpBLASIdLl8ELl32EEEvPyPKlS3_PKT_S6_S6_PS4_S4_S4
STREAM <0x2627540>: >>>>>>>> DUMP KERNEL ARGS: size: 128 <<<<<<<<<

00 00 80 ba 69 2b 00 00 00 89 20 08 6a 2b 00 00
80 8b 20 08 6a 2b 00 00 f8 ff 1f 08 6a 2b 00 00
e8 16 23 08 6a 2b 00 00 70 b6 21 08 6a 2b 00 00
70 b6 21 08 6a 2b 00 00 00 00 00 00 00 00 f0 bf
00 00 00 00 00 00 f0 3f 00 00 00 00 00 00 00 00
00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00
00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00
00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00

STREAM <0x2627540>: >>>>>>>> DUMP KERNEL ARGS PTR INFO <<<<<<<<<
STREAM <0x2627540>: ptr arg index: 0, ptr: 0x2b69ba800000
STREAM <0x2627540>: host ptr: 0x2b69ba800000, device ptr: 0x2b69ba800000, unaligned ptr: 0x2b69ba800000
STREAM <0x2627540>: size byte: 8
STREAM <0x2627540>: ptr arg index: 1, ptr: 0x2b6a08208900
STREAM <0x2627540>: host ptr: 0x2b6a08200000, device ptr: 0x2b6a08200000, unaligned ptr: 0x2b6a08200000
STREAM <0x2627540>: size byte: 15887499264
STREAM <0x2627540>: ptr arg index: 2, ptr: 0x2b6a08208b80
STREAM <0x2627540>: host ptr: 0x2b6a08200000, device ptr: 0x2b6a08200000, unaligned ptr: 0x2b6a08200000
STREAM <0x2627540>: size byte: 15887499264
STREAM <0x2627540>: ptr arg index: 4, ptr: 0x2b6a082316e8
STREAM <0x2627540>: host ptr: 0x2b6a08200000, device ptr: 0x2b6a08200000, unaligned ptr: 0x2b6a08200000
STREAM <0x2627540>: size byte: 15887499264
STREAM <0x2627540>: ptr arg index: 5, ptr: 0x2b6a0821b670
STREAM <0x2627540>: host ptr: 0x2b6a08200000, device ptr: 0x2b6a08200000, unaligned ptr: 0x2b6a08200000
STREAM <0x2627540>: size byte: 15887499264
STREAM <0x2627540>: ptr arg index: 6, ptr: 0x2b6a0821b670
STREAM <0x2627540>: host ptr: 0x2b6a08200000, device ptr: 0x2b6a08200000, unaligned ptr: 0x2b6a08200000
STREAM <0x2627540>: size byte: 15887499264

=========> STREAM <0xd79a400>: VMFault HSA QUEUE ANALYSIS <=========
=========> STREAM <0xd60d0d0>: VMFault HSA QUEUE ANALYSIS <=========
=========> STREAM <0xd719f50>: VMFault HSA QUEUE ANALYSIS <=========
=========> STREAM <0xdbee710>: VMFault HSA QUEUE ANALYSIS <=========
=========> STREAM <0x26a8e60>: VMFault HSA QUEUE ANALYSIS <=========
=========> STREAM <0xd699b20>: VMFault HSA QUEUE ANALYSIS <=========
=========> STREAM <0xd89ace0>: VMFault HSA QUEUE ANALYSIS <=========
=========> STREAM <0xd81a870>: VMFault HSA QUEUE ANALYSIS <=========

KERNEL VMFault Analysis END !!!! <<<<<<

hsaQueueDispatchCallBack: resultStatus=41,lastHipError=719,internalErrorID=131336

出现错误的核方法如下：

global void csr64DynamicVectorBLAS(unsigned long long* __restrict cudaRowCounter, const int64_t* __restrict rowOffsets, const int64_t* __restrict colIndexValues,
const T* __restrict numericalValues, const T* __restrict vectorX, const T* __restrict inVectorY, T* vectorY, const T alpha, const T beta)

{
int64_t i;
T sum;
int64_t row;
int64_t rowStart, rowEnd;
const int64_t laneId = threadIdx.x % THREADS_PER_VECTOR; /lane index in the vector/
const int64_t vectorId = threadIdx.x / THREADS_PER_VECTOR; /vector index in the block/

shared volatile int64_t space[MAX_NUM_VECTORS_PER_BLOCK][2];

/get the row index/
if (laneId == 0) {
row = atomicAdd(cudaRowCounter, 1);
}
/broadcast the value to other lanes from lane 0/
row = __shfl((int)row, 0, THREADS_PER_VECTOR);

/check the row range/
while (row < _cudaNumRows) {

/use two threads to fetch the row offset/
if (laneId < 2) {
space[vectorId][laneId] = rowOffsets[row + laneId];
}
rowStart = space[vectorId][0];
rowEnd = space[vectorId][1];

/there are non-zero elements in the current row/
sum = 0;
/compute dot product/
if (THREADS_PER_VECTOR == 32) {

/ensure aligned memory access/
i = rowStart - (rowStart & (THREADS_PER_VECTOR - 1)) + laneId;

/process the unaligned part/
if (i >= rowStart && i < rowEnd) {
sum += numericalValues[i] * DOUBLE_VECTOR_GET(vectorX, colIndexValues[i]);
}

/process the aligned part/
for (i += THREADS_PER_VECTOR; i < rowEnd; i += THREADS_PER_VECTOR) {
sum += numericalValues[i] * DOUBLE_VECTOR_GET(vectorX, colIndexValues[i]);
}
} else {
/regardless of the global memory access alignment/
for (i = rowStart + laneId; i < rowEnd; i +=
THREADS_PER_VECTOR) {
sum += numericalValues[i] * DOUBLE_VECTOR_GET(vectorX, colIndexValues[i]);
}
}
/intra-vector reduction/
sum *= alpha;
for (i = THREADS_PER_VECTOR >> 1; i > 0; i >>= 1) {
sum += shfl_down_64bits<T>(sum, i, THREADS_PER_VECTOR);
}

/save the results and get a new row/
if (laneId == 0) {
/save the results/
vectorY[row] = sum + beta * DOUBLE_VECTOR_GET(inVectorY, row);

/get a new row index/
row = atomicAdd(cudaRowCounter, 1);
}
row = __shfl((int)row, 0, THREADS_PER_VECTOR);
}/while/
}

调用HIP核函数的方式：
hipLaunchKernelGGL(HIP_KERNEL_NAME(spmv_csr::csr64DynamicVectorBLAS<double, 2, MAX_NUM_THREADS_PER_BLOCK / 2>), numThreadBlocks, numThreadsPerBlock, 0, 0, _cudaRowCounters[i], _rowOffsets[i], _colIndexValues[i], _numericalValues[i], _vectorX[i], _vectorY[i], _vectorY[i], _alpha, _beta);
请麻烦帮忙分析一下这是什么错误？一般是由什么导致的。