CUDA笔记3

这篇具有很好参考价值的文章主要介绍了CUDA笔记3。希望对大家有所帮助。如果存在错误或未考虑完全的地方,请大家不吝赐教,您也可以点击"举报违法"按钮提交疑问。

CUDA编程-入门、加法、乘法

例子1:入门

#include <stdio.h>

__global__ void hello_from_gpu()
{
    const int bid = blockIdx.x;
    const int tid = threadIdx.x;
    printf("Hello world from block %d and thread %d!\n",bid,tid);
}
int main(void)
{
    hello_from_gpu<<<2,3>>>();
    cudaDeviceSynchronize();
    return 0;
}
Hello world from block 1 and thread 0!
Hello world from block 1 and thread 1!
Hello world from block 1 and thread 2!
Hello world from block 0 and thread 0!
Hello world from block 0 and thread 1!
Hello world from block 0 and thread 2!

每个线程执行一次“Hello world”

例子2:数组相加

#include <math.h>
#include <stdio.h>

void __global__ add(const double *x, const double *y, double *z, int count)
{
    const int n = blockDim.x * blockIdx.x + threadIdx.x;
	if( n < count)
	{
        printf("Block ID: %d, blockDim.x: %d,Thread ID: %d, n: %d\n",
               blockIdx.x, blockDim.x, threadIdx.x, n);
	    z[n] = x[n] + y[n];
        printf("z[%d]: %f\n", n, z[n]);
	}
    // 第二种xy输出
    // Add a barrier to ensure all threads have printed before exiting the kernel
    // __syncthreads();

    // // Only have one thread (threadIdx.x == 0) print the entire arrays
    // if (threadIdx.x == 0 && blockIdx.x == 0)
    // {
    //     for (int i = 0; i < count; ++i)
    //     {
    //         printf("Final Values - x[%d]: %f, y[%d]: %f, z[%d]: %f\n", i, x[i], i, y[i], i, z[i]);
    //     }
    // }

}
// void check(const double *z, const int N)
// {
//     bool error = false;
//     for (int n = 0; n < N; ++n)
//     {
//         if (fabs(z[n] - 3) > (1.0e-10))
//         {
//             error = true;
//         }
//     }
//     printf("%s\n", error ? "Errors" : "Pass");
// }
void check(const double *z, const double *h_x, const double *h_y, const int N)
{
    bool error = false;
    for (int n = 0; n < N; ++n)
    {
        // 修改此处的期望值计算方式
        // 不会执行,即error=false,返回Pass
        if (fabs(z[n] - (h_x[n] + h_y[n])) > (1.0e-10))
        {
            error = true;
        }
    }
    printf("%s\n", error ? "Errors" : "Pass");
}

int main(void)
{
    const int N = 6;
    const int M = sizeof(double) * N;
    double *h_x = (double*) malloc(M);
    double *h_y = (double*) malloc(M);
    double *h_z = (double*) malloc(M);

    for (int n = 0; n < N; ++n)
    {
        h_x[n] = 1;
        h_y[n] = n;
        //第一种xy输出
        printf("h_x[%d]: %f, h_y[%d]: %f\n", n, h_x[n], n, h_y[n]);
    }

    double *d_x, *d_y, *d_z;
    cudaMalloc((void **)&d_x, M);
    cudaMalloc((void **)&d_y, M);
    cudaMalloc((void **)&d_z, M);
    cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);
    cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice);

    const int block_size = 2;
    /*N是数组大小
    如果N能够被block_size整除,那么grid_size就等于N/block_size,
    那么每个线程块处理一个数据。
    如果N不能被block_size整除,那么(N+block_size-1)/block_size
    的计算会确保生成的grid_size足够大,
    保所有的输入数据都能够被覆盖到,并且没有数据被漏掉
    如果 N = 10,block_size = 3,则 (N + block_size - 1) / block_size = 4。
    这确保了有足够的线程块来处理所有的10个数据元素。*/
    const int grid_size = (N + block_size - 1) / block_size;
    add<<<grid_size, block_size>>>(d_x, d_y, d_z,N);

    cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);
    // check(h_z, N);
    check(h_z, h_x, h_y, N);

    free(h_x);
    free(h_y);
    free(h_z);
    cudaFree(d_x);
    cudaFree(d_y);
    cudaFree(d_z);
    return 0;
}
// xy值
h_x[0]: 1.000000, h_y[0]: 0.000000
h_x[1]: 1.000000, h_y[1]: 1.000000
h_x[2]: 1.000000, h_y[2]: 2.000000
h_x[3]: 1.000000, h_y[3]: 3.000000
h_x[4]: 1.000000, h_y[4]: 4.000000
h_x[5]: 1.000000, h_y[5]: 5.000000
//线程块,线程,n的值
Block ID: 1, blockDim.x: 2, Thread ID: 0, n: 2
Block ID: 1, blockDim.x: 2, Thread ID: 1, n: 3
Block ID: 2, blockDim.x: 2, Thread ID: 0, n: 4
Block ID: 2, blockDim.x: 2, Thread ID: 1, n: 5
Block ID: 0, blockDim.x: 2, Thread ID: 0, n: 0
Block ID: 0, blockDim.x: 2, Thread ID: 1, n: 1
//输出的z值
z[2]: 3.000000
z[3]: 4.000000
z[4]: 5.000000
z[5]: 6.000000
z[0]: 1.000000
z[1]: 2.000000

