cuda 学习笔记4

一 基本函数

在这里插入图片描述
在这里插入图片描述

在GPU上开辟空间，无论定义的数据是float还是int ,还是****gpu_int,分配空间的函数都是下面固定的形式 (void**)&

在这里插入图片描述

1.函数定义，global void 是配套使用的，是在GPU上定义，也就是GPU上执行，CPU上调用的函数，因为CPU不能识别GPU上运算得到的结果，也就是说在CPU上调用这个函数，是不可能存在return的结果的，所以没有返回值

在这里插入图片描述


#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

// __是GOU函数的标识符，加了__标识这个函数是在GPU上被调用的

//这个函数在GPU上是无法调用的，所谓在GPU上，就是定义的在GPU上的函数内调用，因为该函数当前是在CPU上
int add_one(int a)
{
    return a + 1;
}

//这个函数是被定义在GPU上，由GPU本身调用的函数，所以加了这个__devide__标识之后，就可以在__global__ void show(int *a)函数里面调用该函数
__device__ int add_one(int a)
{
    return a + 1;
}

//但是加了__devide__标识之后，由于表示只能在GPU上调用该函数，所以在mian函数里面如果想要实现a[i] = add_one(a[i]);是不可以的
//所以需要再加一个标识符
__host__ __device__ int add_one(int a)
{
    return a + 1;
}


__global__ void show(int *a)
{
    for (int i = 0; i < 10; i++)
    {
        a[i] = add_one(a[i]);
        printf(" %d ", a[i]);
   }
    printf("\n");
}


#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

// __是GOU函数的标识符，加了__标识这个函数是在GPU上被调用的

//这个函数在GPU上是无法调用的，所谓在GPU上，就是定义的在GPU上的函数内调用，因为该函数当前是在CPU上
//int add_one(int a)
//{
//    return a + 1;
//}
//
这个函数是被定义在GPU上，由GPU本身调用的函数，所以加了这个__devide__标识之后，就可以在__global__ void show(int *a)函数里面调用该函数
//__device__ int add_one(int a)
//{
//    return a + 1;
//}

//但是加了__devide__标识之后，由于表示只能在GPU上调用该函数，所以在mian函数里面如果想要实现a[i] = add_one(a[i]);是不可以的
//所以需要再加一个标识符
__host__ __device__ int add_one(int a)
{
    return a + 1;
}


__global__ void show(int *a)
{
    for (int i = 0; i < 10; i++)
    {
      //  a[i] = add_one(a[i]);
        printf(" %d ", a[i]);
   }
    printf("\n");
}

__global__ void int_gpu(int* a)
{
    for (int i = 0; i < 10; i++)
    {
        a[i] = 100;
    }
}

int main()
{
    
    int cpu[10] = { 10,  10, 10, 10, 10, 10, 10, 10, 10, 10};

    //在GPU上分配空间存储CPU上的数据
    int* gpu_int;
    cudaMalloc((void**)&gpu_int, 10*sizeof(int)); //将指针指向GPU的一个内存地址，
    show << <1, 1 >> > (gpu_int);//一个网格里面只有一个block,也就是只有一个线程

    //GPU上数组初始化
    cudaMemset(gpu_int, 0, 10 * sizeof(int));

    // 将CPU上的数据拷贝到GPU上
    cudaMemcpy(gpu_int, cpu, 10*sizeof(int), cudaMemcpyHostToDevice);    //将数据从GPU拷贝到CPU,同时指定拷贝的长度
    
    // 在CPU上调用GPU上的函数进行计算，定义在GPU上的函数就是在GPU上进行运算
    show << <1, 1 >> > (gpu_int);//一个网格里面只有一个block,也就是只有一个线程
    int_gpu << <1, 1 >> > (gpu_int);
    show << <1, 1 >> > (gpu_int);

    // 将GPU上的数据拷贝到cPU上
    cudaMemcpy(cpu, gpu_int, 10*sizeof(int), cudaMemcpyDeviceToHost);    //将数据从GPU拷贝到CPU,同时指定拷贝的长度
    cudaFree(gpu_int);
    cudaDeviceSynchronize();

    printf(" cpu \n");
    for (int i = 0; i < 10; i++)
    {
        
        printf(" %d ", cpu[i]);
    }

    return 0;
}

二 gird, block, thread 之间的关系和理解

在这里插入图片描述
![在这里插入图片描述](https://img-blog.csdnimg.cn/direct/46d1cf35506748029f14d1215f493bcb.png#pic_center =700x
)

1 同步的使用时机

在这里插入图片描述

2 尽量避免直接从globla memory上频繁读写，可以将数据拷贝到share memory再进行对同一数据频繁的读写操作

local memory，数据读取是很快的
global memeory,从这里读取数据是很慢的
share memory, 是block自己共享的，每个block自己可以读自己内部的数据，一个block内部的线程可以访问自己block的share memory

3

threadIdx是指当前线程在当前线程块里面是排几号，是相对block来说的
blockIdx是当前线程块在当前grid里面，x这个维度是第几个线程块
blockDim是维度的意思
在这里插入图片描述

在这里插入图片描述
下面这种情况对应的就是grid是2维，block是二维的情况，含有Dim的是对整个grid维度或者blcok维度的定义，也就是有几个
在这里插入图片描述

在这里插入图片描述

矩阵运算

案例1 矩阵维度为20*20，获取矩阵位置坐标，x,y并返回x+y

在这里插入图片描述

在这里插入图片描述


#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include < iostream>

using namespace std;

//20*20的矩阵， 坐标（x,y）位置赋值x+y

__device__ int coord_int(int x, int y)
{
    return x + y;
}

__global__ void Matrix_init(int *a, int m, int n)
{
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < m && y < n)
    {
        a[y * n + x] = coord_int(x, y);
    }
}

void show(int* a, int m, int n)
{
    for (int i = 0; i < n; i++)
    {
        for ( int j = 0; j < m; j++)
        {
            cout << a[i * m + j] << " ";

        }
        cout << endl;
    }
}

int main()
{
    int* gpu_int;
    cudaMalloc((void**)&gpu_int, 400 * sizeof(int));
    int cpu_int[400] = { 0 };
    show(cpu_int, 20, 20);
    dim3 blockdim(8, 8);  //线程数要大于400
    dim3 griddim(3, 3);
    //为什么数组要开辟一个新的空间且要用指针，而数据不用，因为数据编译的时候是可以直接读的，不需要开辟空间
    Matrix_init << <griddim, blockdim >> > (gpu_int, 20, 20);
    cudaMemcpy(cpu_int, gpu_int, 400 * sizeof(int), cudaMemcpyDeviceToHost);
    cudaFree(gpu_int);
    show(cpu_int, 20, 20);
    return 0;
}

案例2 矩阵乘法和矩阵加法

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <iostream>

using namespace std;

//矩阵20*20
__host__ __device__ int add_one(int a)
{
    return a + 1;
}

__global__ void Matrix_add(int* a, int* b, int* c, int m, int n)
{
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    if (x < m && y < n)
    {
        c[y * n + x] = a[y * n + x] + b[y * n + x];
    }
}

//矩阵乘法：前一个矩阵的行*后一个矩阵的列，然后相加
//只适用与方阵
__global__ void Matrix_multi(int* a, int* b, int* c, int m)
{
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    int output = 0;
    //c[y * m + x]   = 0; 
    if (x < m && y < m)
    {
        for (int i = 0; i < m; i++)
        {
            output +=  a[y * m + i] + b[i * m + x];
            //c[y * m + x] += a[y * m + i] + b[i * m + x];  这里相当于全局内存反复读取c数组，会导致速度慢很多
        }
        c[y * m + x] = output;
    }
}

