cuda & cublas: 使用 cublas 后调用全局函数

Question

我编写的程序涉及一些矩阵向量乘法和最小二乘求解，所有这些都使用 cublas & cula API。该程序将迭代多次。在每一步中，我必须将一个矩阵的特定行全部设置为零。

我尝试将整个矩阵（50*1000 或更大）复制到 cpu 并将一行设置为零，然后将矩阵复制回来，但这太耗时了，因为程序会迭代 10 次或更多次。所以我决定写一个核函数。

这样的全局函数：

__global__ void Setzero(float* A, int index) /* A is the matrix and in col-major , index is the row I want to set zero */
{
    int ind=blockDim.x*blockIdx.x+threadIdx.x;
    if( ((ind%N)==index ) && (ind<50000) )  //notice matrix is in col-major ,matrix size is 50000
    {   
    A[ind]=0.0;
        ind+=blockDim.x*blockIdx.x;
    }
    else    ;
        __syncthreads();   
}

问题是我何时执行此操作（在调用函数之前使用 cublas）：

cudaMalloc((void**)&A_Gpu_trans,sizeof(float)*50000);
cudaMemcpy(A_Gpu_trans,A_trans,sizeof(float)*M*N,cudaMemcpyHostToDevice);
cublasSgemv_v2(handle,CUBLAS_OP_N,1000,50,&al,A_Gpu_trans,1000,err_gpu,1,&beta,product,1);
dim3 dimBlock(16,1);
dim3 dimGrid((50000-1)/16+1,1);
Setzero<<<dimGrid,dimBlock>>>(A_Gpu_trans,Index);

返回错误：

a __host__ function("Setzero") redeclared with __global__.

还有一个错误：

MSB3721：命令“C:\Program Files\NVIDIA GPU 计算 工具包\CUDA\v5.5\bin\nvcc.exe" -gencode=arch=compute_10,code=\"sm_10,compute_10\" --use-local-env --cl-version 2010 -ccbin "D:\Program Files\Microsoft Visual Studio 10.0\VC\bin" -I" C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.5\include" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.5\include" -G --keep-dir 调试 -maxrregcount=0 --machine 32 --compile -cudart static -g -DWIN32 -D_DEBUG -D_CONSOLE -D_MBCS -Xcompiler "/EHsc /W3 /nologo /Od /Zi /RTC1 /MDd " -o Debug\kernel.cu.obj "C: \Users\Administrator\documents\visual studio 2010\Projects\OOmp\OOmp\kernel.cu"”返回2。

奇怪的是，当我只使用 cublas & cula API 时，我可以得到正确的答案。

Answer 1

虽然您没有在问题中显示它，但您显然在代码中的某处有另一个名为 Setzero 的主机函数。简单的解决方案是将内核重命名为其他名称。

CUDA 工具链发出错误的根本原因是运行时 API 中的 Setzero<<< >>> 内核调用语法导致 CUDA 前端创建一个与内核同名的主机函数，并带有匹配的参数列表并替换内核启动以调用该函数。此主机函数包含启动内核所需的 API 调用。通过使用与内核同名的另一个主机函数，您会破坏此过程并导致您看到的编译错误。

Answer 2

另外，你的函数既错误又非常低效......

你不能在这样的条件下调用同步线程，这可能会导致挂起。这里似乎也完全没有必要。

更重要的是，您为每个矩阵条目启动一个线程，而实际上只有 1/N 的线程执行任何操作。

更好的方法是只启动与将设置为零的条目相对应的线程。像这样的：

__global__ void Setzero(float* A, int index) 
{
  int ind=blockDim.x*blockIdx.x+threadIdx.x;
  if (ind < M)   
    A[index+N*ind]=0.0;
}

然后你启动 M 个线程（或者更确切地说，ceil(M/256) 个线程块，每个线程有 256 个线程，或者你想要的任何块大小）。

例如：

int block_size = 256; // usually a good choice
int num_blocks = (M + block_size - 1) / block_size;
Setzero<<<num_blocks, block_size>>>(A, index);