4、跨步循环

一、介绍

更新的内核设置了网格中线程总数（blockDim.x * gridDim.x）的步幅。 CUDA内核中的这种类型的循环通常称为网格跨步循环。

数据集比网格要大，要使用跨步循环

二、代码

#include <stdio.h>
#include <stdlib.h> //malloc头文件

void cpu(int *a, int N)
{
    for (int i = 0; i < N; i++)
    {
        a[i] = i;
    }
}

__global__ void gpu(int *a, int N)
{
    int threadi = blockIdx.x * blockDim.x + threadIdx.x; //计算出当前是第几个线程
    int stride = gridDim.x * blockDim.x; //当前所有线程块*每个线程块有多少线程=所有线程数
    //printf("当前是第threadi = %d个线程, 所有线程数stride = %d\n", threadi,stride);
    //printf("当前块索引blockIdx.x = %d,该块线程数blockDim.x= %d,当前块中索引threadIdx.x = %d,网格块数gridDim.x=%d\n", blockIdx.x,blockDim.x,threadIdx.x,gridDim.x);
    for (int i = threadi; i < N; i += stride)   
    {
        a[i] *= 2;
        //printf("%d\n", a[i]);
    }
}

bool check(int *a, int N) //测试GPU函数中数据是否正确执行
{
    for (int i = 0; i < N; i++)
    {
        if (a[i] != 2 * i)  return false;
    }
    return true;
}

int main()
{
    const int N = 500;
    size_t size = N * sizeof(int);  //需分配的size大小
    int *a;
    cudaMallocManaged(&a, size);    //分配malloc，既可以CPU使用，也可以GPU使用
    cpu(a, N);  //数组地址和大小
    
    size_t threads = 256;   //一个block256个线程
    size_t blocks = 1;
    gpu<<<blocks, threads>>>(a, N);
    cudaDeviceSynchronize();
    
    check(a, N) ? printf("OK!") : printf("error!");
    cudaFree(a);
}