//cpu上的函数
void show(int* a, int m, int n)
{
    for (int i = 0; i < m; i++)
    {
        for (int j = 0; j < n; j++)
        {
            //  a[i] = add_one(a[i]);
            cout << a[i * m + j] << " ";
        }
        cout << endl;
    }
    
}

__global__ void int_gpu(int* a,int m, int  n)
{
    for (int i = 0; i < m; i++)
    {
        for (int j = 0; j < n; j++)
        {
            a[i * m + j] = 10;
        }
    }
}

//__global__ void my_gpu_multi(int* a, int* b, int* c, int m)
//{
//    __share__ int 
//}


int main()
{

    int cpu[400] = { 0 };
    show(cpu, 20, 20);
 

    //在GPU上分配空间存储CPU上的数据
    int* gpu_int;
    cudaMalloc((void**)&gpu_int, 400 * sizeof(int)); //将指针指向GPU的一个内存地址，
    int* gpu_add;
    cudaMalloc((void**)&gpu_add, 400 * sizeof(int)); //将指针指向GPU的一个内存地址，
    int* gpu_multi;
    cudaMalloc((void**)&gpu_multi, 400 * sizeof(int)); //将指针指向GPU的一个内存地址，



    //GPU上数组初始化
    cudaMemset(gpu_int, 0, 400 * sizeof(int));
    cudaMemset(gpu_add, 0, 400 * sizeof(int));
    cudaMemset(gpu_multi, 0, 400 * sizeof(int));

    // 将CPU上的数据拷贝到GPU上
    cudaMemcpy(gpu_int, cpu, 400 * sizeof(int), cudaMemcpyHostToDevice);    //将数据从GPU拷贝到CPU,同时指定拷贝的长度
    cudaMemcpy(gpu_add, cpu, 400 * sizeof(int), cudaMemcpyHostToDevice);    //将数据从GPU拷贝到CPU,同时指定拷贝的长度
    cudaMemcpy(gpu_multi, cpu, 400 * sizeof(int), cudaMemcpyHostToDevice);    //将数据从GPU拷贝到CPU,同时指定拷贝的长度

    // 在CPU上调用GPU上的函数进行计算，定义在GPU上的函数就是在GPU上进行运算
    dim3 blockdim(8, 8);
    dim3 griddim(3, 3);


    int_gpu << <griddim, blockdim >> > (gpu_int, 20, 20);
    Matrix_add << <griddim, blockdim >> > (gpu_int, gpu_int, gpu_add, 20,20);
    Matrix_multi << <griddim, blockdim >> > (gpu_int, gpu_int, gpu_multi,  20);

    // 将GPU上的数据拷贝到cPU上
    cudaMemcpy(cpu, gpu_int, 400 * sizeof(int), cudaMemcpyDeviceToHost);    //将数据从GPU拷贝到CPU,同时指定拷贝的长度
    show(cpu, 20, 20);//一个网格里面只有一个block,也就是只有一个线程
    cudaMemcpy(cpu, gpu_add, 400 * sizeof(int), cudaMemcpyDeviceToHost);    //将数据从GPU拷贝到CPU,同时指定拷贝的长度
    show(cpu, 20, 20);//一个网格里面只有一个block,也就是只有一个线程
    cudaMemcpy(cpu, gpu_multi, 400 * sizeof(int), cudaMemcpyDeviceToHost);    //将数据从GPU拷贝到CPU,同时指定拷贝的长度
    show(cpu, 20, 20);//一个网格里面只有一个block,也就是只有一个线程

    cudaFree(gpu_int);
    cudaFree(gpu_add);
    cudaFree(gpu_multi);
    cudaDeviceSynchronize();

    return 0;
}

案例3 用共享内存实现矩阵乘法

在这里插入图片描述

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <iostream>
#include <curand.h>
#include <curand_kernel.h>


using namespace std;
#define ARRAY_SIZE 16
#define ARRAY_LENGTH 256
#define BLOCK_SIZE 16



//GPU上初始化,N是
__global__ void gpu_initial(float *a ,int N) {

    int x = threadIdx.x + blockDim.x * blockIdx.x;
    curandState state;
    long seed = N;
    curand_init(seed, x,0,&state);
    if (x < N) a[x] = curand_uniform(&state);
}

//在CPU上使用随机数初始化，N是要初始化的数组长度
void cpu_initial(float *a, int N)
{
    curandGenerator_t gen;
    curandCreateGenerator(&gen, CURAND_RNG_PSEUDO_DEFAULT);  //告诉电脑要用什么样的方法生成随机数
    curandSetPseudoRandomGeneratorSeed(gen, 11ULL);   //指定随机数种子
    curandGenerateUniform(gen, a, N);  //curandGenerateUniform， 生成均匀分布的0-1的随机数，让a里面的数据从1-N的数据都进行赋值
}

void show(float* a, int m, int n)
{
    for (int i = 0; i < m; i++)
    {
        for (int j = 0; j < n; j++)
        {
            //  a[i] = add_one(a[i]);
            cout << a[i * m + j] << " ";
        }
        cout << endl;
    }
    
}

//矩阵乘法：前一个矩阵的行*后一个矩阵的列，然后相加
//只适用与方阵
__global__ void Matrix_multi(float* a, float* b, float* c, int m)
{
    int x = threadIdx.x;
    int y = threadIdx.y;

    int output = 0;
    if (x < m && y < m)
    {
        for (int i = 0; i < m; i++)
        {
            output += a[y * m + i] * b[i * m + x];
        }
        c[y * m + x] = output;
    }
}

__global__ void my_gpu_multi(float* a, float* b, float* c, int m)
{
    int x = threadIdx.x;
    int y = threadIdx.y;
    __shared__ float a_share[256];
    __shared__ float b_share[256];
    if (x < m && y < m)
    {
        a_share[y * m + x] = a[y * m + x];
        b_share[y * m + x] = b[y * m + x];
    }

    __syncthreads();
    float output = 0;
    if (x < m && y < m)
    {
        for (int i = 0; i < m; i++)
        {
            output += a_share[y*m + i] * b_share[i*m + x];
        }
        c[y * m + x] = output;
    }
    
}

@fighting 
共享内存大小不足：
共享内存的大小由 __shared__ float a_share[256]; 和 __shared__ float b_share[256]; 决定。如果 m 值大于 16，那么共享内存大小可能不足，因为 16x16 = 256。如果 m 大于 16，应该调整共享内存大小。
同步问题：
在 my_gpu_multi 中使用了 __syncthreads() 来确保所有线程都完成了数据拷贝。但如果有些线程的计算还没完成就进入下一步，可能会导致不一致的结果。
线程边界检查：
如果 m 的值较大而 block 尺寸 (m, m) 超出了 GPU 硬件的限制，可能会导致一些线程未能正确启动，导致结果不一致。

