CUDA Samples: Dot Product

以下CUDA sample是分别用C++和CUDA实现的两个非常大的向量实现点积操作，并对其中使用到的CUDA函数进行了解说，各个文件内容如下：

common.hpp:

#ifndef FBC_CUDA_TEST_COMMON_HPP_
#define FBC_CUDA_TEST_COMMON_HPP_#include<random>template< typename T >
static inline int check_Cuda(T result, const char * const func, const char * const file, const int line)
{if (result) {fprintf(stderr, "Error CUDA: at %s: %d, error code=%d, func: %s\n", file, line, static_cast<unsigned int>(result), func);cudaDeviceReset(); // Make sure we call CUDA Device Reset before exitingreturn -1;}
}template< typename T >
static inline int check(T result, const char * const func, const char * const file, const int line)
{if (result) {fprintf(stderr, "Error: at %s: %d, error code=%d, func: %s\n", file, line, static_cast<unsigned int>(result), func);return -1;}
}#define checkCudaErrors(val) check_Cuda((val), __FUNCTION__, __FILE__, __LINE__)
#define checkErrors(val) check((val), __FUNCTION__, __FILE__, __LINE__)#define CHECK(x) { \if (x) {} \else { fprintf(stderr, "Check Failed: %s, file: %s, line: %d\n", #x, __FILE__, __LINE__); return -1; } \
}#define PRINT_ERROR_INFO(info) { \fprintf(stderr, "Error: %s, file: %s, func: %s, line: %d\n", #info, __FILE__, __FUNCTION__, __LINE__); \return -1; }#define EPS 1.0e-4 // ε(Epsilon),非常小的数static inline void generator_random_number(float* data, int length, float a = 0.f, float b = 1.f)
{std::random_device rd; std::mt19937 generator(rd()); // 每次产生不固定的不同的值//std::default_random_engine generator; // 每次产生固定的不同的值std::uniform_real_distribution<float> distribution(a, b);for (int i = 0; i < length; ++i) {data[i] = distribution(generator);}
}#endif // FBC_CUDA_TEST_COMMON_HPP_

funset.cpp:

#include "funset.hpp"
#include <random>
#include <iostream>
#include <vector>
#include <memory>
#include "common.hpp"int test_dot_product()
{const int length{ 10000000 };std::unique_ptr<float[]> A(new float[length]);std::unique_ptr<float[]> B(new float[length]);generator_random_number(A.get(), length, -10.f, 10.f);generator_random_number(B.get(), length, -10.f, 10.f);float elapsed_time1{ 0.f }, elapsed_time2{ 0.f }; // millisecondsfloat value1{ 0.f }, value2{ 0.f };int ret = dot_product_cpu(A.get(), B.get(), &value1, length, &elapsed_time1);if (ret != 0) PRINT_ERROR_INFO(long_vector_add_cpu);ret = dot_product_gpu(A.get(), B.get(), &value2, length, &elapsed_time2);if (ret != 0) PRINT_ERROR_INFO(matrix_mul_gpu);if (fabs(value1 - value2) > EPS) {fprintf(stderr, "Result verification failed value1: %f, value2: %f\n", value1, value2);}fprintf(stderr, "cpu run time: %f ms, gpu run time: %f ms\n", elapsed_time1, elapsed_time2);return 0;
}

dot_product.cpp:

#include "funset.hpp"
#include <chrono>int dot_product_cpu(const float* A, const float* B, float* value, int elements_num, float* elapsed_time)
{auto start = std::chrono::steady_clock::now();*value = 0.f;for (int i = 0; i < elements_num; ++i) {(*value) += A[i] * B[i];}auto end = std::chrono::steady_clock::now();auto duration = std::chrono::duration_cast<std::chrono::nanoseconds>(end - start);*elapsed_time = duration.count() * 1.0e-6;return 0;
}

dot_product.cu:

