CUDA 三重循环

Question

我是 CUDA 的新手，我正在努力将 C 代码转换为 CUDA C，它成功构建但一直崩溃。三重循环功能肯定是错误的，我不知道我应该改变什么。

函数调用：

for (z=0;z<=max;z++)    
    {
    correlationsum=coefficient(x, n, dim, z);
    printf("result for epsilon %d returns %d\n", z, correlation_sum);
    }

函数

long coefficient(int vctr[40000], long numberofpoints, int coefficientrow, int epsilon)
{
long i, j, k, sum, numberofpairs;
long sq_epsilon;
sq_epsilon=epsilon*epsilon;
numberofpairs=0;
for (i=1;i<=numberofpoints-coefficientrow;i++)
    {
    sum=0;
    for (j=i+1;j<=numberofpoints+1-coefficientrow;j++)
        {
        for (k=0;k<coefficientrow;k++)
            {
            sum=sum+(vctr[i+k]-vctr[j+k])*(vctr[i+k]-vctr[j+k]);                
            }
        if(sum<sq_epsilon)  
            {
            numberofpairs++;
            sum=0;
            }
        }
    }
return (numberofpairs);
}

我在限制 GPU 部分的功能时遇到问题，所以它不会越界（例如 k 小于上面的系数行）。我看到可以分配 block/threadids 并使用 if 函数。我已经尝试过了，但是在三重 for 循环中有点……奇怪。

这里几乎是完整的代码。

    #define THREADS 1024
__global__ void coefficient(int *vctr, int numberofpoints, int coefficient_row, int epsilon, int *numbofpairs){
        int i = blockIdx.x * blockDim.x + threadIdx.x;
        int j = blockIdx.y * blockDim.y + threadIdx.y;
        int k = blockIdx.z * blockDim.z + threadIdx.z;
        int sum;
        numbofpairs = 0;
        int sq_epsilon = epsilon*epsilon;


        if (i <= numberofpoints - coefficient_row)
        {
            sum = 0;
            if (j <= numberofpoints + 1 - coefficient_row)
            {
                if (k < coefficient_row)
                    sum = sum + (vctr[i + k] - vctr[j + k])*(vctr[i + k] - vctr[j + k]);
                if (sum < sq_epsilon){
                    numbofpairs++;
                    sum = 0;
    }}}}    

int main()
{
int n, dim, max, z;
int *d_n, *d_dim, *d_z, *d_x, *d_numbofpairs;
int x[40000], correlation_sum = 0;
    n=10;  
    max=10;
    dim=3;  

    cudaMalloc((void **)&d_n, sizeof(int));
    cudaMalloc((void **)&d_dim, sizeof(int));
    cudaMalloc((void **)&d_z, sizeof(int));
    cudaMalloc((void **)&d_x, sizeof(int));
    cudaMalloc((void **)&d_numbofpairs, sizeof(int));

    cudaMemcpy(d_n, &n, sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_dim, &dim, sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_x, &x, sizeof(int), cudaMemcpyHostToDevice);

    for (z = 0; z <= max; z++)
    {
        cudaMemcpy(d_z, &z, sizeof(int), cudaMemcpyHostToDevice);
        coefficient << <1, THREADS >> >(d_x, *d_n, *d_dim, *d_z, d_numbofpairs);
        cudaMemcpy(&correlation_sum, d_numbofpairs, sizeof(int), cudaMemcpyDeviceToHost);
        printf("result for epsilon %d returns %d\n", z, correlation_sum);
    }
    cudaFree(d_n);
    cudaFree(d_dim);
    cudaFree(d_z);
    cudaFree(d_x);
    cudaFree(d_numbofpairs);
    return 0;
}

我需要一些帮助或提示要更改的内容、错误的内容以及为什么它总是崩溃，以便我可以修复它。谢谢！

编辑：我完成了一些部分，对不起我的错。至于threads和blocks，我很疑惑，GPU显示1024个threads per block，不知道是不是。

Answer 1

所以 "crash" 是一个段错误。段错误是 host 代码中的问题，而不是内核代码（尽管它可能在您对 CUDA API 的使用中）。

您的代码存在各种问题。