int main()
{
    int m = 6;
    int N = m * m;
    float* p_d, * p_da, * p_db, * p_h;

    //cpu上开辟空间
    p_h = (float*)malloc(N * sizeof(float));

    //GPU上开辟空间
    cudaMalloc((void**)&p_d, N * sizeof(float)); //将指针指向GPU的一个内存地址，
    cudaMalloc((void**)&p_da, N * sizeof(float)); //将指针指向GPU的一个内存地址，
    cudaMalloc((void**)&p_db, N * sizeof(float)); //将指针指向GPU的一个内存地址，

    //数组初始化，使用两种不同的方式进行初始化
    gpu_initial <<<16, 16 >>>(p_da, N);  //直接在GPU上初始化
    cpu_initial(p_db, N);   //在CPU上初始化,但是由于是调用的Gpu上的函数进行的初始化，所以最后还是等价于在GPU上初始化的


    cudaMemcpy(p_h, p_da, N * sizeof(float), cudaMemcpyDeviceToHost);    //将数据从GPU拷贝到CPU,同时指定拷贝的长度
    printf("\n p_da");
    for (int i = 0; i < N; i++)
    {
        if (i % m == 0)  printf("\n ");
        cout << p_h[i] << " ";
    }

    cudaMemcpy(p_h, p_db, N * sizeof(float), cudaMemcpyDeviceToHost);    //将数据从GPU拷贝到CPU,同时指定拷贝的长度
    printf("\n p_db");
    for (int i = 0; i < N; i++)
    {
        if (i % m == 0)  printf("\n ");
        cout << p_h[i] << " ";
    }
    printf("\n ");


    //share memory只能在一个block上使用，所以只能定义一个block
    dim3 blockdim(m, m);
    //Matrix_multi <<<1, blockdim>>> (p_da, p_db, p_d, m);
     将GPU上的数据拷贝到cPU上
    //cudaMemcpy(p_h, p_d, N * sizeof(float), cudaMemcpyDeviceToHost);    //将数据从GPU拷贝到CPU,同时指定拷贝的长度
    //printf("funtuon Matrix_multi: ------------\n ");
    //show(p_h, m, m);//一个网格里面只有一个block,也就是只有一个线程



    my_gpu_multi <<<1, blockdim>>> (p_da, p_db, p_d, m);
    cudaMemcpy(p_h, p_d, N * sizeof(float), cudaMemcpyDeviceToHost);    //将数据从GPU拷贝到CPU,同时指定拷贝的长度
    printf("funtuon my_gpu_multi: ------------\n ");
    show(p_h, m, m);


    cudaFree(p_da);
    cudaFree(p_db);
    cudaFree(p_d);
    free(p_h);

    return 0;
}

常用官方库的使用

1. cuda案例

cuda函数说明官网
在这里插入图片描述

在这里插入图片描述


//Example 1. Application Using C and cuBLAS: 1-based indexing
//-----------------------------------------------------------
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <cuda_runtime.h>
#include "cublas_v2.h"
#define M 6
#define N 5
#define IDX2F(i,j,ld) ((((j)-1)*(ld))+((i)-1))

static __inline__ void modify(cublasHandle_t handle, float* m, int ldm, int n, int p, int q, float alpha, float beta) {
    cublasSscal(handle, n - q + 1, &alpha, &m[IDX2F(p, q, ldm)], ldm);
    cublasSscal(handle, ldm - p + 1, &beta, &m[IDX2F(p, q, ldm)], 1);
}

int main(void) {
    // cudaError_t可以判断cuda有没有错误，如果定义的cudaStat不是cudaSucess，就会提示错误信息位置
    cudaError_t cudaStat;
    cublasStatus_t stat;
    cublasHandle_t handle;
    int i, j;
    float* devPtrA;
    float* a = 0;
    a = (float*)malloc(M * N * sizeof(*a));
    if (!a) {
        printf("host memory allocation failed");
        return EXIT_FAILURE;
    }
    for (j = 1; j <= N; j++) {
        for (i = 1; i <= M; i++) {
            a[IDX2F(i, j, M)] = (float)((i - 1) * N + j);
        }
    }
    cudaStat = cudaMalloc((void**)&devPtrA, M * N * sizeof(*a));
    if (cudaStat != cudaSuccess) {
        printf("device memory allocation failed");
        free(a);
        return EXIT_FAILURE;
    }
    stat = cublasCreate(&handle);
    if (stat != CUBLAS_STATUS_SUCCESS) {
        printf("CUBLAS initialization failed\n");
        free(a);
        cudaFree(devPtrA);
        return EXIT_FAILURE;
    }
    stat = cublasSetMatrix(M, N, sizeof(*a), a, M, devPtrA, M);
    if (stat != CUBLAS_STATUS_SUCCESS) {
        printf("data download failed");
        free(a);
        cudaFree(devPtrA);
        cublasDestroy(handle);
        return EXIT_FAILURE;
    }
    modify(handle, devPtrA, M, N, 2, 3, 16.0f, 12.0f);
    stat = cublasGetMatrix(M, N, sizeof(*a), devPtrA, M, a, M);
    if (stat != CUBLAS_STATUS_SUCCESS) {
        printf("data upload failed");
        free(a);
        cudaFree(devPtrA);
        cublasDestroy(handle);
        return EXIT_FAILURE;
    }
    cudaFree(devPtrA);
    cublasDestroy(handle);
    for (j = 1; j <= N; j++) {
        for (i = 1; i <= M; i++) {
            printf("%7.0f", a[IDX2F(i, j, M)]);
        }
        printf("\n");
    }
    free(a);
    return EXIT_SUCCESS;
}

2.cuda自带函数：实现矩阵乘法运算 cublasSgemm

在这里插入图片描述
lad是A的行数，ldb是B的行数
在这里插入图片描述


//Example 1. Application Using C and cuBLAS: 1-based indexing
//-----------------------------------------------------------
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <cuda_runtime.h>
#include "cublas_v2.h"
#include <ctime>
#include <iostream>

using namespace std;


int main()
{
	srand(time(0));

	// C = A*B
	int M = 2;  //矩阵A的行，矩阵C的行
	int N = 3;  //矩阵A的列，矩阵B的行
	int K = 4;  //矩阵B的列，矩阵C的列

	
	float* h_A = (float*)malloc(sizeof(float) * M * N);
	float* h_B = (float*)malloc(sizeof(float) * N * K);
	float* h_C = (float*)malloc(sizeof(float) * M * K);


	//*******************************  矩阵初始化
	cout << " A ----------- " << endl;
	for (int i = 0; i < M * N; i++)
	{
		h_A[i] = rand() % 10;
		cout << h_A[i] << " ";
		if ((i + 1) % N == 0)
			cout << endl;
	}
	cout << endl;

	cout << " B ----------- " << endl;
	for (int i = 0; i < N * K; i++)
	{
		h_B[i] = rand() % 10;
		cout << h_B[i] << " ";
		if ((i + 1) % K == 0)
			cout << endl;
	}
	cout << endl;

	//******************************* 在GPU上开辟空间存储数据
	float *d_A, *d_B, *d_C;
	cudaMalloc( (void**)&d_A, sizeof(float) * M * N);
	cudaMalloc((void**)&d_B, sizeof(float) * N * K);
	cudaMalloc((void**)&d_C, sizeof(float) * M * K);

	//******************************* 将数据从CPU拷贝到GPU
	cudaMemcpy(d_A, h_A, sizeof(float) * M * N, cudaMemcpyHostToDevice);
	cudaMemcpy(d_B, h_B, sizeof(float) * N * K, cudaMemcpyHostToDevice);

	float alpha = 1;
	float beta = 0;

	cublasHandle_t handle;
	cublasCreate(&handle);
	cublasSgemm(handle,
		CUBLAS_OP_N,  //数据不转置
		CUBLAS_OP_N,
		K,    //矩阵B的列
		M,    //矩阵A的行
		N,    //矩阵A的列
		&alpha,
		d_B,
		K,
		d_A,
		N,
		&beta,
		d_C,
		K);

	cudaMemcpy(h_C, d_C, sizeof(float) * M * K, cudaMemcpyDeviceToHost);


	cout << " c ----------- " << endl;
	for (int i = 0; i < M * K; i++)
	{
		cout << h_C[i] << " ";
		if ((i + 1) % K == 0)
			cout << endl;
	}
	cout << endl;

	cudaFree(d_A);
	cudaFree(d_B);
	cudaFree(d_C);

	return 0;
}

3.cuda自带函数：实现矩阵每个数的翻倍， cublasSccal

把间隔设置为1，就可以实现把当前矩阵每一个元素都乘以这么一个常量
在这里插入图片描述


//Example 1. Application Using C and cuBLAS: 1-based indexing
//-----------------------------------------------------------
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <cuda_runtime.h>
#include "cublas_v2.h"
#include <ctime>
#include <iostream>

using namespace std;


