cuda 学习笔记4

news2024/11/25 4:32:19

一 基本函数

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

在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;
}

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

如若内容造成侵权/违法违规/事实不符，请联系多彩编程网进行投诉反馈，一经查实，立即删除！

相关文章

关于0xc000007b的一种解决方案

关于0xc000007b的一种解决方案

今天我在安装qview并运行时时&#xff0c;遇到了这个问题。 我在网上查找了许多解决方案&#xff0c;但它们大多都说是某些dll缺失或错误引起的。 这些说法应该是正确的&#xff0c;但我用了dll修复工具后&#xff0c;一点用都没有。 后来捣鼓半天后&#xff0c;我发现很可能…
阅读更多...
Golang 百题(实战快速掌握语法)_2

Golang 百题(实战快速掌握语法)_2

返回集合中满足指定条件的最后一个元素 本实验将实现判断给定集合中的元素是否符合&#xff0c;并返回符合的最后一个元素。 知识点 forfmt.Error 适合人群 本课程属于基础课程。需要用户掌握 Go 语言编程基础知识、计算机基础知识和 Linux 环境的基本用法。 许可证 内容…
阅读更多...
【可控图像生成系列论文（五）】ControlNet 和 IP-Adapter 之间的区别有哪些？

【可控图像生成系列论文（五）】ControlNet 和 IP-Adapter 之间的区别有哪些？

系列文章目录 【可控图像生成系列论文&#xff08;一&#xff09;】 简要介绍了 MimicBrush 的整体流程和方法&#xff1b;【可控图像生成系列论文&#xff08;二&#xff09;】 就MimicBrush 的具体模型结构、训练数据和纹理迁移进行了更详细的介绍。【可控图像生成系列论文&…
阅读更多...
MySQL高级-索引-设计原则小结

MySQL高级-索引-设计原则小结

文章目录 1、设计原则2、索引小结2.1、索引概述2.2、索引结构2.3、索引分类2.4、索引语法2.5、SQL性能分析2.6、索引使用2.7、索引设计原则 1、设计原则 针对于数据量较大&#xff0c;且查询比较频繁的表建立索引。针对于常作为查询条件&#xff08;where&#xff09;、排序&am…
阅读更多...
东软睿驰总裁兼CTO杜强受邀出席 CICV 2024智能网联汽车技术首脑（CTO）闭门峰会

东软睿驰总裁兼CTO杜强受邀出席 CICV 2024智能网联汽车技术首脑（CTO）闭门峰会

近日&#xff0c;第十一届国际智能网联汽车技术年会&#xff08;CICV 2024&#xff09;在北京举办&#xff0c;会议期间组织智能网联汽车技术首脑&#xff08;CTO&#xff09;闭门峰会&#xff0c;邀请40余位技术领袖围绕智能网联汽车产业生态建设以及智能网联汽车数据、算力和…
阅读更多...
使用k8s变更线上版本号

使用k8s变更线上版本号

第一步&#xff0c;在镜像仓库中找到历史版本号 第二步&#xff0c;在rancher中在工作负载里 第三步&#xff0c;在rancher找到这个版本号&#xff0c;可以更改之前的版本号 这样就可以很方便的退回到之前的版本了
阅读更多...
【技巧】如何检查多个GPU之间是否支持P2P通信

【技巧】如何检查多个GPU之间是否支持P2P通信

转载请注明出处&#xff1a;小锋学长生活大爆炸[xfxuezhagn.cn] 如果本文帮助到了你&#xff0c;欢迎[点赞、收藏、关注]哦~ 需要用到cuda_samples&#xff1a;GitHub - NVIDIA/cuda-samples 该工具的详细解释可以看这个&#xff1a; 【知识】详细介绍 CUDA Samples 示例工程…
阅读更多...
时间序列分析入门：概念、模型与应用【ARMA、ARIMA模型】

时间序列分析入门：概念、模型与应用【ARMA、ARIMA模型】

在这篇博客中&#xff0c;我们将全面探讨时间序列分析的基本概念和分类&#xff0c;深入理解平稳性及其检验方法&#xff0c;并介绍自回归模型&#xff08;AR&#xff09;、滑动平均模型&#xff08;MA&#xff09;、自回归滑动平均模型&#xff08;ARMA&#xff09;以及自回归…
阅读更多...
吐血推荐！3款视频生成工具，全部国产，都免费

吐血推荐！3款视频生成工具，全部国产，都免费

AI视频大模型的爆发&#xff0c;让创作爆款视频不再是专业人士的能力。 今天二师兄给大家推荐3款免费的视频生成工具。 01 可灵 推荐指数 &#xff1a; 五颗星 先看效果 可灵大模型测试 可灵大模型是快手AI团队自主研发的视频生成大模型&#xff0c;具备强大的视频创作能力&a…
阅读更多...
FlinkX学习

FlinkX学习