这可能会引起麻烦：
```
int x[40000]
```
这会创建一个基于堆栈的大型分配。相反，我建议进行动态分配：
```
int *x = (int *)malloc(40000*sizeof(int));
```
动态分配有更高的大小限制。
从您的内核使用情况可以很清楚地看出您打算使用整个 x 向量。因此，d_x 的设备分配不正确：
```
cudaMalloc((void **)&d_x, sizeof(int));
```
我们需要在设备上分配与主机上相同的大小：
```
cudaMalloc((void **)&d_x, 40000*sizeof(int));
```
对应于 2，您可能希望将整个 x 向量复制到设备（这不是很清楚，因为您的代码没有显示 [=20= 的初始化]), and 你在这里错误地获取了 x 的地址，但是 x 已经一个指针:
```
cudaMemcpy(d_x, &x, sizeof(int), cudaMemcpyHostToDevice);
```
所以我们想要这样的东西：
```
cudaMemcpy(d_x, x, 40000*sizeof(int), cudaMemcpyHostToDevice);
```
您的其他内核参数似乎是标量参数。您也大多不正确地处理这些问题：
```
__global__ void coefficient(int *vctr, int numberofpoints, int coefficient_row, int epsilon, int *numbofpairs){
```
对于像上面指定的 numberofpoints 这样的参数（单向传递给函数），我们只是 按值传递 我们在调用内核时想要的主机数量，就像我们使用普通的 C 函数一样。所以这个内核调用是不正确的（即使它看起来可以编译）：
```
coefficient << <1, THREADS >> >(d_x, *d_n, *d_dim, *d_z, d_numbofpairs);
```
相反，我们只想按值传递主机变量：
```
coefficient << <1, THREADS >> >(d_x, n, dim, z, d_numbofpairs);
```
因为 d_numbofpairs 是双向的，所以你的用法是正确的。
我还建议将 proper cuda error checking 添加到您的代码中。

这是一个修复了上述错误的完整示例。我认为结果当然是假的，因为输入数据（例如 x）未初始化。

$ cat t724.cu
#include <stdio.h>

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)


#define THREADS 1024

__global__ void coefficient(int *vctr, int numberofpoints, int coefficient_row, int epsilon, int *numbofpairs){
        int i = blockIdx.x * blockDim.x + threadIdx.x;
        int j = blockIdx.y * blockDim.y + threadIdx.y;
        int k = blockIdx.z * blockDim.z + threadIdx.z;
        int sum;
        numbofpairs = 0;
        int sq_epsilon = epsilon*epsilon;


        if (i <= numberofpoints - coefficient_row)
        {
            sum = 0;
            if (j <= numberofpoints + 1 - coefficient_row)
            {
                if (k < coefficient_row)
                    sum = sum + (vctr[i + k] - vctr[j + k])*(vctr[i + k] - vctr[j + k]);
                if (sum < sq_epsilon){
                    numbofpairs++;
                    sum = 0;
    }}}}

int main()
{
  int n, dim, max, z;
  int  *d_x, *d_numbofpairs;
  int correlation_sum = 0;
  int *x = (int *)malloc(40000*sizeof(int));
  if (x == NULL) {printf("malloc fail\n"); return -1;}
    n=10;
    max=10;
    dim=3;

    cudaMalloc((void **)&d_x, sizeof(int));
    cudaCheckErrors("cudaMalloc 1 fail");
    cudaMalloc((void **)&d_numbofpairs, sizeof(int));
    cudaCheckErrors("cudaMalloc 2 fail");
    cudaMemcpy(d_x, x, sizeof(int), cudaMemcpyHostToDevice);
    cudaCheckErrors("cudaMemcpy 1 fail");

    for (z = 0; z <= max; z++)
    {
        coefficient << <1, THREADS >> >(d_x, n, dim, z, d_numbofpairs);
        cudaMemcpy(&correlation_sum, d_numbofpairs, sizeof(int), cudaMemcpyDeviceToHost);
        cudaCheckErrors("cudaMemcpy 2/kernel fail");
        printf("result for epsilon %d returns %d\n", z, correlation_sum);
    }
    cudaFree(d_x);
    cudaFree(d_numbofpairs);
    return 0;
}
$ nvcc -o t724 t724.cu
$ ./t724
result for epsilon 0 returns 3
result for epsilon 1 returns 3
result for epsilon 2 returns 3
result for epsilon 3 returns 3
result for epsilon 4 returns 3
result for epsilon 5 returns 3
result for epsilon 6 returns 3
result for epsilon 7 returns 3
result for epsilon 8 returns 3
result for epsilon 9 returns 3
result for epsilon 10 returns 3
$

请注意，我没有对您的内核代码进行任何更改。

CUDA 三重循环

CUDA triple loop

c

for-loop

cuda