int main()
{
	srand(time(0));

	// C = A*B
	int M = 2;  //矩阵A的行，矩阵C的行
	int N = 3;  //矩阵A的列，矩阵B的行

	
	float* h_A = (float*)malloc(sizeof(float) * M * N);


	//*******************************  矩阵初始化
	cout << " A ----------- " << endl;
	for (int i = 0; i < M * N; i++)
	{
		h_A[i] = rand() % 10;
		cout << h_A[i] << " ";
		if ((i + 1) % N == 0)
			cout << endl;
	}
	cout << endl;

	//******************************* 在GPU上开辟空间存储数据
	float *d_A, *d_B, *d_C;
	cudaMalloc( (void**)&d_A, sizeof(float) * M * N);


	//******************************* 将数据从CPU拷贝到GPU
	cudaMemcpy(d_A, h_A, sizeof(float) * M * N, cudaMemcpyHostToDevice);

	float alpha = 2.2;
	cublasHandle_t handle;
	cublasStatus_t stat;

	//函数用于初始化CUBLAS库并创建一个CUBLAS上下文。上下文被表示为一个cublasHandle_t类型的句柄。
	//handle是一个指向cublasHandle_t类型的指针，cublasCreate函数会将创建的句柄存储在这个位置。
	//如果初始化成功，stat将返回CUBLAS_STATUS_SUCCESS，否则会返回一个错误码。
	stat = cublasCreate(&handle);  
	if (stat != CUBLAS_STATUS_SUCCESS) {
		printf("CUBLAS initialization failed\n");
		return EXIT_FAILURE;
	}

	stat = cublasSscal(handle,6, &alpha, d_A,1);
	if (stat != CUBLAS_STATUS_SUCCESS) {
		printf("CUBLAS scaling failed\n");
		return EXIT_FAILURE;
	}

	cudaMemcpy(h_A, d_A, sizeof(float) * M * N, cudaMemcpyDeviceToHost);


	cout << " c ----------- " << endl;
	for (int i = 0; i < M * N; i++)
	{
		cout << h_A[i] << " ";
		if ((i + 1) % N == 0)
			cout << endl;
	}
	cout << endl;

	cudaFree(d_A);


	return 0;
}

在这里插入图片描述

4. 傅里叶变换 https://docs.nvidia.com/cuda/cufft/index.html

cufftExecC2C() and cufftExecZ2Z()是一样的，前者是浮点型，后者是double，后者精度更高
在这里插入图片描述
在这里插入图片描述


//Example 1. Application Using C and cuBLAS: 1-based indexing
//-----------------------------------------------------------
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <cuda_runtime.h>
#include "cublas_v2.h"
#include <ctime>
#include <iostream>
#include <cufft.h>

using namespace std;


int main()
{
	const int Nt = 256;
	const int BACTH = 1;

	//BACTH 用户霹雳那个处理一批一维的数据，假设数据是512，当BATCH =2,则将0-255,256-512作为两个一维信号做FFT变换
	//定义cufftDoubleComplex类型的指针，device_in记录输入值，device_out计记录计算结果
	cufftDoubleComplex *host_in, * host_out, * device_in, * device_out;


	//这种方式开辟空空间再向GPU传输数据是要快于传统的malloc方式的，两者是等价的
	/*	float* host_in = (float*)malloc(sizeof(float) * Nt);
	float* host_out = (float*)malloc(sizeof(float) * Nt);
   */
	cudaMallocHost((void**)&host_in, sizeof(cufftDoubleComplex) * Nt);
	cudaMallocHost((void**)&host_out, sizeof(cufftDoubleComplex) * Nt);


	for (int i = 0; i < Nt; i++)
	{
		host_in[i].x = i + 1;
		host_in[i].y = i + 1;
	}

	// *************  在GPU上开辟空间
	cudaMalloc((void**)&device_in, sizeof(cufftDoubleComplex) * Nt);
	cudaMalloc((void**)&device_out, sizeof(cufftDoubleComplex) * Nt);

	//从CPU上数据拷贝数据到GPU上
	cudaMemcpy(device_in, host_in, Nt* sizeof(cufftDoubleComplex), cudaMemcpyHostToDevice);

	//创建cufft句柄
	cufftHandle cufftForwrdHandle;
	cufftPlan1d(&cufftForwrdHandle, Nt, CUFFT_Z2Z, BACTH);   //传参

	//CUFFT_Z2Z因为前面定义的数据类型是Double-Complex,所以这里指定的数据类型是这个
	// typedef enum cufftType_t {
	//	CUFFT_R2C = 0x2a,     // Real to Complex (interleaved)
	//	CUFFT_C2R = 0x2c,     // Complex (interleaved) to Real
	//	CUFFT_C2C = 0x29,     // Complex to Complex, interleaved
	//	CUFFT_D2Z = 0x6a,     // Double to Double-Complex
	//	CUFFT_Z2D = 0x6c,     // Double-Complex to Double
	//	CUFFT_Z2Z = 0x69      // Double-Complex to Double-Complex
	//} cufftType;

	//执行fft正变换
	cufftExecZ2Z(cufftForwrdHandle, device_in, device_out, CUFFT_FORWARD);  //正变换是CUFFT_FORWARD，反变换是CUFFT_INVERSE

	// 从GPU 上数据拷贝数据到CPU上
	cudaMemcpy(host_out, device_out, Nt * sizeof(cufftDoubleComplex), cudaMemcpyDeviceToHost);

	//设置输出精度--输出正变换的结果
	cout << " 正变换的结果: " << endl;
	for (int i = 0; i < Nt; i++)
	{
		cout << host_out[i].x << " + j*" << host_out[i].y << endl;
	}
	cudaFree(device_in);
	cudaFree(device_out);
	return 0;
}

三 常见的可以采用并行的模式

1, 一对一 ，例如输入x,函数输出y = x*x

在这里插入图片描述
在这里插入图片描述
在这里插入图片描述

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <iostream>
#include <curand.h>
#include <curand_kernel.h>


using namespace std;
#define ARRAY_SIZE 16
#define ARRAY_LENGTH 256
#define BLOCK_SIZE 16



//GPU上初始化,N是
__global__ void gpu_initial(float *a ,int N) {

    int x = threadIdx.x + blockDim.x * blockIdx.x;
    curandState state;
    long seed = N;
    curand_init(seed, x,0,&state);
    if (x < N) a[x] = curand_uniform(&state);
}



__global__ void square(float *d_out,int N)
{
    int myID = threadIdx.x + blockDim.x * blockIdx.x;

    if (myID < N)
    {
        float data = d_out[myID];
        d_out[myID] = data * data;
    }
}

int main()
{
    int m = 6;
    int N = m * m;
    float* d_out, *h_in, *h_out;

    //cpu上开辟空间
    h_in = (float*)malloc(N * sizeof(float));
    h_out = (float*)malloc(N * sizeof(float));

    //GPU上开辟空间
    cudaMalloc((void**)&d_out, N * sizeof(float)); //将指针指向GPU的一个内存地址，

    //数组初始化，使用两种不同的方式进行初始化
    //直接在GPU上初始化, 4 是 numBlocks，表示启动了 4 个块（block）,16 是 numThreadsPerBlock，表示每个块中有 16 个线程（thread)
    //这种初始化方式适用于一维数组
    gpu_initial <<<4, 16 >>>(d_out, N); 
    cudaMemcpy(h_in, d_out, N * sizeof(float), cudaMemcpyDeviceToHost);

    printf("\n h_in");
    for (int i = 0; i < N; i++)
    {
        if (i % m == 0)  printf("\n ");
        cout << h_in[i] << " ";
    }

    square << <4, 16 >> > (d_out, N);


    //将数据从GPU拷贝到CPU,同时指定拷贝的长度
    cudaMemcpy(h_out, d_out, N * sizeof(float), cudaMemcpyDeviceToHost);


    printf("\n h_out");
    for (int i = 0; i < N; i++)
    {
        if (i % m == 0)  printf("\n ");
        cout << h_out[i] << " ";
    }

    cudaFree(d_out);
    free(h_in);
    free(h_out);

    return 0;
}