Block ID: 0, blockDim.x: 2, Thread ID: 0, n: 0,即第1线程块,第1个线程,根据以下公式:

  const int n = blockDim.x * blockIdx.x + threadIdx.x; 
  z[n] = x[n] + y[n];

算出n=0,进一步,z[0]=x[0]+y[0]=1+0=1
即第一步算出索引,第二步带入索引值求出z,其余同理。
请注意,由于线程和块的索引计算,每个线程处理不同的数据元素,从而实现了并行计算。在实际的大规模并行计算中,这种并行性能够带来显著的加速。
如Cuda笔记1所讲的,所有线程执行相同的核函数,且并行执行。
每个线程负责处理一个计算任务。在GPU中,大量的线程可以同时执行相同的指令,每个线程独立地处理不同的数据或任务,这被称为SIMD(Single Instruction, Multiple Data)架构。(分治思想,将大的任务分成多个一样的小任务,每个thread处理一个计算任务)
例如,假设有一个大数组需要排序,可以将数组分割成若干段,每个线程负责对其中一段进行排序。这样,每个线程独立执行排序任务,最后再将各段有序的数组合并成整体有序的数组。

通过规约(Reduction)操作理解GPU内存体系(源于知乎-紫气东来)

reduce 算法可以描述为
CUDA笔记3,CUDA,笔记,算法
其中 ×○可表示为求 sum,min,max,avg 等操作,最后获得的输出相比于输入一般维度上会递减。
在GPU中,reduce采用了一种树形的计算方式,例如下面的求和问题:
CUDA笔记3,CUDA,笔记,算法
由于GPU没有针对global数据的同步操作,只能针对block的数据进行同步。所以,一般而言将reduce分为两个阶段,其示意图如下:
CUDA笔记3,CUDA,笔记,算法
假设给定一个长度为N的数组,需要计算该数组的所有元素之和。首先需要将数组分为m个小份。而后,在第一阶段中,开启m个block计算出m个小份的reduce值。最后,在第二阶段中,使用一个block将m个小份再次进行reduce,得到最终的结果。
chatgpt的回答:
在GPU编程中,Reduce(归约)是一种并行计算模式,用于将一个大型数据集的元素合并为一个单一的值。Reduce操作通常用于计算总和、平均值、最大值、最小值等。在GPU中,由于存在大量的并行处理单元(CUDA核心),可以通过归约操作充分利用并行性,加速数据处理。

典型的GPU中的Reduce流程如下:

数据分割: 将大型数据集划分成小块,每个块由一个线程块(Thread Block)处理。

并行计算: 每个线程块中的线程并行地对其分配的数据块执行归约操作。这可以是对数据进行相加、取最大值、取最小值等。

层级归约: 由于线程块的数量通常大于硬件所支持的线程数,需要进行多层级的归约。每一层级的归约将当前层级的中间结果传递到下一层级,直到最终得到一个单一的结果。

最终结果: 最后一层的归约结果即为整个数据集的最终结果。

CUDA编程中,可以使用**共享内存(shared memory)**来加速归约操作,因为共享内存是线程块内的线程共享的。这减少了对全局内存的访问次数,提高了性能。

归约操作的实现可能因问题的不同而有所差异,但其核心思想是通过并行计算和层级归约来充分利用GPU的并行计算能力。

例子3:矩阵乘法(全局内存vs共享内存)

1)全局内存

#include <stdio.h>
#include <math.h>
#include "error.cuh"

#define BLOCK_SIZE 16


