CUDA编程练习（三）向量内积

Simple_Bright&Man

已于 2024-06-04 20:55:12 修改

阅读量436

点赞数 10

分类专栏： CUDA编程文章标签：算法 python c++

于 2024-05-28 22:28:55 首次发布

本文链接：https://blog.youkuaiyun.com/SimpleBright_Man/article/details/139279382

版权

1 想法

相乘没什么好说的，取值是平行的，a取i，b也取i

不行，脑子开始内讧了。
有以下规约想法：

元素太多，每个线程在数组上跨grid规模取值求和，一口气就能把元素数量；
线程足够多，直接block内折半向下归约，公共内存里面会出现block数量的元素（或者放在共享内存开辟的第一个值的位置）；
原子加法，或者压到单block内接着归约；
考虑线程束，线程束规约一次性压缩32倍（？？？有用么这，不还是SM一波一波的一次性压缩一半，记得复习）

代码写满了注释，要理解应该是只考验耐心:>

2 代码

#include <iostream>
#include <assert.h>
#include <cstdlib>
#include <cmath>

#define VECTOR_LENGTH 100000000
#define THREAD_PER_BLOCK 256
#define BLOCK_NUM 64
#define MAX_ERR 1e-4
#define CUDA_CALL(x) do { if((x) != cudaSuccess) { printf("Error at %s:%d\n",__FILE__,__LINE__); return EXIT_FAILURE;}} while(0)

void cuda_error_check(const char* prefix, const char * postfix) {
    if (cudaPeekAtLastError()!=cudaSuccess) {
        printf("\n%s%s%s", prefix, cudaGetErrorString(cudaGetLastError()), postfix);
        cudaDeviceReset();
    }
}

__global__ void dummy_vectorDotProduct(float *out, float *a, float *b, int n) 
{
    float sum=0;  // 寄存器
    for(int i = 0; i < n; i++)
    {
        sum = sum +  a[i] * b[i];
    }
    *out = sum;
}

__global__ void sharedAtomic_vectorDotProduct(float *out, float *a, float *b, int n) // 考虑到共享内存，一个block一个结果
{
    int leap= gridDim.x*blockDim.x;
    int bid = blockIdx.x*blockDim.x+threadIdx.x;
    int tid = threadIdx.x;
    
    __shared__ float ssTmp[THREAD_PER_BLOCK];  // 共存
    ssTmp[tid] = 0;
    // 同步修改
    __syncthreads();
    
    for(int i = bid; i < n; i+=leap)
    {
        ssTmp[tid] += a[i] * b[i];  // 全部加到一个grid里面，
        // __syncthreads();  // 猜猜这个东西有多重要？加了就会导致GPU原地爆炸？BUT WHY？？