2. 卷积操作

在这里插入图片描述

  • 我写的
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <iostream>
#include <curand.h>
#include <curand_kernel.h>


using namespace std;
#define ARRAY_SIZE 16
#define ARRAY_LENGTH 256
#define BLOCK_SIZE 16



//GPU上初始化,N是
__global__ void gpu_initial(float* a, int N) {

    int x = threadIdx.x + blockDim.x * blockIdx.x;
    curandState state;
    long seed = N;
    curand_init(seed, x, 0, &state);
    if (x < N) a[x] = curand_uniform(&state);
}


void show(float* a, int m, int n)
{
    for (int i = 0; i < m; i++)
    {
        for (int j = 0; j < n; j++)
        {
            //  a[i] = add_one(a[i]);
            cout << a[i * m + j] << "      ";
        }
        cout << endl;
    }

}


//设置n*n大小的block，把对应坐标的矩阵数据读取到共享内存里，然后其中的(n-2)*(n-2)个线程进行卷积运算
__global__ void Conv_method1(float* p_d, float*kernel, int block_size, int kernel_Len)
{
    int x = threadIdx.x;  //现在是一个一维的一个block
    int y = threadIdx.y;
    extern __shared__ float s_pd[];

    if (x < block_size && y < block_size)
    {
        s_pd[y * block_size + x] = p_d[y * block_size + x];
        //printf("s_pd[y * block_size + x] = %f \n" , s_pd[y * block_size + x]);
    }
    __syncthreads();

    //for (int i = 0; i < kernel_Len; i++) {
    //    printf("kernel[%d] = %f", i, kernel[i]);
    //}
    float out = 0;
    if (x > 0 && y >0 && x < block_size-1  && y < block_size-1)
    {
         out = s_pd[(y-1)* block_size + x - 1] * kernel[0] + s_pd[(y - 1) * block_size + x] * kernel[1] +  s_pd[(y - 1) * block_size + x+1] * kernel[2] + 
                s_pd[y * block_size + x - 1] * kernel[3]    + s_pd[y * block_size + x] * kernel[4]       +  s_pd[y * block_size + x + 1] * kernel[5] + 
                s_pd[(y+1)* block_size + x - 1] * kernel[6] + s_pd[(y + 1) * block_size + x] * kernel[7] + s_pd[(y + 1) * block_size + x + 1] * kernel[8];
         
         //printf(" x = %d,y = %d,out = %f \n",x,y,out);
    }
    p_d[y * block_size + x] = out;
    __syncthreads();
}


int main()
{
    int m = 6;
    int N = m*m;
    int kernelLen = 9;
    float* p_d, * h_in, * h_out;
    float *h_kenel, * d_kernel;

    //cpu上开辟空间
    h_in = (float*)malloc(N * sizeof(float));
    h_out = (float*)malloc(N * sizeof(float));
    h_kenel = (float*)malloc(9 * sizeof(float));

    //GPU上开辟空间
    cudaMalloc((void**)&p_d, N * sizeof(float)); //将指针指向GPU的一个内存地址
    cudaMalloc((void**)&d_kernel, 9 * sizeof(float)); //将指针指向GPU的一个内存地址

     使用循环赋值  
    printf("h_kenel  \n ");
    float values[] = { 1, 1, 1, 1, -8, 1, 1,1,1};
    for (int i = 0; i < kernelLen; i++) {
        h_kenel[i] = values[i];
        cout << h_kenel[i] << " ";
    }
    printf(" \n ");

    cudaMemcpy(d_kernel, h_kenel, kernelLen * sizeof(float), cudaMemcpyHostToDevice);

    gpu_initial << <16, 16 >> > (p_d, N);  //直接在GPU上初始化
    cudaMemcpy(h_in, p_d, N * sizeof(float), cudaMemcpyDeviceToHost);    //将数据从GPU拷贝到CPU,同时指定拷贝的长度


    printf("h_in  \n ");
    show(h_in, m, m);

    dim3 griddim(1, 1);
    dim3 blockdim(m, m);
    Conv_method1 << <griddim, blockdim, sizeof(float)* N >> > (p_d, d_kernel, m, kernelLen);


    //将数据从GPU拷贝到CPU,同时指定拷贝的长度
    cudaMemcpy(h_out, p_d, N * sizeof(float), cudaMemcpyDeviceToHost);


    printf("\n h_out \n  ");
    show(h_out,m,m);

    cudaFree(p_d);
    free(h_in);
    free(h_out);

    return 0;
}
  • 别人写的
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <iostream>
#include <fstream>
#include <iomanip>
#include <time.h>

using namespace std;

// 调用CUDA函数并检查是否出现错误
#define CUDA_CALL(x) {\
	const cudaError_t a = (x);\
	if (a != cudaSuccess) { \
		fprintf(stderr, "\nCUDA Error: %s (err_num=%d)\nfile %s, line %d\n", cudaGetErrorString(a), a, __FILE__, __LINE__);\
		cudaDeviceReset(); exit(1);\
	} \
}

// 调用核函数并检查是否出现错误
__host__ void cuda_error_check(const char* kernelName) {
	if (cudaPeekAtLastError() != cudaSuccess) {
		printf("\n%s %s\n", kernelName, cudaGetErrorString(cudaGetLastError()));
		cudaDeviceReset();
		exit(1);
	}
}

// 设置参数
const int TILE_W = 32;			// block的x维大小
const int TILE_H = 32;			// block的y维大小
const int DATA_W = 320;			// 输入矩阵的x维大小
const int DATA_H = 640;			// 输入矩阵的y维大小
const int KERNEL_RADIUS = 1;	// 卷积核的半径

// 卷积核写入GPU常量内存中
__constant__ int KERNEL[2 * KERNEL_RADIUS + 1][2 * KERNEL_RADIUS + 1] =
{ 1, 1, 1,
1, -8, 1,
1, 1, 1 };

// 通过共享内存，进行卷积运算
__global__ void convolution(float* dst, float* src) {
	// 线程在所在block中的x,y坐标
	int tidx = threadIdx.x;
	int tidy = threadIdx.y;

	// 线程应该读取到shared memory中对应的矩阵元素坐标
	int readx = (blockDim.x - 2 * KERNEL_RADIUS) * blockIdx.x + (threadIdx.x - KERNEL_RADIUS);
	int ready = (blockDim.y - 2 * KERNEL_RADIUS) * blockIdx.y + (threadIdx.y - KERNEL_RADIUS);

	// 除去不需要加载内存的地方
	if (readx >= DATA_W + KERNEL_RADIUS || ready >= DATA_H + KERNEL_RADIUS)
		return;

	// 声明block的共享内存
	__shared__ float src_s[TILE_H][TILE_W];
	// 把当前block需要处理的区域读取到shared memory，输入矩阵周围的记为0
	if (readx >= 0 && readx < DATA_W && ready >= 0 && ready < DATA_H) {
		src_s[tidy][tidx] = src[ready * DATA_W + readx];
	}
	else {
		src_s[tidy][tidx] = 0;
	}
	// 同步block中所有线程，保证共享内存完全读入矩阵
	__syncthreads();

	// 卷积计算
	float output = 0;
	int kernel_w = 2 * KERNEL_RADIUS + 1;
	if (tidx < blockDim.x - 2 * KERNEL_RADIUS && readx < DATA_W - KERNEL_RADIUS &&
		tidy < blockDim.y - 2 * KERNEL_RADIUS && ready < DATA_H - KERNEL_RADIUS) {
		for (int i = 0; i < kernel_w; i++) {
			for (int j = 0; j < kernel_w; j++) {
				output += src_s[tidy + j][tidx + i] * KERNEL[j][i];
			}
		}
		// 写入dst对应坐标
		dst[(ready + KERNEL_RADIUS) * DATA_W + (readx + KERNEL_RADIUS)] = output;
	}
}