__global__ void gpu_matrix_mult(int *a,int *b, int *c, int m, int n, int k)
{ 
    //不是thread的一维索引,而是复合线程索引值,更符合矩阵运算等
    //索引就是变值,不要想复杂了,就相当于一维数组,二维数组中的索引一回事
    //一个索引对应一个线程(一个线程的位置,线程是最小单位),然后用这个线程计算,GPU只是很多线程同时计算而已。
    //计算当前线程在矩阵中的行索引。
    int row = blockIdx.y * blockDim.y + threadIdx.y; 
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int sum = 0;
    //确保当前线程的索引在矩阵的有效范围内,以避免越界访问。
    if( col < k && row < m) 
    {
        //对于每一列(或每一行,具体取决于矩阵乘法的顺序),
        //将对应位置的元素相乘并累加到 sum 中。
         //读取m的一行,n的一列
        for(int i = 0; i < n; i++) 
        {
            /*在实际的 GPU 编程中,选择线程块和网格的配置通常取决于算法和数据的特性,
            以及对硬件架构的优化。
            这儿矩阵相乘,因此选择二维的线程块和网格,线程块的 y 索引和 x 索引
            可以自然地映射到结果矩阵中的行和列。
            */ 
            sum += a[row * n + i] * b[i * k + col];
            //即每个(row,col)位置的线程处理一个计算任务
            /*
            与CPU的对比,
            void cpu_matrix_mult(int *h_a, int *h_b, int *h_result, int m, int n, int k) {
              for (int i = 0; i < m; ++i) 
              {
                 for (int j = 0; j < k; ++j) 
                 {
                    int tmp = 0.0;
                    for (int h = 0; h < n; ++h) 
                    {
                        tmp += h_a[i * n + h] * h_b[h * k + j];
                    }
                    h_result[i * k + j] = tmp;
                }
            }
        }
            事实上,外面的两个for循环用线程的动态索引取代了,而线程不是按顺序执行的
            而是多个线程同时执行的
            cpu中的i和j分别相当于GPU中的row,col 
            */
        }
        //row,col确定当前线程在输出矩阵 C 中的二维索引
        //用的一维表示二维a[i*M+j] i是行下标,M是一行的元素个数,j是列下标
        c[row * k + col] = sum;
    }
} 

