cuda编程实例

矢量求和

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

#define N   100

__global__ void add( int *a, int *b, int *c ) {
    int tid = blockIdx.x;    // this thread handles the data at its thread id
    if (tid < N)
        c[tid] = a[tid] + b[tid];
}

int main( void ) {
    int a[N], b[N], c[N];
    int *dev_a, *dev_b, *dev_c;

    // allocate the memory on the GPU
    CHECK( cudaMalloc( (void**)&dev_a, N * sizeof(int) ) );
    CHECK( cudaMalloc( (void**)&dev_b, N * sizeof(int) ) );
    CHECK( cudaMalloc( (void**)&dev_c, N * sizeof(int) ) );

    // fill the arrays 'a' and 'b' on the CPU
    for (int i=0; i<N; i++) {
        a[i] = -i;
        b[i] = i * i;
    }

    // copy the arrays 'a' and 'b' to the GPU
    CHECK( cudaMemcpy( dev_a, a, N * sizeof(int),
                              cudaMemcpyHostToDevice ) );
    CHECK( cudaMemcpy( dev_b, b, N * sizeof(int),
                              cudaMemcpyHostToDevice ) );

    add<<<N,1>>>( dev_a, dev_b, dev_c );

    // copy the array 'c' back from the GPU to the CPU
    CHECK( cudaMemcpy( c, dev_c, N * sizeof(int),
                              cudaMemcpyDeviceToHost ) );

    // display the results
    for (int i=0; i<N; i++) {
        printf( "%d + %d = %d\n", a[i], b[i], c[i] );
    }

    // free the memory allocated on the GPU
    CHECK( cudaFree( dev_a ) );
    CHECK( cudaFree( dev_b ) );
    CHECK( cudaFree( dev_c ) );

    return 0;
}
                     
                                                     

julia集

在这里插入图片描述

在这里插入代码片

任意长度矢量求和

在这里插入图片描述

#include "error.cuh"
#include <stdio.h>
#include <cuda_runtime.h>
#define N   (33 * 1024)

__global__ void add( int *a, int *b, int *c ) {

//获取线程id,网格大小1D,线程块大小1D,gridDim.x=128,blockDim.x=128
//数据量:33792,线程数:16384

        int tid = threadIdx.x + blockIdx.x * blockDim.x;

//方法1:
        for(int i=0;i<N;i++){
                if(tid==i%blockDim.x*gridDim.x){
                        c[i] = a[i] + b[i];

                }
        }
/*

//方法2:
        while (tid < N) {
                c[tid] = a[tid] + b[tid];
                tid += blockDim.x * gridDim.x;
        }
*/
}
int main( void ) {
    //定义主机指针和设备指针
    int *a, *b, *c;
    int *dev_a, *dev_b, *dev_c;

    //CPU上动态分配内存
    a = (int*)malloc( N * sizeof(int) );
    b = (int*)malloc( N * sizeof(int) );
    c = (int*)malloc( N * sizeof(int) );

    //GPU上动态分配内存
    CHECK( cudaMalloc( (void**)&dev_a, N * sizeof(int) ) );
    CHECK( cudaMalloc( (void**)&dev_b, N * sizeof(int) ) );
    CHECK( cudaMalloc( (void**)&dev_c, N * sizeof(int) ) );

    //CPU上为a和b赋值
    for (int i=0; i<N; i++) {
        a[i] = i;
        b[i] = 2 * i;
    }

    //将a和b的数据从主机拷贝到设备
    CHECK( cudaMemcpy( dev_a, a, N * sizeof(int),
                              cudaMemcpyHostToDevice ) );
    CHECK( cudaMemcpy( dev_b, b, N * sizeof(int),
                              cudaMemcpyHostToDevice ) );
    //调用核函数,设置线程,网格大小1D,线程块大小1D,32
    add<<<128,128>>>( dev_a, dev_b, dev_c );

    //将数据从设备传递到主机上
    CHECK( cudaMemcpy( c, dev_c, N * sizeof(int),
                              cudaMemcpyDeviceToHost ) );

    //结果校验
    bool success = true;
    for (int i=0; i<N; i++) {
        if ((a[i] + b[i]) != c[i]) {
            success = false;
        }
    }
    if (success)
        printf( "We did it!\n" );
    else
        printf("Not Pass!Error!!!\n");

    //释放GPU内存空间
    CHECK( cudaFree( dev_a ) );
    CHECK( cudaFree( dev_b ) );
    CHECK( cudaFree( dev_c ) );

    //释放CPU内存空间
    free( a );
    free( b );
    free( c );

    return 0;
}

矩阵处理

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

#define N 1000
#define BLOCK_SIZE 32

__managed__ int input_Matrix[N][N];
__managed__ int output_GPU[N][N];
__managed__ int output_CPU[N][N];
__managed__ int output_GPU[N][N];
__managed__ int output_CPU[N][N];
__global__ void kernel(int input_M[N][N], int output_M[N][N])
{

  int x = blockIdx.x * blockDim.x + threadIdx.x;//<<gradx,blockx>>
  int y = blockIdx.y * blockDim.y + threadIdx.y;//<grady,blocky>>
//x维度变化快相当于内层循环
  if(x<N && y<N)
  {
      if(x%2==0 && y%2==0)
      {
          output_M[y][x] = input_M[y][x]*input_M[y][x];
      }
      else
      {
          output_M[y][x] = input_M[y][x]-1;
      }
  }
}
void cpu_kernel(int intput_M[N][N], int output_CPU[N][N])
{
    for(int i=0; i<N; i++)
     {
        for(int j=0; j<N; j++)
        {
            if(j%2==0 && i%2==0)
            {
                output_CPU[i][j] = intput_M[i][j]*intput_M[i][j];
            }
            else
            {
                output_CPU[i][j] = intput_M[i][j]-1;
            }
        }
    }
}