int main() {

	const int INPUTSIZE = DATA_H * DATA_W;

	printf("---------- initilizing ----------\n");
	clock_t tt = clock();

	// CPU输入输出矩阵的声明
	float* h_src = (float*)malloc(INPUTSIZE * sizeof(float));
	float* h_dst = (float*)malloc(INPUTSIZE * sizeof(float));

	// 输入矩阵中元素全部设为1
	for (int i = 0; i < DATA_W; i++) {
		for (int j = 0; j < DATA_H; j++) {
			h_src[i + j * DATA_W] = (float)1;
		}
	}

	// 将输入矩阵输出
	ofstream ofs("input_output.txt");
	for (int j = 0; j < DATA_H; j++) {
		for (int i = 0; i < DATA_W; i++) {
			ofs << setw(5) << h_src[i + j * DATA_W];
		}
		ofs << '\n';
	}
	ofs << '\n';

	// 设定使用第一个GPU
	CUDA_CALL(cudaSetDevice(0));

	// GPU输入输出矩阵的声明
	float* d_src = 0;
	float* d_dst = 0;

	// 初始化
	CUDA_CALL(cudaMalloc(&d_src, INPUTSIZE * sizeof(float)));
	CUDA_CALL(cudaMalloc(&d_dst, INPUTSIZE * sizeof(float)));

	CUDA_CALL(cudaMemcpy(d_src, h_src, INPUTSIZE * sizeof(float), cudaMemcpyHostToDevice));
	CUDA_CALL(cudaMemcpy(d_dst, h_dst, INPUTSIZE * sizeof(float), cudaMemcpyHostToDevice));

	// 输出内存初始化所需时间
	printf("Initializaion time(ms) : %f\n", ((float)clock() - tt) / CLOCKS_PER_SEC);

	printf("---------- calculating ----------\n");

	// 设定核函数中grid和block的参数
	//dim3 gridDim = (11, 22);
	/*gridDim定义了网格的维度。在这个例子中，网格有两个维度：
		x 维度有 11 个块
		y 维度有 22 个块*/
	dim3 gridDim = { (DATA_W + (TILE_W - 2 * KERNEL_RADIUS - 1)) / (TILE_W - 2 * KERNEL_RADIUS),
		(DATA_H + (TILE_H - 2 * KERNEL_RADIUS - 1)) / (TILE_H - 2 * KERNEL_RADIUS) };

	printf("---------- calculating ----(DATA_W + (TILE_W - 2 * KERNEL_RADIUS - 1)) / (TILE_W - 2 * KERNEL_RADIUS) = %d, (DATA_H + (TILE_H - 2 * KERNEL_RADIUS - 1)) / (TILE_H - 2 * KERNEL_RADIUS) = %d--\n", 
		(DATA_W + (TILE_W - 2 * KERNEL_RADIUS - 1)) / (TILE_W - 2 * KERNEL_RADIUS),
		(DATA_H + (TILE_H - 2 * KERNEL_RADIUS - 1)) / (TILE_H - 2 * KERNEL_RADIUS));
	dim3 blockDim = { TILE_W, TILE_H };

	// 调用核函数并计时
	cudaEvent_t start, stop;
	cudaEventCreate(&start);
	cudaEventCreate(&stop);
	cudaEventRecord(start, 0);
	convolution << <gridDim, blockDim >> > (d_dst, d_src);

	// 检查核函数调用是否出现错误
	cuda_error_check("convolution_shared_memory");

	// CPU等待GPU完成核函数的计算
	CUDA_CALL(cudaDeviceSynchronize());

	// 输出核函数调用时长
	cudaEventRecord(stop, 0);
	cudaEventSynchronize(stop);
	float elapsedTime;
	cudaEventElapsedTime(&elapsedTime, start, stop);
	printf("Kernel time(ms) : %f\n", elapsedTime);

	cudaEventDestroy(start);
	cudaEventDestroy(stop);

	// CPU获取GPU核函数计算结果
	CUDA_CALL(cudaMemcpy(h_dst, d_dst, INPUTSIZE * sizeof(float), cudaMemcpyDeviceToHost));

	// 输出卷积后的结构
	for (int j = 0; j < DATA_H; j++) {
		for (int i = 0; i < DATA_W; i++) {
			ofs << setw(5) << h_dst[i + j * DATA_W];
		}
		ofs << '\n';
	}
	ofs.close();

	// 释放GPU内存
	cudaFree(d_src);
	cudaFree(d_dst);
	//cudaFree(KERNEL);

	// 释放CPU内存
	free(h_src);
	free(h_dst);

	// 清空重置GPU
	CUDA_CALL(cudaDeviceReset());

	printf("Total calculation time(ms) : %f\n", ((float)clock() - tt) / CLOCKS_PER_SEC);

}

3.

在这里插入图片描述

4.

在这里插入图片描述

5

在这里插入图片描述

动态共享内存的初始化和使用

在这里插入图片描述

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <iostream>
#include <curand.h>
#include <curand_kernel.h>


using namespace std;
#define ARRAY_SIZE 16
#define ARRAY_LENGTH 256
#define BLOCK_SIZE 16



//GPU上初始化,N是
__global__ void gpu_initial(float* a, int N) {

    int x = threadIdx.x + blockDim.x * blockIdx.x;
    curandState state;
    long seed = N;
    curand_init(seed, x, 0, &state);
    if (x < N) a[x] = curand_uniform(&state);
}


//使用共享内存
__global__ void sortFun(float* p_d, int N)
{
    int x = threadIdx.x;  //现在是一个一维的一个block
    //定义动态内存,动态内存的长度是在初始化的时候 sortFun << <1, N, sizeof(float)*N >> > (d_in, N);sizeof(float)*N 指定共享内存的字节大小
    extern __shared__ float temp_d[];  
    temp_d[x] = p_d[x];
    __syncthreads();  //相当于是把所有数据都加载到temp_d之后才开始后面的排序操作


    //奇偶排序，长度为N的数组需要排序N次，所以这里的for循环里面的N是对排序次数的循环，而不是数据
    for (int i = 0; i < N; i++)  
    {
        int j = i % 2;  //先奇偶比，然后下一次才是偶奇比，也就是偶数次是奇偶比，奇数次是偶奇比
        int idx = 2 * x + j;
        if (idx + 1 < N && temp_d[idx] < temp_d[idx + 1])
        {
            float tep = temp_d[idx];
            temp_d[idx] = temp_d[idx + 1];
            temp_d[idx + 1] = tep;
        }
        __syncthreads();  //每排序一次，要等所有数据都判断完毕才进行下一轮的排序
    }
    p_d[x] = temp_d[x];
    __syncthreads();  //所有数据都排序完毕且结束，程序结束
}



int main()
{
    int m = 6;
    int N = m;
    float* p_d, *h_in, *h_out;

    //cpu上开辟空间
    h_in = (float*)malloc(N * sizeof(float));
    h_out = (float*)malloc(N * sizeof(float));

    //GPU上开辟空间
    cudaMalloc((void**)&p_d, N * sizeof(float)); //将指针指向GPU的一个内存地址
     使用循环赋值  
    //int values[] = { 1, 5, 4, 2, 6, 3};
    //for (int i = 0; i < N; i++) {
    //    h_in[i] = values[i];
    //}
    //cudaMemcpy(p_d, h_in, N * sizeof(float), cudaMemcpyHostToDevice);

    gpu_initial << <16, 16 >> > (p_d, N);  //直接在GPU上初始化
    cudaMemcpy(h_in, p_d, N * sizeof(float), cudaMemcpyDeviceToHost);    //将数据从GPU拷贝到CPU,同时指定拷贝的长度

    printf("h_in  \n ");
    for (int i = 0; i < N; i++)
    {
        cout << h_in[i] << " ";
    }

    sortFun <<<1, N, sizeof(float)*N >>> (p_d, N);
  //  sort << <1, N, N * sizeof(float) >> > (p_d, N);

    //将数据从GPU拷贝到CPU,同时指定拷贝的长度
    cudaMemcpy(h_out, p_d, N * sizeof(float), cudaMemcpyDeviceToHost);


    printf("\n h_out \n  ");
    for (int i = 0; i < N; i++)
    {
        cout << h_out[i] << " ";
    }

    cudaFree(p_d);
    free(h_in);
    free(h_out);

    return 0;
}