__global__ void gpu_matrix_mult_shared(int *d_a, int *d_b, int *d_result, int n) 
{
    __shared__ int tile_a[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ int tile_b[BLOCK_SIZE][BLOCK_SIZE];

    int row = blockIdx.y * BLOCK_SIZE + threadIdx.y;
    int col = blockIdx.x * BLOCK_SIZE + threadIdx.x;
    int tmp = 0;
    int idx;

    for (int sub = 0; sub < gridDim.x; ++sub) 
    {
        idx = row * n + sub * BLOCK_SIZE + threadIdx.x;
        tile_a[threadIdx.y][threadIdx.x] = row<n && (sub * BLOCK_SIZE + threadIdx.x)<n? d_a[idx]:0;
        idx = (sub * BLOCK_SIZE + threadIdx.y) * n + col;
        tile_b[threadIdx.y][threadIdx.x] = col<n && (sub * BLOCK_SIZE + threadIdx.y)<n? d_b[idx]:0;

        __syncthreads();
        for (int k = 0; k < BLOCK_SIZE; ++k) 
        {
            tmp += tile_a[threadIdx.y][k] * tile_b[k][threadIdx.x];
        }
        __syncthreads();
    }
    if(row < n && col < n)
    {
        d_result[row * n + col] = tmp;
    }
}

void cpu_matrix_mult(int *h_a, int *h_b, int *h_result, int m, int n, int k) {
    for (int i = 0; i < m; ++i) 
    {
        for (int j = 0; j < k; ++j) 
        {
            int tmp = 0.0;
            for (int h = 0; h < n; ++h) 
            {
                tmp += h_a[i * n + h] * h_b[h * k + j];
            }
            h_result[i * k + j] = tmp;
        }
    }
}

int main(int argc, char const *argv[])
{
    int m=1000;
    int n=1000;
    int k=1000;

    int *h_a, *h_b, *h_c, *h_cc, *h_cs;
    //void **) &h_a将分配的内存地址赋值给指针h_a 
    //sizeof(int)*m*n)分配一个 int 类型的数组,其大小为 m*n 个元素
    //即主机端分配一个 int 类型的二维数组,大小为 m 行、n 列。
    CHECK(cudaMallocHost((void **) &h_a, sizeof(int)*m*n));
    CHECK(cudaMallocHost((void **) &h_b, sizeof(int)*n*k));
    CHECK(cudaMallocHost((void **) &h_c, sizeof(int)*m*k));
    CHECK(cudaMallocHost((void **) &h_cc, sizeof(int)*m*k));
    CHECK(cudaMallocHost((void **) &h_cs, sizeof(int)*m*k));
    

    /*CUDA事件
    cudaEvent_t: 这是CUDA库中表示事件的数据类型。
    CUDA事件通常用于测量 GPU 执行时间,用于确定 GPU 执行某个任务所花费的时间。
    cudaEventCreate表示用于创建 CUDA 事件
    原型“cudaError_t cudaEventCreate(cudaEvent_t* event);”
    
    */
    cudaEvent_t start, stop,stop_share;
    CHECK(cudaEventCreate(&start));
    CHECK(cudaEventCreate(&stop));
    CHECK(cudaEventCreate(&stop_share));

    //将整个二维矩阵h_a赋值为1
    //数组是在内存按顺序存储,因此二维其实也是一维
    for (int i = 0; i < m; ++i) {
        for (int j = 0; j < n; ++j) {
            h_a[i * n + j] = 1;
        }
    }

    for (int i = 0; i < n; ++i) {
        for (int j = 0; j < k; ++j) {
            h_b[i * k + j] = 0;
        }
    }

    int *d_a, *d_b, *d_c, *d_c_share;
    CHECK(cudaMalloc((void **) &d_a, sizeof(int)*m*n));
    CHECK(cudaMalloc((void **) &d_b, sizeof(int)*n*k));
    CHECK(cudaMalloc((void **) &d_c, sizeof(int)*m*k));
    CHECK(cudaMalloc((void **) &d_c_share, sizeof(int)*m*k));
    /*  cudaDeviceSynchronize()目的是确保设备上的所有操作都已完成。
    这段程序片段的目的是确保CUDA分配的设备内存已经成功完成。*/
    cudaDeviceSynchronize(); 
    // cudaEventRecord是CUDA中记录事件的函数,用于将当前时间戳记录到由start表示的事件中。
    CHECK(cudaEventRecord(start));
    // copy matrix A and B from host to device memory(cudaMemcpy全局内存)
    CHECK(cudaMemcpy(d_a, h_a, sizeof(int)*m*n, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_b, h_b, sizeof(int)*n*k, cudaMemcpyHostToDevice));

    unsigned int grid_rows = (m + BLOCK_SIZE - 1) / BLOCK_SIZE;
    unsigned int grid_cols = (k + BLOCK_SIZE - 1) / BLOCK_SIZE;
    /*
    dim3三维网格和线程块,z=1,blockidx.z=0,threadidx.z=0
    创建了一个二维的线程块与线程网格
    */
    dim3 dimGrid(grid_cols, grid_rows);
    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
    

    gpu_matrix_mult<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m,n,k);    

    CHECK(cudaMemcpy(h_c, d_c, (sizeof(int)*m*k), cudaMemcpyDeviceToHost));
    //cudaThreadSynchronize();
    CHECK(cudaEventRecord(stop));
    CHECK(cudaEventSynchronize(stop));
    
    gpu_matrix_mult_shared<<<dimGrid, dimBlock>>>(d_a, d_b, d_c_share, n);
    CHECK(cudaMemcpy(h_cs, d_c_share, (sizeof(int)*m*k), cudaMemcpyDeviceToHost));
    
    CHECK(cudaEventRecord(stop_share));
    CHECK(cudaEventSynchronize(stop_share));
    
    float elapsed_time, elapsed_time_share;
    CHECK(cudaEventElapsedTime(&elapsed_time, start, stop));
    CHECK(cudaEventElapsedTime(&elapsed_time_share, stop, stop_share));
    printf("Time_global = %g ms.\n", elapsed_time);
    printf("Time_share = %g ms.\n", elapsed_time_share);

    CHECK(cudaEventDestroy(start));
    CHECK(cudaEventDestroy(stop));    

    //cpu_matrix_mult(h_a, h_b, h_c, m, n, k);

    int ok = 1;
    /*for (int i = 0; i < m; ++i)
    { 
        for (int j = 0; j < k; ++j)
        {
            if(fabs(h_cs[i*k + j] - 0)>(1.0e-10))
            {
                printf("hcs: %d hc: %d  ",h_cs[i*k + j], h_c[i*k + j]);
                ok = 0;
            }
        }
    }

    if(ok)
    {
        printf("Pass!!!\n");
    }
    else
    {
        printf("Error!!!\n");
    }
    */
    // free memory
    CHECK(cudaFree(d_a));
    CHECK(cudaFree(d_b));
    CHECK(cudaFree(d_c));
    CHECK(cudaFreeHost(h_a));
    CHECK(cudaFreeHost(h_b));
    CHECK(cudaFreeHost(h_c));
    return 0;
}

重点:
①在实际的 GPU 编程中,选择线程块和网格的配置通常取决于算法和数据的特性,以及对硬件架构的优化。这儿矩阵相乘,因此选择二维的线程块和网格,线程块的 y 索引和 x 索引可以自然地映射到结果矩阵中的行和列。
②即每个(row,col)位置的线程处理一个计算任务。事实上,CPU外面的两个for循环用GPU线程的动态索引取代了,而线程不是按顺序执行的。而是多个线程同时执行的。CPU中的i和j分别相当于GPU中的row,col
③定义MK个thread,刚好动态映射_a,d_b的mn、nk,以及输出结果d_c中的mk个值
④copy matrix A and B from host to device memory(cudaMemcpy全局内存)
⑤逻辑上是二维的,实际内存物理上是一维的,因此要用一维索引形式表示

2)共享内存