#include "funset.hpp"
#include <iostream>
#include <algorithm>
#include <memory>
#include <cuda_runtime.h> // For the CUDA runtime routines (prefixed with "cuda_")
#include <device_launch_parameters.h>
#include "common.hpp"/* __global__: 函数类型限定符;在设备上运行;在主机端调用,计算能力3.2及以上可以在
设备端调用;声明的函数的返回值必须是void类型;对此类型函数的调用是异步的,即在
设备完全完成它的运行之前就返回了;对此类型函数的调用必须指定执行配置,即用于在
设备上执行函数时的grid和block的维度,以及相关的流(即插入<<<   >>>运算符);
a kernel,表示此函数为内核函数(运行在GPU上的CUDA并行计算函数称为kernel(内核函
数),内核函数必须通过__global__函数类型限定符定义);*/
__global__ static void dot_product(const float* A, const float* B, float* partial_C, int elements_num)
{/* __shared__: 变量类型限定符；使用__shared__限定符，或者与__device__限定符连用，此时声明的变量位于block中的共享存储器空间中，与block具有相同的生命周期，仅可通过block内的所有线程访问；__shared__和__constant__变量默认为是静态存储；在__shared__前可以加extern关键字，但表示的是变量大小由执行参数确定；__shared__变量在声明时不能初始化；可以将CUDA C的关键字__shared__添加到变量声明中，这将使这个变量驻留在共享内存中；CUDA C编译器对共享内存中的变量与普通变量将分别采取不同的处理方式 */__shared__ float cache[256]; // == threadsPerBlock/* gridDim: 内置变量,用于描述线程网格的维度,对于所有线程块来说,这个变量是一个常数,用来保存线程格每一维的大小,即每个线程格中线程块的数量.一个grid最多只有二维,为dim3类型；blockDim: 内置变量,用于说明每个block的维度与尺寸.为dim3类型,包含了block在三个维度上的尺寸信息;对于所有线程块来说,这个变量是一个常数,保存的是线程块中每一维的线程数量;blockIdx: 内置变量,变量中包含的值就是当前执行设备代码的线程块的索引;用于说明当前thread所在的block在整个grid中的位置,blockIdx.x取值范围是[0,gridDim.x-1],blockIdx.y取值范围是[0, gridDim.y-1].为uint3类型,包含了一个block在grid中各个维度上的索引信息;threadIdx: 内置变量,变量中包含的值就是当前执行设备代码的线程索引;用于说明当前thread在block中的位置;如果线程是一维的可获取threadIdx.x,如果是二维的还可获取threadIdx.y,如果是三维的还可获取threadIdx.z;为uint3类型,包含了一个thread在block中各个维度的索引信息 */int tid = threadIdx.x + blockIdx.x * blockDim.x;int cacheIndex = threadIdx.x;float tmp{ 0.f };while (tid < elements_num) {tmp += A[tid] * B[tid];tid += blockDim.x * gridDim.x;}// 设置cache中相应位置上的值// 共享内存缓存中的偏移就等于线程索引；线程块索引与这个偏移无关，因为每// 个线程块都拥有该共享内存的私有副本cache[cacheIndex] = tmp;/* __syncthreads: 对线程块中的线程进行同步；CUDA架构将确保，除非线程块中的每个线程都执行了__syncthreads()，否则没有任何线程能执行__syncthreads()之后的指令;在同一个block中的线程通过共享存储器(shared memory)交换数据，并通过栅栏同步(可以在kernel函数中需要同步的位置调用__syncthreads()函数)保证线程间能够正确地共享数据；使用clock()函数计时，在内核函数中要测量的一段代码的开始和结束的位置分别调用一次clock()函数，并将结果记录下来。由于调用__syncthreads()函数后，一个block中的所有thread需要的时间是相同的，因此只需要记录每个block执行需要的时间就行了，而不需要记录每个thread的时间 */__syncthreads();// 对于规约运算来说，以下code要求threadPerBlock必须是2的指数int i = blockDim.x / 2;while (i != 0) {if (cacheIndex < i)cache[cacheIndex] += cache[cacheIndex + i];// 在循环迭代中更新了共享内存变量cache，并且在循环的下一次迭代开始之前，// 需要确保当前迭代中所有线程的更新操作都已经完成__syncthreads();i /= 2;}// 只有cacheIndex == 0的线程执行这个保存操作，这是因为只有一个值写入到// 全局内存，因此只需要一个线程来执行这个操作，当然你也可以选择任何一个// 线程将cache[0]写入到全局内存if (cacheIndex == 0)partial_C[blockIdx.x] = cache[0];
}int dot_product_gpu(const float* A, const float* B, float* value, int elements_num, float* elapsed_time)
{/* cudaEvent_t: CUDA event types,结构体类型, CUDA事件,用于测量GPU在某个任务上花费的时间,CUDA中的事件本质上是一个GPU时间戳,由于CUDA事件是在GPU上实现的,因此它们不适于对同时包含设备代码和主机代码的混合代码计时*/cudaEvent_t start, stop;// cudaEventCreate: 创建一个事件对象,异步启动cudaEventCreate(&start);cudaEventCreate(&stop);// cudaEventRecord: 记录一个事件,异步启动,start记录起始时间cudaEventRecord(start, 0);size_t lengthA{ elements_num * sizeof(float) }, lengthB{ elements_num * sizeof(float) };float *d_A{ nullptr }, *d_B{ nullptr }, *d_partial_C{ nullptr };// cudaMalloc: 在设备端分配内存cudaMalloc(&d_A, lengthA);cudaMalloc(&d_B, lengthB);/* cudaMemcpy: 在主机端和设备端拷贝数据,此函数第四个参数仅能是下面之一:(1). cudaMemcpyHostToHost: 拷贝数据从主机端到主机端(2). cudaMemcpyHostToDevice: 拷贝数据从主机端到设备端(3). cudaMemcpyDeviceToHost: 拷贝数据从设备端到主机端(4). cudaMemcpyDeviceToDevice: 拷贝数据从设备端到设备端(5). cudaMemcpyDefault: 从指针值自动推断拷贝数据方向,需要支持统一虚拟寻址(CUDA6.0及以上版本)cudaMemcpy函数对于主机是同步的 */cudaMemcpy(d_A, A, lengthA, cudaMemcpyHostToDevice);cudaMemcpy(d_B, B, lengthB, cudaMemcpyHostToDevice);const int threadsPerBlock{ 256 };const int blocksPerGrid = std::min(64, (elements_num + threadsPerBlock - 1) / threadsPerBlock);size_t lengthC{ blocksPerGrid * sizeof(float) };cudaMalloc(&d_partial_C, lengthC);/* <<< >>>: 为CUDA引入的运算符,指定线程网格和线程块维度等,传递执行参数给CUDA编译器和运行时系统,用于说明内核函数中的线程数量,以及线程是如何组织的;尖括号中这些参数并不是传递给设备代码的参数,而是告诉运行时如何启动设备代码,传递给设备代码本身的参数是放在圆括号中传递的,就像标准的函数调用一样;不同计算能力的设备对线程的总数和组织方式有不同的约束;必须先为kernel中用到的数组或变量分配好足够的空间,再调用kernel函数,否则在GPU计算时会发生错误,例如越界等;使用运行时API时,需要在调用的内核函数名与参数列表直接以<<<Dg,Db,Ns,S>>>的形式设置执行配置,其中：Dg是一个dim3型变量,用于设置grid的维度和各个维度上的尺寸.设置好Dg后,grid中将有Dg.x*Dg.y个block,Dg.z必须为1;Db是一个dim3型变量,用于设置block的维度和各个维度上的尺寸.设置好Db后,每个block中将有Db.x*Db.y*Db.z个thread;Ns是一个size_t型变量,指定各块为此调用动态分配的共享存储器大小,这些动态分配的存储器可供声明为外部数组(extern __shared__)的其他任何变量使用;Ns是一个可选参数,默认值为0;S为cudaStream_t类型,用于设置与内核函数关联的流.S是一个可选参数,默认值0. */dot_product << < blocksPerGrid, threadsPerBlock >> >(d_A, d_B, d_partial_C, elements_num);/* cudaDeviceSynchronize: kernel的启动是异步的, 为了定位它是否出错, 一般需要加上cudaDeviceSynchronize函数进行同步; 将会一直处于阻塞状态,直到前面所有请求的任务已经被全部执行完毕,如果前面执行的某个任务失败,将会返回一个错误；当程序中有多个流,并且流之间在某一点需要通信时,那就必须在这一点处加上同步的语句,即cudaDeviceSynchronize；异步启动reference: https://stackoverflow.com/questions/11888772/when-to-call-cudadevicesynchronize *///cudaDeviceSynchronize();std::unique_ptr<float[]> partial_C(new float[blocksPerGrid]);cudaMemcpy(partial_C.get(), d_partial_C, lengthC, cudaMemcpyDeviceToHost);*value = 0.f;for (int i = 0; i < blocksPerGrid; ++i) {(*value) += partial_C[i];}// cudaFree: 释放设备上由cudaMalloc函数分配的内存cudaFree(d_A);cudaFree(d_B);cudaFree(d_partial_C);// cudaEventRecord: 记录一个事件,异步启动,stop记录结束时间cudaEventRecord(stop, 0);// cudaEventSynchronize: 事件同步,等待一个事件完成,异步启动cudaEventSynchronize(stop);// cudaEventElapseTime: 计算两个事件之间经历的时间,单位为毫秒,异步启动cudaEventElapsedTime(elapsed_time, start, stop);// cudaEventDestroy: 销毁事件对象,异步启动cudaEventDestroy(start);cudaEventDestroy(stop);return 0;
}