在这里插入图片描述

6 !!! 数组求和

在这里插入图片描述

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <iostream>
#include <curand.h>
#include <curand_kernel.h>


using namespace std;
#define ARRAY_SIZE 16
#define ARRAY_LENGTH 256
#define BLOCK_SIZE 16



//GPU上初始化,N是
__global__ void gpu_initial(float* a, int N) {

    int x = threadIdx.x + blockDim.x * blockIdx.x;
    curandState state;
    long seed = N;
    curand_init(seed, x, 0, &state);
    if (x < N) a[x] = curand_uniform(&state);
}


void show(float* a, int m, int n)
{
    for (int i = 0; i < m; i++)
    {
        for (int j = 0; j < n; j++)
        {
            //  a[i] = add_one(a[i]);
            cout << a[i * m + j] << "      ";
        }
        cout << endl;
    }
}


__global__ void reduce_kernel (float *data_in, float *data_out)
{
    //现在是N个1维的block
    int threadIDinAllBlock = threadIdx.x + blockDim.x*blockIdx.x;  
    int threadIDinOneBlock = threadIdx.x;
    int blockID = blockIdx.x;
    extern __shared__ float s_d[];

    //共享内存是在一个blcok,但是传入的数据是分散到1024个block上
    s_d[threadIDinOneBlock] = data_in[threadIDinAllBlock];
    __syncthreads();

    //for (int i = 0; i < N; i++)
    //{
    //    printf("here s_d[%d] = %f \n",i, s_d[i]);
    //}
    //printf(" \n");

    //对半求和，前一半数据的第一位数和后半数据的第一位数相加，。。。
    //这里是算完一个block
    for (int s = blockDim.x/2; s > 0; s >>= 1)  //s >>= 1 是除以2的意思,数据依次是1024,512,256.。。。。
    {
        if (threadIDinOneBlock < s)
        {
            s_d[threadIDinOneBlock] += s_d[threadIDinOneBlock + s];
        }
        __syncthreads();
    }
    
    //输出1024个blcok的各自的结果
    if (threadIDinOneBlock == 0)
    {
        data_out[blockID] = s_d[0];
        /*printf("here s_d[0] = %f \n", s_d[0]);*/
        printf("here  data_out[%d] = %f \n", blockID, data_out[blockID]);
    }
}

void cpu_reduce(float *d_in, float* d_mid, float* d_out, int N)
{
    int threadNum = N;
    int blocks = N / threadNum;
    //printf("blocks   = %d \n", blocks);

    reduce_kernel << <blocks, threadNum, threadNum * sizeof(float) >> > (d_in, d_mid);
    reduce_kernel << <1, threadNum, threadNum * sizeof(float) >> > (d_mid, d_out);
}


int main()
{
    int m = 50;
    int N = m*m;
    float* d_in,*d_mid, *d_out, * h_in, * h_out;

    //cpu上开辟空间
    h_in = (float*)malloc(N * sizeof(float));
    h_out = (float*)malloc(N * sizeof(float));

    //GPU上开辟空间
    cudaMalloc((void**)&d_in, N * sizeof(float)); //将指针指向GPU的一个内存地址
    cudaMalloc((void**)&d_mid, m * sizeof(float));//一个block的大小
    cudaMalloc((void**)&d_out, N * sizeof(float));

    gpu_initial << <m, m >> > (d_in, N);  //直接在GPU上初始化
    cudaMemcpy(h_in, d_in, N * sizeof(float), cudaMemcpyDeviceToHost);    //将数据从GPU拷贝到CPU,同时指定拷贝的长度
    clock_t t_start_cpu = clock();
    float sum = 0;
    printf("in mian  \n");
    for (int i = 0; i < N; i++)
    {
        sum += h_in[i];
        //cout << h_in[i] << " ";
    }
    cout << " \n";
    cout << "sum_ cpu = " << sum << endl;
    cout << " time in cpu = " << (double)(clock() - t_start_cpu) / CLOCKS_PER_SEC << endl;
    cout << " \n";


    cout << "----------------------- \n";
    clock_t t_start_gpu = clock();
    cpu_reduce(d_in, d_mid,d_out,N);
    cudaDeviceSynchronize();
    cudaMemcpy(h_out, d_out,N * sizeof(float), cudaMemcpyDeviceToHost);
    cout << " SUM gpu : h_out = " << h_out[0] << endl;
    cout << " time in GPU = " << (double)(clock() - t_start_gpu) / CLOCKS_PER_SEC << endl;

    cudaFree(d_in);
    cudaFree(d_out);
    cudaFree(d_mid);
    free(h_in);
    free(h_out);

    return 0;
}

本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若转载，请注明出处：/a/750444.html

如若内容造成侵权/违法违规/事实不符，请联系我们进行投诉反馈qq邮箱809451989@qq.com，一经查实，立即删除！

相关文章

【软件测试】白盒测试与接口测试详解

【软件测试】白盒测试与接口测试详解

&#x1f345; 视频学习&#xff1a;文末有免费的配套视频可观看 &#x1f345; 点击文末小卡片&#xff0c;免费获取软件测试全套资料&#xff0c;资料在手&#xff0c;涨薪更快 一、什么是白盒测试 白盒测试是一种测试策略&#xff0c;这种策略允许我们检查程序的内部结构&a…
阅读更多...
短视频带货实战营(高阶课)，从0到1做个赚钱的抖音号（17节课）

短视频带货实战营(高阶课)，从0到1做个赚钱的抖音号（17节课）

课程目录&#xff1a; 1-短视频带贷先导课_1.mp4 2-账号搭建_1.mp4 3-账号养号涨粉套路_1.mp4 4-开通橱窗_1.mp4 5-管家式选品_1.mp4 6-六个能赚钱的赛道_1.mp4 7-选品之精选联盟_1.mp4 8-好物分享的三种形式_1.mp4 9-短视频之图文课_1.mp4 10-短视频之剪辑课_1.mp4 …
阅读更多...
el-upload+python fastAPI实现上传文件

el-upload+python fastAPI实现上传文件

el-upload通过action指定后端接口&#xff0c;并通过name指定传输的文件包裹在什么变量名中 <el-uploadclass"upload-demo"dragaction"https://ai.zscampus.com/toy/upload"multiplename"fileList":limit"10"accept".xlsx, .x…
阅读更多...
昇思25天学习打卡营第9天|使用静态图加速

昇思25天学习打卡营第9天|使用静态图加速

一、简介&#xff1a; AI编译框架分为两种运行模式&#xff0c;分别是动态图模式以及静态图模式。MindSpore默认情况下是以动态图模式运行&#xff0c;但也支持手工切换为静态图模式。两种运行模式的详细介绍如下&#xff1a; &#xff08;1&#xff09;动态图&#xff1a; …
阅读更多...
如何使用Hugging Face Transformers为情绪分析微调BERT？

如何使用Hugging Face Transformers为情绪分析微调BERT？