理论:
CUDA笔记3,CUDA,笔记,算法
整体思路:
申请很多线程,每个线程求取pd中的每个值,就需要从M和N读取一行和一列的值。那么要求出pdsut中每个值,就要读取m和n中by=1行,bx=1中的所有值。即每一行需要多很多次,放到共享内存。但是共享内存小,因此分块。做子矩阵(蓝色)线程。
1、在共享内存申请空间
2、pd相当于一个block,这里面的每个thread只读取一个元素,放入share memory中,所有线程同时行动
3、计算时,每个线程读取share memory读取行的元素(蓝色块中的黄色行),结果存在temp中。然后黄色矩阵的行,结果放在temp中。计算完后,放到pd中。

#include <stdio.h>
#include <math.h>
#include "error.cuh"

#define BLOCK_SIZE 16


__global__ void gpu_matrix_mult(int *a,int *b, int *c, int m, int n, int k)
{ 
    //不是thread的一维索引,而是复合线程索引值,更符合矩阵运算等
    //索引就是变值,不要想复杂了,就相当于一维数组,二维数组中的索引一回事
    //一个索引对应一个线程(一个线程的位置,线程是最小单位),然后用这个线程计算,GPU只是很多线程同时计算而已。
    //计算当前线程在矩阵中的行索引。
    int row = blockIdx.y * blockDim.y + threadIdx.y; 
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int sum = 0;
    //确保当前线程的索引在矩阵的有效范围内,以避免越界访问。
    if( col < k && row < m) 
    {
        //对于每一列(或每一行,具体取决于矩阵乘法的顺序),
        //将对应位置的元素相乘并累加到 sum 中。
         //读取m的一行,n的一列
        for(int i = 0; i < n; i++) 
        {
            /*在实际的 GPU 编程中,选择线程块和网格的配置通常取决于算法和数据的特性,
            以及对硬件架构的优化。
            这儿矩阵相乘,因此选择二维的线程块和网格,线程块的 y 索引和 x 索引
            可以自然地映射到结果矩阵中的行和列。
            */ 
            sum += a[row * n + i] * b[i * k + col];
            //即每个(row,col)位置的线程处理一个计算任务
            /*
            与CPU的对比,
            void cpu_matrix_mult(int *h_a, int *h_b, int *h_result, int m, int n, int k) {
              for (int i = 0; i < m; ++i) 
              {
                 for (int j = 0; j < k; ++j) 
                 {
                    int tmp = 0.0;
                    for (int h = 0; h < n; ++h) 
                    {
                        tmp += h_a[i * n + h] * h_b[h * k + j];
                    }
                    h_result[i * k + j] = tmp;
                }
            }
        }
            事实上,外面的两个for循环用线程的动态索引取代了,而线程不是按顺序执行的
            而是多个线程同时执行的
            cpu中的i和j分别相当于GPU中的row,col 
            */
        }
        //row,col确定当前线程在输出矩阵 C 中的二维索引
        //用的一维表示二维a[i*M+j] i是行下标,M是一行的元素个数,j是列下标
        c[row * k + col] = sum;
    }
} 

