c++ - CUDA : stencil code not working

Question

首先:如果这个话题做得不好，我很抱歉(这是我的第一个话题)我目前正在尝试学习 NVIDIA 上的 GPU 计算，但是 我对 CUDA 的 __syncthreads() 方法有疑问，我认为它不起作用。我尝试在网上搜索，但没有找到解决方法。

__global__ void stencil_1d(int *in, int *out) {
    __shared__ int temp[BLOCK_SIZE + 2 * RADIUS];   // Création de la mémoire partagée avec tout les threads d'un même block

    int gindex = threadIdx.x + blockIdx.x * blockDim.x;
    int lindex = threadIdx.x + RADIUS;

    /*for (int i = 0; i < N*sizeof(int); i++)
        printf("%i ---i=%i \n", in[i], i);*/

    // Déplacer les éléments d'entrées sur la mémoire partagée
    temp[lindex] = in[gindex];
    if (threadIdx.x < RADIUS) {
        temp[lindex - RADIUS] = in[gindex - RADIUS]; // Récuprère le Halo avant les valeurs du block
        temp[lindex + BLOCK_SIZE] = in[gindex + BLOCK_SIZE]; // Récupère le Halo après les valeurs du block
    }

    __syncthreads(); // Synchronisation des Threads d'un même Block

    int result = 0;
    for (int offset = -RADIUS; offset <= RADIUS ; offset++)
        result += temp[lindex + offset];

    out[gindex] = result;
}

当我取消对 for 的注释时，程序正常工作，但目前没有 for 我的程序在 out 变量中返回 -842150451。

主要代码:

int main()
{
    int size = N * sizeof(int);

    /******************** Utilisation de la fonction GPU stencil_1d ********************/

    int *in, *out; // Variable sur la mémoire Host
    int *d_in, *d_out;  //Variable sur la mémoire Device

    // Allocation de la mémore aux variables sur le Device
    cudaMalloc((void**)&d_in, size);
    cudaMalloc((void**)&d_out, size);

    // Allocation de la mémoire aux variables de Host
    in = (int*)malloc(size); random_ints(in, N);
    out = (int*)malloc(size);

    // Copie des valeurs des variables de Host vers Device
    cudaMemcpy(d_in, in, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_out, out, size, cudaMemcpyHostToDevice);

    // Exécution de la fonction sur le Device (ici 3 Blocks, 10 Threads)
    stencil_1d <<<N/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_in, d_out);

    // Récupération de la variable out de Device vers Host
    cudaMemcpy(out, d_out, size, cudaMemcpyDeviceToHost);

    // Affichage du résultat
    for(int i=0; i<size; i++)
        printf("%i ---i=%i \n", out[i], i);

    // Libération de la mémoire prise par les variables sur Host
    free(in); free(out);

    // Libération de la mémoire prise par les variables sur le Device
    cudaFree(d_in); cudaFree(d_out);

    return 0;
}

如果忘记了:定义:

#define N 30
#define THREADS_PER_BLOCK 10
#define BLOCK_SIZE (N/THREADS_PER_BLOCK)
#define RADIUS 3

random_ints 代码:

void random_ints(int *var, int n) // Attribue une valeur à toutes le composantes des variables
{
    int i;
    for (i = 0; i < n; i++)
        var[i] = 1;
}

预先感谢您的回答。

Answer 1

This code最初是为教学而设计的；它有一些缺陷。

首先，任何时候您在使用 CUDA 代码时遇到问题，我建议 proper CUDA error checking并使用 cuda-memcheck 运行您的代码(参见下面我使用 cuda-memcheck 的示例)。如果您在向其他人寻求帮助之前这样做，并在您的问题中提供错误信息，它可能会帮助其他人帮助您。

如果您使用 cuda-memcheck 运行此代码，它将指示访问全局内存和共享内存的错误。

1. 您选择的 BLOCK_SIZE 不正确。这应该设置为等于 THREADS_PER_BLOCK，而不是 (N/THREADS_PER_BLOCK)。看来您打算使用 3 个 block (每个 block 10 个线程)来运行此内核，因此我们将使用它。

