一、介绍
更新的内核设置了网格中线程总数(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);
}