__global__ void gpu_matrix_mult_shared(int *d_a, int *d_b, int *d_result, int M, int N, int K) 
{
    /*
    __shared__: 这是一个CUDA关键字,
    表示声明的变量是共享内存中的变量。共享内存是一块对于线程块中的所有线程可见的内存。
    int tile_a[BLOCK_SIZE][BLOCK_SIZE];这是一个二维整数数组,
    表示在共享内存中分配的二维数组。tile_a 被用于存储一个矩阵的子矩阵。
    怎么理解被用于存储一个矩阵的子矩阵?
    矩阵乘法通常涉及到将大矩阵分解为子矩阵,
    然后每个线程块负责计算一个或多个子矩阵的乘法。
    在矩阵乘法的上下文中,tile_a 存储了一个矩阵的子矩阵。当每个线程块负责计算某个子矩阵的乘法时,
    将这个子矩阵加载到 tile_a 中,以便线程块中的每个线程可以快速访问和共享这部分数据。
    
    */
    __shared__ int tile_a[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ int tile_b[BLOCK_SIZE][BLOCK_SIZE];

    int row = blockIdx.y * BLOCK_SIZE + threadIdx.y;
    int col = blockIdx.x * BLOCK_SIZE + threadIdx.x;
    int tmp = 0;
    int idx;
    //读取每个矩阵小块
    // for (int sub = 0; sub < gridDim.x; ++sub) 
    for (int sub = 0; sub <= N/BLOCK_SIZE; ++sub) 
    {
        // idx = row * n + sub * BLOCK_SIZE + threadIdx.x;
        // tile_a[threadIdx.y][threadIdx.x] = row<n && (sub * BLOCK_SIZE + threadIdx.x)<n? d_a[idx]:0;
        // idx = (sub * BLOCK_SIZE + threadIdx.y) * n + col;
        // tile_b[threadIdx.y][threadIdx.x] = col<n && (sub * BLOCK_SIZE + threadIdx.y)<n? d_b[idx]:0;
        //详细版
        int r = row;
        int c = sub * BLOCK_SIZE + threadIdx.x;
        idx = r * N + c;

        if (r >= M || c >= N)
        {
            tile_a[threadIdx.y][threadIdx.x]=0;
        }
        else
        {
            tile_a[threadIdx.y][threadIdx.x]=d_a[idx];
        }

        r = sub * BLOCK_SIZE + threadIdx.y;
        c = col;
        idx = r * K + c;
        if (c >= K || c >= N)
        {
            tile_b[threadIdx.y][threadIdx.x]=0;
        }
        else
        {
            tile_b[threadIdx.y][threadIdx.x]=d_b[idx];
        }

        __syncthreads();

        //相乘累加到一个temp中
        for (int k = 0; k < BLOCK_SIZE; ++k) 
        {
            tmp += tile_a[threadIdx.y][k] * tile_b[k][threadIdx.x];
        }
        __syncthreads();
    }

    if(row < M && col < K)
    {
        d_result[row * K + col] = tmp;
    }
}

void cpu_matrix_mult(int *h_a, int *h_b, int *h_result, int m, int n, int k) {
    for (int i = 0; i < m; ++i) 
    {
        for (int j = 0; j < k; ++j) 
        {
            int tmp = 0.0;
            for (int h = 0; h < n; ++h) 
            {
                tmp += h_a[i * n + h] * h_b[h * k + j];
            }
            h_result[i * k + j] = tmp;
        }
    }
}

int main(int argc, char const *argv[])
{
    int m=1000;
    int n=1000;
    int k=1000;

    int *h_a, *h_b, *h_c, *h_cc, *h_cs;
    //void **) &h_a将分配的内存地址赋值给指针h_a 
    //sizeof(int)*m*n)分配一个 int 类型的数组,其大小为 m*n 个元素
    //即主机端分配一个 int 类型的二维数组,大小为 m 行、n 列。
    CHECK(cudaMallocHost((void **) &h_a, sizeof(int)*m*n));
    CHECK(cudaMallocHost((void **) &h_b, sizeof(int)*n*k));
    CHECK(cudaMallocHost((void **) &h_c, sizeof(int)*m*k));
    CHECK(cudaMallocHost((void **) &h_cc, sizeof(int)*m*k));
    CHECK(cudaMallocHost((void **) &h_cs, sizeof(int)*m*k));
    

    /*CUDA事件
    cudaEvent_t: 这是CUDA库中表示事件的数据类型。
    CUDA事件通常用于测量 GPU 执行时间,用于确定 GPU 执行某个任务所花费的时间。
    cudaEventCreate表示用于创建 CUDA 事件
    原型“cudaError_t cudaEventCreate(cudaEvent_t* event);”
    
    */
    cudaEvent_t start, stop,stop_share;
    CHECK(cudaEventCreate(&start));
    CHECK(cudaEventCreate(&stop));
    CHECK(cudaEventCreate(&stop_share));

    //将整个二维矩阵h_a赋值为1
    //数组是在内存按顺序存储,因此二维其实也是一维
    for (int i = 0; i < m; ++i) {
        for (int j = 0; j < n; ++j) {
            h_a[i * n + j] = 1;
        }
    }

    for (int i = 0; i < n; ++i) {
        for (int j = 0; j < k; ++j) {
            h_b[i * k + j] = 0;
        }
    }

    int *d_a, *d_b, *d_c, *d_c_share;
    CHECK(cudaMalloc((void **) &d_a, sizeof(int)*m*n));
    CHECK(cudaMalloc((void **) &d_b, sizeof(int)*n*k));
    CHECK(cudaMalloc((void **) &d_c, sizeof(int)*m*k));
    CHECK(cudaMalloc((void **) &d_c_share, sizeof(int)*m*k));
    /*  cudaDeviceSynchronize()目的是确保设备上的所有操作都已完成。
    这段程序片段的目的是确保CUDA分配的设备内存已经成功完成。*/
    cudaDeviceSynchronize(); 
    // cudaEventRecord是CUDA中记录事件的函数,用于将当前时间戳记录到由start表示的事件中。
    CHECK(cudaEventRecord(start));
    // copy matrix A and B from host to device memory(cudaMemcpy全局内存)
    CHECK(cudaMemcpy(d_a, h_a, sizeof(int)*m*n, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_b, h_b, sizeof(int)*n*k, cudaMemcpyHostToDevice));

    unsigned int grid_rows = (m + BLOCK_SIZE - 1) / BLOCK_SIZE;
    unsigned int grid_cols = (k + BLOCK_SIZE - 1) / BLOCK_SIZE;
    /*
    dim3三维网格和线程块,z=1,blockidx.z=0,threadidx.z=0
    创建了一个二维的线程块与线程网格
    */
    dim3 dimGrid(grid_cols, grid_rows);
    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
    

    gpu_matrix_mult<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m,n,k);    

    CHECK(cudaMemcpy(h_c, d_c, (sizeof(int)*m*k), cudaMemcpyDeviceToHost));
    /*使用这两个函数可以计算CUDA代码的执行时间,
    方法是在cudaEventRecord(start)和cudaEventRecord(stop)之间执行CUDA代码
    cudaEventSynchronize(stop):等待stop事件完成,确保GPU执行到此处时
    已完成前面的所有操作。
    这是一种同步操作,保证我们能够准确地测量CUDA代码的执行时间。
    */
    CHECK(cudaEventRecord(stop));
    CHECK(cudaEventSynchronize(stop));
    
   
    gpu_matrix_mult_shared<<<dimGrid, dimBlock>>> (d_a, d_b, d_c_share, m,n,k);
    CHECK(cudaMemcpy(h_cs, d_c_share, (sizeof(int)*m*k), cudaMemcpyDeviceToHost));
    
    CHECK(cudaEventRecord(stop_share));
    CHECK(cudaEventSynchronize(stop_share));
    
    //通过cudaEventElapsedTime计算两个事件之间的时间差。
    float elapsed_time, elapsed_time_share;
    CHECK(cudaEventElapsedTime(&elapsed_time, start, stop));
    CHECK(cudaEventElapsedTime(&elapsed_time_share, stop, stop_share));
    printf("Time_global = %g ms.\n", elapsed_time);
    printf("Time_share = %g ms.\n", elapsed_time_share);
    
    /*
cudaEventDestroy函数用于销毁CUDA事件。在CUDA程序中,为了避免内存泄漏,
应该在不再需要使用CUDA事件的地方调用cudaEventDestroy来释放相应的资源。*/
    CHECK(cudaEventDestroy(start));
    CHECK(cudaEventDestroy(stop));    

    //cpu_matrix_mult(h_a, h_b, h_c, m, n, k);

    // int ok = 1;
    /*for (int i = 0; i < m; ++i)
    { 
        for (int j = 0; j < k; ++j)
        {
            if(fabs(h_cs[i*k + j] - 0)>(1.0e-10))
            {
                printf("hcs: %d hc: %d  ",h_cs[i*k + j], h_c[i*k + j]);
                ok = 0;
            }
        }
    }

    if(ok)
    {
        printf("Pass!!!\n");
    }
    else
    {
        printf("Error!!!\n");
    }
    */
    // free memory
    CHECK(cudaFree(d_a));
    CHECK(cudaFree(d_b));
    CHECK(cudaFree(d_c));
    CHECK(cudaFreeHost(h_a));
    CHECK(cudaFreeHost(h_b));
    CHECK(cudaFreeHost(h_c));
    return 0;
}