情绪分析指用于判断文本中表达的情绪的自然语言处理(NLP)技术&#xff0c;它是客户反馈评估、社交媒体情绪跟踪和市场研究等现代应用背后的一项重要技术。情绪可以帮助企业及其他组织评估公众意见、提供改进的客户服务&#xff0c;并丰富产品或服务。 BERT的全称是来自Transfo…
阅读更多...
The First Descendant第一后裔联机失败、联机报错这样处理

The First Descendant第一后裔联机失败、联机报错这样处理

第一后裔/The First Descendant是一款免费的多人合作射击游戏&#xff0c;玩家将进入一片混乱的英格里斯大陆&#xff0c;扮演继承者后裔&#xff0c;通过各种主支线任务和故事剧情触发&#xff0c;最终揭开自身的秘密&#xff0c;并带领大家一起抵抗邪恶势力的入侵。为了避免玩…
阅读更多...
【Java Web】三大域对象

【Java Web】三大域对象

目录 一、域对象概述 二、三大域对象 三、域对象使用相关API 一、域对象概述 一些可用于存储数据和传递数据的对象被称为域对象&#xff0c;根据传递数据范围的不同&#xff0c;我们称之为不同的域&#xff0c;不同的域对象代表不同的域&#xff0c;共享数据的范围也不同。 二、…
阅读更多...
ISP IC/FPGA设计-第一部分-SC130GS摄像头分析-IIC通信（1）

ISP IC/FPGA设计-第一部分-SC130GS摄像头分析-IIC通信（1）

1.摄像头模组 SC130GS通过一个引脚&#xff08;SPI_I2C_MODE&#xff09;选择使用IIC或SPI配置接口&#xff0c;通过查看摄像头模组的原理图&#xff0c;可知是使用IIC接口&#xff1b; 通过手册可知IIC设备地址通过一个引脚控制&#xff0c;查看摄像头模组的原理图&#xff…
阅读更多...
数据库调优厂商 OtterTune 宣布停止运营

数据库调优厂商 OtterTune 宣布停止运营

昨天刷到消息&#xff0c;得知数据库优化厂商 OtterTune 停止了运营。OtterTune 的成员主要来自 CMU Andy Pavlo 教授领导的数据库实验室。公司正式成立于 2021 年 5 月&#xff0c;融资了 1450 万美金。 按照 Andy 教授的说法&#xff0c;公司是被一个收购 offer 搞砸了。同时…
阅读更多...
npm-check【实用教程】升级项目中的依赖

npm-check【实用教程】升级项目中的依赖

安装 npm-check npm i -g npm-check检查项目中的依赖 npm-check会显示项目中没有使用&#xff0c;以及有新版本的依赖 升级项目中的依赖 npm-check -u方向键上下可以移动图中左侧的箭头空格键可选中/取消选中标注为 Major Update 和 Non-semver 类的版本&#xff0c;需去官网查…
阅读更多...
用MySQL和navicatpremium做一个项目—（财务管理系统）。

用MySQL和navicatpremium做一个项目—（财务管理系统）。

1 ER图缩小的话怕你们看不清&#xff0c;所以截了两张图 2 vsdx绘图结果 3DDL和DML,都有点长分了好多次上传&#xff0c;慢慢看 DDL -- 用户表 CREATE TABLE users (user_id INT AUTO_INCREMENT PRIMARY KEY COMMENT 用户ID,username VARCHAR(50) NOT NULL UNIQUE COMMENT 用…
阅读更多...
奔驰汽车的通信如此固若金汤的原因

奔驰汽车的通信如此固若金汤的原因

随着摄像系统、距离控制、航线保持等功能以及制动辅助系统、制动力分配系统、车身侧倾干预与缓解系统等功能的飞速发展，汽车的系统功能之间已经不再独立，而是呈现互相合作的关系，各功能之间的无缝集成更是各大整车厂追求的目标。俗话说，外练筋骨皮，内练一口气，有了各式安…
阅读更多...
第4章 客户端-客户端案例分析

第4章 客户端-客户端案例分析

1 Redis内存陡增 1.1.现象 服务端现象&#xff1a;Redis主节点内存陡增&#xff0c;几乎用满maxmemory&#xff0c;而从节点内存并没有变化&#xff08;正常情况下主从节点内存使用量基本相同&#xff09;。 客户端现象&#xff1a;客户端产生了OOM异常&#xff0c;也就是Redis…
阅读更多...
桃园三结义 | 第1集 | 三人一条心，黄土变成金，有你带着俺，大事定能成功啊！| 正所谓择木之禽，得其良木，择主之臣，得遇明主 | 三国演义 | 群雄逐鹿

桃园三结义 | 第1集 | 三人一条心，黄土变成金，有你带着俺，大事定能成功啊！| 正所谓择木之禽，得其良木，择主之臣，得遇明主 | 三国演义 | 群雄逐鹿

&#x1f64b;大家好&#xff01;我是毛毛张! &#x1f308;个人首页&#xff1a; 神马都会亿点点的毛毛张 &#x1f4cc;这篇博客是毛毛张结合三国演义原著分享三国演义文学剧本中的经典台词和语句&#xff0c;本篇分享的是《三国演义》第Ⅰ部分《群雄逐鹿》的第1️⃣集《桃…
阅读更多...
eNSP中VRRP的配置和使用

eNSP中VRRP的配置和使用

一、基础配置 1.新建拓扑图 2.配置vlan a.CORE-S1 <Huawei>system-view [Huawei]sysname CORE-S1 [CORE-S1]vlan 10 [CORE-S1-vlan10]vlan 20 [CORE-S1-vlan20]vlan 30 b.CORE-S2 <Huawei>system-view [Huawei]sysname CORE-S2 [CORE-S2]vlan 10 [CORE…
阅读更多...
2024年6月27日 (周四) 叶子游戏新闻

2024年6月27日 (周四) 叶子游戏新闻

老板键工具来唤去: 它可以为常用程序自定义快捷键&#xff0c;实现一键唤起、一键隐藏的 Windows 工具&#xff0c;并且支持窗口动态绑定快捷键&#xff08;无需设置自动实现&#xff09;。 喜马拉雅下载工具: 字面意思 Steam国产“类8番”游戏《永恒逃脱&#xff1a;暗影城堡》…
阅读更多...
软件协同开发是一种通过团队合作来创建软件的开发方法

软件协同开发是一种通过团队合作来创建软件的开发方法

软件协同开发是一种通过团队合作来创建软件的开发方法。与传统的瀑布模型相比&#xff0c;软件协同开发强调团队成员之间的合作和沟通&#xff0c;以实现更高效的开发过程和更优质的软件产品。 在软件协同开发中&#xff0c;团队成员通过一系列工具和技术来协同工作。这些工具…
阅读更多...
视觉灵感的探索和分享平台

视觉灵感的探索和分享平台

做设计没灵感&#xff1f;大脑一片空白&#xff1f;灵感是创作的源泉&#xff0c;也是作品的灵魂所在。工作中缺少灵感&#xff0c;这是每个设计师都会经历的苦恼&#xff0c;那当我们灵感匮乏的时候&#xff0c;该怎么办呢&#xff1f;别急&#xff0c;即时设计、SurfCG、Lapa…
阅读更多...
一种PCB外壳设计方法的尝试

一种PCB外壳设计方法的尝试

一个异性PCB的板框&#xff0c;外壳&#xff0c;PCB设计&#xff1a; 正常情况下先由机械工程师用CAD设计出板框导出DXF文件&#xff0c;之后基于此DXF文件作为板框进行PCB设计和外壳设计&#xff0c;但对硬件工程师来讲有时候直接在PCB软件上进行简单的板框设计显得更方便&am…
阅读更多...
apktool反编译apk工具

apktool反编译apk工具

Android apk安装包反编译——apktool工具-CSDN博客 Android 如何反编译APK获取源码_android studio apk反源码-CSDN博客
阅读更多...
最新文章

Copyright @ 2022~2023