Sep 26, 2018 原创文章
CUDA并行编程学习（2）-- 使用CUDA计算一个数组的距离值
代码下载：https://github.com/myurtoglu/cudaforengineers/tree/master/dist_v1_cuda
基于《CUDA高性能并行计算》中的例程
#include <stdio.h>
#define N 64
#define TPB 32
int main()
{
const float ref = 0.5f;
float *d_out = 0;
cudaMalloc(&d_out, N*sizeof(float));
distanceKernel<<<N/TPB, TPB>>>(d_out, ref, N);
cudaFree(d_out); // Free the memory
return 0;
}
Main函数
1、声明float 型指针 d_out
2、使用cudaMalloc 在设备端申请长度为N = 64 的float型数组 d_out
3、使用 核函数 distanceKernel 计算距离。
distanceKernel<<<N/TPB, TPB>>>(d_out, ref, N); 调用核函数，使用N/TPB个线程块，其中每个线程块包含TPB个线程。
4、使用cudaFree 释放设备端中的float型数组 d_out
__device__ float scale(int i, int n)
{
return ((float)i)/(n - 1);
}
1、标识符__device__，为设备端调用并设备端的执行的函数
__device__ float distance(float x1, float x2)
{
return sqrt((x2 - x1)*(x2 - x1));
}
__global__ void distanceKernel(float *d_out, float ref, int len)
{
const int i = blockIdx.x*blockDim.x + threadIdx.x;
const float x = scale(i, len);
d_out[i] = distance(x, ref);
printf("i = %2d: dist from %f to %f is %f.\n", i, ref, x, d_out[i]);
}
核函数 distanceKernel
1、标识符__global__，为核函数的标识符
2、使用核函数的内建线程索引和维度索引变量计算索引i
const int i = blockIdx.x*blockDim.x + threadIdx.x;
其中：blockIdx （uint3向量类型）为这个线程块在网格中的索引；blockDim（dim3向量类型）为每个线程块中的线程数目；threadIdx（uint3向量类型）为这个线程在线程块中的索引