解释部分代码:
CUDA笔记3,CUDA,笔记,算法

参考链接:
https://zhuanlan.zhihu.com/p/654027980
https://www.bilibili.com/video/BV17T4y117vK?p=5文章来源地址https://www.toymoban.com/news/detail-822058.html

到了这里,关于CUDA笔记3的文章就介绍完了。如果您还想了解更多内容,请在右上角搜索TOY模板网以前的文章或继续浏览下面的相关文章,希望大家以后多多支持TOY模板网!

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处: 如若内容造成侵权/违法违规/事实不符,请点击违法举报进行投诉反馈,一经查实,立即删除!

领支付宝红包 赞助服务器费用

相关文章

  • CUDA学习笔记(七)Kernel性能调节

    本篇博文转载于https://www.cnblogs.com/1024incn/tag/CUDA/,仅用于学习。 这部分主要介绍并行分析,涉及掌握nvprof的几个metric参数,具体的这些调节为什么会影响性能会在后续博文解释。 下面是我们的kernel函数sumMatrixOnGPUD: 我们指定一个比较大的数据矩阵,包含16384个元素: 下面的

    2024年02月07日
    浏览(31)
  • CUDA 的随机数算法 API

    参考自 Nvidia cuRand 官方 API 文档 如下是是在 dropout 优化中手写的 uniform_random 的 Kernel: 我们首先来看 curand_init 函数的签名和语义: 给定相同的seed、sequence、offset 参数下, curand_init 会保证产生相同的其实状态 state。另外此函数会在调用 2^67 ⋅ sequence + offset 次 cu_rand API 之后「

    2024年02月03日
    浏览(45)
  • [CUDA 学习笔记] 如何优化 CUDA 矩阵乘内核以获得类似 cuBLAS 的性能: 工作日志

    注: 本文主要是对博文 “How to Optimize a CUDA Matmul Kernel for cuBLAS-like Performance: a Worklog - SIBOEHM” 的翻译, 并进行了一定的备注和补充 在这篇文章中, 我将迭代优化用 CUDA 编写的矩阵乘法的实现. 我的目标不是构建一个 cuBLAS 替代品, 而是深入了解用于现代深度学习的 GPU 的最重要的

    2024年04月28日
    浏览(38)
  • 不是Nvidia(英伟达)显卡可以安装CUDA跑深度学习算法吗?

    答:不行! Cuda主要是面向Nvidia的GPU的。Intel和AMD的显示芯片都不能进行。所以,想要让cuda环境搭建在自己的Windows系统上利用GPU的并行计算跑深度学习算法,就必须要有Nvidia显卡且要安装CUDA。 下面两张图片是AMD显卡和Nvidia显卡的对照: AMD显卡:😕/img-blog.csdnimg.cn/74ef793a4a0

    2024年02月11日
    浏览(70)
  • pytorch-v2.0.1 cuda arm64 aarch64 torch 2.0.1+cu118 源码编译笔记【2】验证cuda安装 成功

    接上篇 pytorch-v2.0.1 cuda arm64 aarch64 torch 2.0.1+cu118 源码编译笔记_hkNaruto的博客-CSDN博客 由于采用/usr/local/bin/gcc编译,先设置LD_LIBRARY_PATH,再启动python3 import torch报错 # /usr/local/Python-3.10.12/bin/python3 Python 3.10.12 (main, Sep  4 2023, 10:01:29) [GCC 9.2.0] on linux Type \\\"help\\\", \\\"copyright\\\", \\\"credits\\\" or \\\"l

    2024年02月09日
    浏览(42)
  • 基于OpenCV+CUDA实时视频抠绿、背景合成以及抠绿算法小结

    百度百科上描述抠绿“抠绿是指在摄影或摄像时,以绿色为背景进行拍摄,在后期制作时使用特技机的“色键”将绿色背景抠去,改换其他更理想的背景的技术。”绿幕的使用已经非常普遍,大到好莱坞大片,小到自媒体的节目,一些商业娱乐场景,几乎都用使用。但是很多

    2023年04月09日
    浏览(62)
  • 【BBuf的CUDA笔记】九,使用newbing(chatgpt)解析oneflow softmax相关的fuse优化

    随着年纪越来越大,读代码越来越困难,如果你发现看不懂同事写的代码应该怎么办呢?不要担心,大语言模型的时代了来了,chatgpt和gpt4会教会我们怎么读代码。本篇文章就来展示一下使用newbing(chatgpt)来读oneflow softmax相关的fuse优化kernel的过程。本文的代码解释均由chat

    2024年02月01日
    浏览(43)
  • 模型实战(11)之win10下Opencv+CUDA部署yolov5、yolov8算法

    测试环境:AMDRH7000+RTX3050+win10+vs2-10+opencv455+cuda11.7 关于opencv470+contrib+cuda的编译,可以详见:Win10下Opencv+CUDA联合编译详细教程 本文代码同时支持 yolov5 、 yolov8 两个模型,详细过程将在文中给出, 完整代码仓库最后给出 其中,yolov8在opencv-DNN + CUDA下的效果如下: 新建VS项目,名

    2024年02月16日
    浏览(40)
  • aarch64 arm64 部署 stable diffusion webui 笔记 【1】准备 venv 安装pytorch 验证cuda

    aarch64 pytorch(没有aarch64对应版本,自行编译) pytorch-v2.0.1 cuda arm64 aarch64 torch 2.0.1+cu118 源码编译笔记【2】验证cuda安装 成功_hkNaruto的博客-CSDN博客 [root@ceph3 stable-diffusion-webui]# /usr/local/Python-3.10.12/bin/python3 -m venv venv [root@ceph3 stable-diffusion-webui]# source venv/bin/activate (venv)[root@ceph3 stab

    2024年02月07日
    浏览(44)
  • 学习笔记:在Anaconda环境下装Pytorch,CUDA,pillow,numpy等库并寻找适配的Python版本

    目录 1.在Anaconda Prompt创建新环境 2.去pytorch官网上查找环境中Python所对应的pytorch版本并下载相关包 查询电脑独立NVIDIA显卡所适配的CUDA版本 CUDA版本查询 3.部分库版本的安装与修改  本人在Anaconda下创建的新环境为 python 3.7.0 pytorch 1.8.0 pillow 9.5.0 numpy 1.21.5 能够正常运行 如果我这

    2024年02月05日
    浏览(44)

觉得文章有用就打赏一下文章作者

支付宝扫一扫打赏

博客赞助

微信扫一扫打赏

请作者喝杯咖啡吧~博客赞助

支付宝扫一扫领取红包,优惠每天领

二维码1

领取红包

二维码2

领红包