GitHub： https://github.com/fengbingchun/CUDA_Test

CUDA Samples: Dot Product相关推荐

CUDA Samples: dot product(使用零拷贝内存)
以下CUDA sample是分别用C++和CUDA实现的点积运算code,CUDA包括普通实现和采用零拷贝内存实现两种,并对其中使用到的CUDA函数进行了解说,code参考了<GPU高性能编程C ...
CUDA Samples: 获取设备属性信息
通过调用CUDA的cudaGetDeviceProperties函数可以获得指定设备的相关信息,此函数会根据GPU显卡和CUDA版本的不同得到的结果也有所差异,下面code列出了经常用到的设备信息: ...
CUDA Samples: matrix multiplication(C = A * B)
以下CUDA sample是分别用C++和CUDA实现的两矩阵相乘运算code即C= A*B,CUDA中包含了两种核函数的实现方法,第一种方法来自于CUDA Samples\v8.0\0_Simple ...
CUDA Samples：Vector Add
以下CUDA sample是分别用C++和CUDA实现的两向量相加操作,参考CUDA 8.0中的sample:C:\ProgramData\NVIDIA Corporation\CUDA Sample ...
FB面经Prepare: Dot Product
Conduct Dot Product of two large Vectors 1. two pointers 2. hashmap 3. 如果没有额外空间,如果一个很大,一个很小,适合scan小的 ...
安装cuda后却没有CUDA Samples怎么办？
Cuda 11.6版本之后将不再编译cuda,所以必须自己从github下载后自行编译, 下载网址为 cuda samples 对于windows用户, Windows示例是使用Visual Stud ...
向量点积(Dot Product),向量叉积(Cross Product)
参考的是<游戏和图形学的3D数学入门教程>,非常不错的书,推荐阅读,老外很喜欢把一个东西解释的很详细. 1.向量点积(Dot Product) 向量点积的结果有什么意义?事实上,向量的点积 ...
dot product【点积】
(1)概念点积在数学中,又称数量积(dot product; scalar product),是指接受在实数R上的两个向量并返回一个实数值标量的二元运算. 两个向量a = [a1, a2,-, an ...
CUDA Samples 之 Simulations 之 Particles （1）
CUDA Samples 之 Simulations 之 Particles源码学习(1) 自己用C++编程做颗粒堆积,但效率很低,所以想将程序并行,所以开始接触CUDA.但是完全不知道如何搭一个并行 ...

CUDA Samples: Dot Product

CUDA Samples: Dot Product相关推荐

最新文章

热门文章