2. 这些行将越界索引:

   if (threadIdx.x < RADIUS) {
       temp[lindex - RADIUS] = in[gindex - RADIUS]; // Récuprère le Halo avant les valeurs du block
       temp[lindex + BLOCK_SIZE] = in[gindex + BLOCK_SIZE]; // Récupère le Halo après les valeurs du block
   

   例如，在第一个 block 中，第一个线程(threadIdx.x为0)，gindex为0，所以计算gindex - RADIUS 将为 -3。这不可能是正确的。

3. 这个 for 循环是不正确的:

   for(int i=0; i<size; i++)
   

   应该是:

   for(int i=0; i<N; i++)
   

当我修复这些问题时，您的代码运行没有错误并为我产生了一个合理的结果:

$ cat t280.cu
#define N 30
#define THREADS_PER_BLOCK 10
#define BLOCK_SIZE THREADS_PER_BLOCK
#define RADIUS 3

#include <stdio.h>

void random_ints(int *var, int n) // Attribue une valeur à toutes le composantes des variables
{
    int i;
    for (i = 0; i < n; i++)
        var[i] = 1;
}

__global__ void stencil_1d(int *in, int *out) {
    __shared__ int temp[BLOCK_SIZE + 2 * RADIUS];   // Création de la mémoire partagée avec tout les threads d'un même block

    int gindex = threadIdx.x + blockIdx.x * blockDim.x;
    int lindex = threadIdx.x + RADIUS;

    /*for (int i = 0; i < N*sizeof(int); i++)
        printf("%i ---i=%i \n", in[i], i);*/

    // Déplacer les éléments d'entrées sur la mémoire partagée
    temp[lindex] = in[gindex];
    if (threadIdx.x < RADIUS) {
        temp[lindex - RADIUS] = (gindex >= RADIUS)?in[gindex - RADIUS]:0; // Récuprère le Halo avant les valeurs du block
        temp[lindex + BLOCK_SIZE] = ((gindex + BLOCK_SIZE)<N)?in[gindex + BLOCK_SIZE]:0; // Récupère le Halo après les valeurs du block
    }

    __syncthreads(); // Synchronisation des Threads d'un même Block

    int result = 0;
    for (int offset = -RADIUS; offset <= RADIUS ; offset++)
        result += temp[lindex + offset];

    out[gindex] = result;
}

int main()
{
    int size = N * sizeof(int);

    /******************** Utilisation de la fonction GPU stencil_1d ********************/

    int *in, *out; // Variable sur la mémoire Host
    int *d_in, *d_out;  //Variable sur la mémoire Device

    // Allocation de la mémore aux variables sur le Device
    cudaMalloc((void**)&d_in, size);
    cudaMalloc((void**)&d_out, size);

    // Allocation de la mémoire aux variables de Host
    in = (int*)malloc(size); random_ints(in, N);
    out = (int*)malloc(size);

    // Copie des valeurs des variables de Host vers Device
    cudaMemcpy(d_in, in, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_out, out, size, cudaMemcpyHostToDevice);

    // Exécution de la fonction sur le Device (ici 3 Blocks, 10 Threads)
    stencil_1d <<<N/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_in, d_out);

    // Récupération de la variable out de Device vers Host
    cudaMemcpy(out, d_out, size, cudaMemcpyDeviceToHost);

    // Affichage du résultat
    for(int i=0; i<N; i++)
        printf("%i ---i=%i \n", out[i], i);

    // Libération de la mémoire prise par les variables sur Host
    free(in); free(out);

    // Libération de la mémoire prise par les variables sur le Device
    cudaFree(d_in); cudaFree(d_out);

    return 0;
}
$ nvcc -o t280 t280.cu
$ cuda-memcheck ./t280
========= CUDA-MEMCHECK
4 ---i=0
5 ---i=1
6 ---i=2
7 ---i=3
7 ---i=4
7 ---i=5
7 ---i=6
7 ---i=7
7 ---i=8
7 ---i=9
7 ---i=10
7 ---i=11
7 ---i=12
7 ---i=13
7 ---i=14
7 ---i=15
7 ---i=16
7 ---i=17
7 ---i=18
7 ---i=19
7 ---i=20
7 ---i=21
7 ---i=22
7 ---i=23
7 ---i=24
7 ---i=25
7 ---i=26
6 ---i=27
5 ---i=28
4 ---i=29
========= ERROR SUMMARY: 0 errors
$

我们在模板输出的每一端得到 4、5、6 的原因是我们为上面的第 2 项在内核中设置的限制，以避免越界索引。如果需要，您可以更改此边界行为。

还有一条评论:现在您的代码选择了 N 和 THREADS_PER_BLOCK 以便它可以被均分。只要您这样做(并遵守其他限制，例如每 block 1024 的最大线程数)，您应该可以使用此代码。为了充分的灵 active other changes should be made ，但是我在这里描述的应该足以让您克服这些错误。