int main(int argc, char const *argv[])
{

    for (int i = 0; i < N; ++i) {
        for (int j = 0; j < N; ++j)
        {

            input_Matrix[i][j] = rand()%3001;
        }
    }

//CPU
    cudaEvent_t start_cpu,stop_cpu;
    CHECK(cudaEventCreate(&start_cpu));
    CHECK(cudaEventCreate(&start_cpu));
    CHECK(cudaEventCreate(&stop_cpu));
    CHECK(cudaEventRecord(start_cpu));
    printf("\n***********CPU RUN**************\n");
    cpu_kernel(input_Matrix, output_CPU);
    CHECK(cudaEventRecord(stop_cpu));
    CHECK(cudaEventSynchronize(stop_cpu));
    float elapsed_time_cpu;
    CHECK(cudaEventElapsedTime(&elapsed_time_cpu, start_cpu, stop_cpu));
    printf("Time_CPU = %g ms.\n", elapsed_time_cpu);
    CHECK(cudaEventDestroy(start_cpu));
    CHECK(cudaEventDestroy(stop_cpu));

//GPU
    cudaEvent_t start,stop_gpu;
    CHECK(cudaEventCreate(&start));
    CHECK(cudaEventCreate(&stop_gpu));
    CHECK(cudaEventRecord(start));

    unsigned int grid_rows = (N + BLOCK_SIZE - 1) / BLOCK_SIZE;
    unsigned int grid_cols = (N + BLOCK_SIZE - 1) / BLOCK_SIZE;

    dim3 dimGrid(grid_cols, grid_rows);//网格大小,保证gridDim.x/y*blockDim.y,可以被32整除,且大于数组(x,y)的
数据大小
    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
    printf("\n***********GPU RUN**************\n");
    kernel<<<dimGrid, dimBlock>>>(input_Matrix, output_GPU);

    CHECK(cudaDeviceSynchronize());
    CHECK(cudaEventRecord(stop_gpu));
    CHECK(cudaEventSynchronize(stop_gpu));

    float elapsed_time_gpu;
    CHECK(cudaEventElapsedTime(&elapsed_time_gpu, start, stop_gpu));
    printf("Time_GPU = %g ms.\n", elapsed_time_gpu);
    CHECK(cudaEventDestroy(start));
    CHECK(cudaEventDestroy(stop_gpu));
    //校验
    printf("\n***********Check result**************\n");
    int ok=1;
    for (int i = 0; i < N; ++i)
    {
        for (int j = 0; j < N; ++j)
        {
            if(fabs(output_GPU[i][j] - output_CPU[i][j])>(1.0e-10))
            {
                ok = 0;
            }

        }
    }


    if(ok)
    {
        printf("Pass!!!\n");
    }
    else
    {
        printf("Error!!!\n");
    }

    // free memory
    return 0;
}

矩阵转置

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

点积

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

常量内存

在这里插入图片描述

CUDA 编程: GPU 高性能并行计算经典教材》是一本广泛被认可的权威教材,涵盖了 CUDA 编程的重要概念、原理和技术。该书针对使用 NVIDIA GPU 进行并行计算的开发人员和学习者提供了深入的指导和实践经验。 首先,该教材从基础开始,介绍了 CUDA 编程的基本概念和体系结构,包括并行计算模型、线程和块的概念以及内存模型等。这为读者打下了坚实的基础,使他们能够理解并应用 CUDA 的核心原理。 其次,该教材详细介绍了 CUDA 编程的各个环节,包括在 GPU 上设置和管理线程、优化并行性能以及处理复杂数据结构等。通过深入讲解 CUDA 编程中的关键技术和技巧,读者能够更好地理解如何有效地利用 GPU 并行计算的优势。 此外,该教材还提供了丰富的实例和案例,涵盖了广泛的应用领域,如图像处理、机器学习和科学计算等。这些实例不仅有助于读者将理论知识转化为实际应用,还展示了 CUDA 编程在不同领域中的灵活性和高效性。 最后,该教材还介绍了 CUDA 并行计算的最新发展和趋势,如 CUDA C++、GPU 云计算和深度学习等。通过了解这些最新进展,读者能够跟随技术的发展潮流,进一步拓展自己的知识和技能。 综上所述,《CUDA 编程: GPU 高性能并行计算经典教材》作为一本经典教材,全面而详细地介绍了 CUDA 编程的各个方面,既适合初学者入门,也适合有一定经验的开发人员进一步提升技术水平。无论是对于学习 CUDA 编程的人员,还是对于进行 GPU 高性能并行计算的开发者来说,这本教材都是一本不可多得的宝贵资料。

“相关推荐”对你有帮助么?

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论 1
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值