FlinkX学习 FlinkX安装 由于flinkx已经改名chunjun 官网已不存在 (https://gitee.com/lugela/flinkx#flinkx)这里可以看到flinkx的操作文档 1、上传并解压 unzip flinkx-1.10.zip -d /usr/local/soft/2、配置环境变量 FLINKX_HOME/usr/local/soft/flinkx-1.10 export PATH$F…
阅读更多...
PHP-CGI的漏洞（CVE-2024-4577）

PHP-CGI的漏洞（CVE-2024-4577）

通过前两篇文章的铺垫&#xff0c;现在我们可以了解 CVE-2024-4577这个漏洞的原理 漏洞原理 CVE-2024-4577是CVE-2012-1823这个老漏洞的绕过&#xff0c;php cgi的老漏洞至今已经12年&#xff0c;具体可以参考我的另一个文档 简单来说&#xff0c;就是使用cgi模式运行的PHP&…
阅读更多...
零拷贝技术（zero copy），DMA，mmap，sendfile

零拷贝技术（zero copy），DMA，mmap，sendfile

在一些高性能的IO场景下我们经常能听到零拷贝技术&#xff0c;这是个不错的话题。 零拷贝指的是内核态与用户态之间的数据拷贝&#xff0c;而这两个区域的数据拷贝只能依靠CPU&#xff0c;但是CPU最重要的作用应该是运算。 一、DMA的由来 在没有DMA之前&#xff0c;磁盘的IO…
阅读更多...
【NPS】哑终端设备如何实现域VLAN动态分配

【NPS】哑终端设备如何实现域VLAN动态分配

在【NPS】微软NPS配置802.1x&#xff0c;验证域账号&#xff0c;动态分配VLAN&#xff08;有线网络续篇&#xff09;中&#xff0c;已经通过C3PL策略配置实现了802.1x验证没有通过时&#xff0c;自动分配一个Guest VLAN&#xff0c;以确保用户至少能够访问基本的网络服务。问题…
阅读更多...
数字时代的文化革命：Facebook的社会影响

数字时代的文化革命：Facebook的社会影响

随着数字技术的飞速发展和互联网的普及&#xff0c;社交网络如今已成为人们日常生活中不可或缺的一部分。在众多社交平台中&#xff0c;Facebook作为最大的社交网络之一&#xff0c;不仅连接了全球数十亿用户&#xff0c;更深刻影响了人们的社会互动方式、文化认同和信息传播模…
阅读更多...
展开说说：Android列表之RecyclerView

展开说说：Android列表之RecyclerView

RecyclerView 它是从Android5.0出现的全新列表组件&#xff0c;更加强大和灵活。用于显示列表形式 (list) 或者网格形式 (grid) 的数据&#xff0c;替代ListView和GridView成为Android主流的列表组件。可以说Android客户端只要有表格的地方就有RecyclerView。 RecyclerView 内…
阅读更多...
【linux】使用vnc连接远程桌面，需要安装tigervnc，并在服务端期待，然后在客户端使用tigervnc-viewer进行连接即可

【linux】使用vnc连接远程桌面，需要安装tigervnc，并在服务端期待，然后在客户端使用tigervnc-viewer进行连接即可

vnc 远程设置方法 需要服务端安装软件&#xff1a; sudo apt install -y tigervnc-standalone-server# 先配置密码使用&#xff1a; tightvncpasswd启动服务&#xff0c;禁用本机 vncserver -localhost no -geometry 1924x1080 :1客户端安装软件&#xff1a; sudo apt insta…
阅读更多...
JavaScript高级程序设计(第四版)--学习记录之基本引用类型

JavaScript高级程序设计(第四版)--学习记录之基本引用类型

Date Date类型将日期保存为自协调世界时间1970年1月1日午夜至今所经过的毫秒数。 创建日期对象 let now new Date() Date.parse()方法接收一个表示日期的字符串参数&#xff0c;尝试将这个字符串转换为表示该日期的毫秒数。 let time new Date(Date.parse("May 24,2024&…
阅读更多...
Jmeter+InfluxDB+Grafana性能测试数据展示

Jmeter+InfluxDB+Grafana性能测试数据展示

JmeterInfluxDBGrafana提供了一种更好的对Jmeter压测结果的实时监控展示。可以理解为数据源产生的数据加上时间记录并存储&#xff0c;然后使用各种开源图表组件进行展示。实现jmeter报告的更好的可视化展示 1&#xff09;方便测试结果数据落地以及更好的分析 2&#xff09;将…
阅读更多...
超好用的思维导图—万兴亿图脑图 v10解锁版安装教程 (思维导图软件和头脑风暴工具)

超好用的思维导图—万兴亿图脑图 v10解锁版安装教程 (思维导图软件和头脑风暴工具)

前言 万兴亿图脑图 (Wondershare EdrawMind) 是一款多平台协作思维导图软件和头脑风暴工具,亿图思维导图提供丰富的布局,样式,主题及配色方案,集成拥有数万幅原创思维导图作品的思维导图社区,涵盖教育,职场,自我提升等各大领域精华知识.支持会议演示,多端创作,云端存储,导图分…
阅读更多...
BioCLIP：物种图像的基础视觉模型

BioCLIP：物种图像的基础视觉模型

从无人机到个人手机&#xff0c;各种相机收集的自然世界图像是越来越丰富的生物信息来源。从图像中提取生物相关信息用于科学的计算方法和工具激增&#xff0c;尤其是计算机视觉。然而&#xff0c;其中大多数都是为特定任务设计的&#xff0c;不容易适应或扩展到新的问题、环境…
阅读更多...
推荐文章
最新文章

Copyright @ 2022~2023