CUDA编程：从hello world开始

已于 2024-01-09 18:00:50 修改
阅读量937 收藏 9
点赞数 10
分类专栏： cuda编程 文章标签： 深度学习 人工智能 c语言
于 2024-01-03 22:00:06 首次发布
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/qq_43570025/article/details/135374018
版权

目录

第一个程序：用GPU输出Hello World

写一个CUDA C程序步骤

hello.cu

编译

运行

common.h

cuda程序执行步骤

两个数组相加

CPU上的两个数据相加

GPU上的两个数据相加

计时

nvprof

第一个程序：用GPU输出Hello World

写一个CUDA C程序步骤

  1. 用专用扩展名.cu来创建一个源文件

  2. 使用CUDA nvcc编译器来编译程序

  3. 从命令行运行可执行文件，这个文件有可在GPU上运行的内核代码

hello.cu

#include "../common/common.h"
#include <stdio.h>

/*
 * A simple introduction to programming in CUDA. This program prints "Hello
 * World from GPU! from 10 CUDA threads running on the GPU.
 */

__global__ void helloFromGPU()
{
    printf("Hello World from GPU!\n");
}

int main(int argc, char **argv)
{
    printf("Hello World from CPU!\n");

    helloFromGPU<<<1, 10>>>();
    CHECK(cudaDeviceReset());
    return 0;
}

编译

nvcc hello.cu -o hello

运行

./hello

cudaDeviceReset()函数是用来显示的释放和清空当前进程中于当前设备有关的所有资源当重置函数移除后 编译运行则只输出hello world from cpu,printf在gpu上被调用，cudaDeviceReset()函数使这些来自gpu的输出发送到主机，然后在控制台输出。没有调用cudaDeviceReset()函数就不能保证这些可以被显示

关于cudaDeviceReset(); 这句话如果没有，则不能正常的运行，因为这句话包含了隐式同步，GPU和CPU执行程序是异步的，核函数调用后成立刻会到主机线程继续，而不管GPU端核函数是否执行完毕，所以上面的程序就是GPU刚开始执行，CPU已经退出程序了，所以我们要等GPU执行完了，再退出主机线程。

common.h

#include <sys/time.h>
#include <cuda_runtime.h>
#include <stdio.h>

#ifndef _COMMON_H
#define _COMMON_H

cudaError_t ErrorCheck(cudaError_t status,const char* filename,int lineNumber)
{
    if(status != cudaSuccess)
    {
        printf("CUDA API error: \r\ncode=%d, name=%s, description=%s\r\nfile=%s, line=%d\r\n",
                status,cudaGetErrorName(status),cudaGetErrorString(status),filename,lineNumber
        );
    }
    return status;
}

#define CHECK(call)                                                            \
{                                                                              \
    const cudaError_t error = call;                                            \
    if (error != cudaSuccess)                                                  \
    {                                                                          \
        fprintf(stderr, "Error: %s:%d, ", __FILE__, __LINE__);                 \
        fprintf(stderr, "code: %d, reason: %s\n", error,                       \
                cudaGetErrorString(error));                                    \
        exit(1);                                                               \
    }                                                                          \
}


#define CHECK_CUBLAS(call)                                                     \
{                                                                              \
    cublasStatus_t err;                                                        \
    if ((err = (call)) != CUBLAS_STATUS_SUCCESS)                               \
    {                                                                          \
        fprintf(stderr, "Got CUBLAS error %d at %s:%d\n", err, __FILE__,       \
                __LINE__);                                                     \
        exit(1);                                                               \
    }                                                                          \
}

#define CHECK_CURAND(call)                                                     \
{                                                                              \
    curandStatus_t err;                                                        \
    if ((err = (call)) != CURAND_STATUS_SUCCESS)                               \
    {                                                                          \
        fprintf(stderr, "Got CURAND error %d at %s:%d\n", err, __FILE__,       \
                __LINE__);                                                     \
        exit(1);                                                               \
    }                                                                          \
}

#define CHECK_CUFFT(call)                                                      \
{                                                                              \
    cufftResult err;                                                           \
    if ( (err = (call)) != CUFFT_SUCCESS)                                      \
    {                                                                          \
        fprintf(stderr, "Got CUFFT error %d at %s:%d\n", err, __FILE__,        \
                __LINE__);                                                     \
        exit(1);                                                               \
    }                                                                          \
}

#define CHECK_CUSPARSE(call)                                                   \
{                                                                              \
    cusparseStatus_t err;                                                      \
    if ((err = (call)) != CUSPARSE_STATUS_SUCCESS)                             \
    {                                                                          \
        fprintf(stderr, "Got error %d at %s:%d\n", err, __FILE__, __LINE__);   \
        cudaError_t cuda_err = cudaGetLastError();                             \
        if (cuda_err != cudaSuccess)                                           \
        {                                                                      \
            fprintf(stderr, "  CUDA error \"%s\" also detected\n",             \
                    cudaGetErrorString(cuda_err));                             \
        }                                                                      \
        exit(1);                                                               \
    }                                                                          \
}

inline double seconds()
{
    struct timeval tp;
    struct timezone tzp;
    int i = gettimeofday(&tp, &tzp);
    return ((double)tp.tv_sec + (double)tp.tv_usec * 1.e-6);
}

#define CHECK_EXIT(sts, str) \
    if (sts) \
    {       \
        std::cerr << str << " " << __FILE__ << " "  << __LINE__ << std::endl; \
        exit(0);\
    }
#endif // _COMMON_H

cuda程序执行步骤

  1. 设置参与运算的GPU cudaSetDevice

  2. 申请显存 cudaMalloc

  3. 复制数据 cudaMemcpy（参数cudaMemcpyHostToDevice）

  4. 调用kernel函数做运算 xxxKernel

  5. 等待运算结束后同步 cudaDeviceSynchronize（kernel函数的调用是异步的）

  6. 复制运算结果到到内存 cudaMemcpy （参数cudaMemcpyDeviceToHost）

  7. 释放显存 cudaFree

两个数组相加

CPU上的两个数据相加

sumArraysOnHost.c

#include <stdlib.h>
#include <time.h>

/*
 * This example demonstrates a simple vector sum on the host. sumArraysOnHost
 * sequentially iterates through vector elements on the host.
 */

void sumArraysOnHost(float *A, float *B, float *C, const int N)
{
    for (int idx = 0; idx < N; idx++)
    {
        C[idx] = A[idx] + B[idx];
    }

}

void initialData(float *ip, int size)
{
    // generate different seed for random number
    time_t t;
    srand((unsigned) time(&t));

    for (int i = 0; i < size; i++)
    {
        ip[i] = (float)(rand() & 0xFF) / 10.0f;
    }

    return;
}

int main(int argc, char **argv)
{
    int nElem = 1024;
    size_t nBytes = nElem * sizeof(float);

    float *h_A, *h_B, *h_C;
    h_A = (float *)malloc(nBytes);
    h_B = (float *)malloc(nBytes);
    h_C = (float *)malloc(nBytes);

    initialData(h_A, nElem);
    initialData(h_B, nElem);

    sumArraysOnHost(h_A, h_B, h_C, nElem);

    free(h_A);
    free(h_B);
    free(h_C);

    return(0);
}

gcc -O2 -std=c99 -o sumArraysOnHost sumArraysOnHost.c

GPU上的两个数据相加

sumArraysOnGPU-small-case.cu

#include "../common/common.h"
#include <cuda_runtime.h>
#include <stdio.h>

/*
 * This example demonstrates a simple vector sum on the GPU and on the host.
 * sumArraysOnGPU splits the work of the vector sum across CUDA threads on the
 * GPU. Only a single thread block is used in this small case, for simplicity.
 * sumArraysOnHost sequentially iterates through vector elements on the host.
 */

void checkResult(float *hostRef, float *gpuRef, const int N)
{
    double epsilon = 1.0E-8;
    bool match = 1;

    for (int i = 0; i < N; i++)
    {
        if (abs(hostRef[i] - gpuRef[i]) > epsilon)
        {
            match = 0;
            printf("Arrays do not match!\n");
            printf("host %5.2f gpu %5.2f at current %d\n", hostRef[i],
                   gpuRef[i], i);
            break;
        }
    }

    if (match) printf("Arrays match.\n\n");

    return;
}


void initialData(float *ip, int size)
{
    // generate different seed for random number
    time_t t;
    srand((unsigned) time(&t));

    for (int i = 0; i < size; i++)
    {
        ip[i] = (float)(rand() & 0xFF) / 10.0f;
    }

    return;
}


void sumArraysOnHost(float *A, float *B, float *C, const int N)
{
    for (int idx = 0; idx < N; idx++)
        C[idx] = A[idx] + B[idx];
}

__global__ void sumArraysOnGPU(float *A, float *B, float *C, const int N)
{
    int i = threadIdx.x;

    if (i < N) C[i] = A[i] + B[i];
}


int main(int argc, char **argv)
{
    printf("%s Starting...\n", argv[0]);

    // set up device
    int dev = 0;
    CHECK(cudaSetDevice(dev));

    // set up data size of vectors
    int nElem = 1 << 5;
    printf("Vector size %d\n", nElem);

    // malloc host memory
    size_t nBytes = nElem * sizeof(float);

    float *h_A, *h_B, *hostRef, *gpuRef;
    h_A     = (float *)malloc(nBytes);
    h_B     = (float *)malloc(nBytes);
    hostRef = (float *)malloc(nBytes);
    gpuRef  = (float *)malloc(nBytes);

    // initialize data at host side
    initialData(h_A, nElem);
    initialData(h_B, nElem);

    memset(hostRef, 0, nBytes);
    memset(gpuRef,  0, nBytes);

    // malloc device global memory
    float *d_A, *d_B, *d_C;
    CHECK(cudaMalloc((float**)&d_A, nBytes));
    CHECK(cudaMalloc((float**)&d_B, nBytes));
    CHECK(cudaMalloc((float**)&d_C, nBytes));

    // transfer data from host to device
    CHECK(cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_C, gpuRef, nBytes, cudaMemcpyHostToDevice));

    // invoke kernel at host side
    dim3 block (nElem);
    dim3 grid  (1);

    sumArraysOnGPU<<<grid, block>>>(d_A, d_B, d_C, nElem);
    printf("Execution configure <<<%d, %d>>>\n", grid.x, block.x);

    // copy kernel result back to host side
    CHECK(cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost));

    // add vector at host side for result checks
    sumArraysOnHost(h_A, h_B, hostRef, nElem);

    // check device results
    checkResult(hostRef, gpuRef, nElem);

    // free device global memory
    CHECK(cudaFree(d_A));
    CHECK(cudaFree(d_B));
    CHECK(cudaFree(d_C));

    // free host memory
    free(h_A);
    free(h_B);
    free(hostRef);
    free(gpuRef);

    CHECK(cudaDeviceReset());
    return(0);
}

计时

#include "../common/common.h"
#include <cuda_runtime.h>
#include <stdio.h>

/*
 * This example demonstrates a simple vector sum on the GPU and on the host.
 * sumArraysOnGPU splits the work of the vector sum across CUDA threads on the
 * GPU. Only a single thread block is used in this small case, for simplicity.
 * sumArraysOnHost sequentially iterates through vector elements on the host.
 * This version of sumArrays adds host timers to measure GPU and CPU
 * performance.
 */

void checkResult(float *hostRef, float *gpuRef, const int N)
{
    double epsilon = 1.0E-8;
    bool match = 1;

    for (int i = 0; i < N; i++)
    {
        if (abs(hostRef[i] - gpuRef[i]) > epsilon)
        {
            match = 0;
            printf("Arrays do not match!\n");
            printf("host %5.2f gpu %5.2f at current %d\n", hostRef[i],
                   gpuRef[i], i);
            break;
        }
    }

    if (match) printf("Arrays match.\n\n");

    return;
}

void initialData(float *ip, int size)
{
    // generate different seed for random number
    time_t t;
    srand((unsigned) time(&t));

    for (int i = 0; i < size; i++)
    {
        ip[i] = (float)( rand() & 0xFF ) / 10.0f;
    }

    return;
}

void sumArraysOnHost(float *A, float *B, float *C, const int N)
{
    for (int idx = 0; idx < N; idx++)
    {
        C[idx] = A[idx] + B[idx];
    }
}
__global__ void sumArraysOnGPU(float *A, float *B, float *C, const int N)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < N) C[i] = A[i] + B[i];
}

int main(int argc, char **argv)
{
    printf("%s Starting...\n", argv[0]);

    // set up device
    int dev = 0;
    cudaDeviceProp deviceProp;
    CHECK(cudaGetDeviceProperties(&deviceProp, dev));
    printf("Using Device %d: %s\n", dev, deviceProp.name);
    CHECK(cudaSetDevice(dev));

    // set up data size of vectors
    int nElem = 1 << 24;
    printf("Vector size %d\n", nElem);

    // malloc host memory
    size_t nBytes = nElem * sizeof(float);

    float *h_A, *h_B, *hostRef, *gpuRef;
    h_A     = (float *)malloc(nBytes);
    h_B     = (float *)malloc(nBytes);
    hostRef = (float *)malloc(nBytes);
    gpuRef  = (float *)malloc(nBytes);

    double iStart, iElaps;

    // initialize data at host side
    iStart = seconds();
    initialData(h_A, nElem);
    initialData(h_B, nElem);
    iElaps = seconds() - iStart;
    printf("initialData Time elapsed %f sec\n", iElaps);
    memset(hostRef, 0, nBytes);
    memset(gpuRef,  0, nBytes);

    // add vector at host side for result checks
    iStart = seconds();
    sumArraysOnHost(h_A, h_B, hostRef, nElem);
    iElaps = seconds() - iStart;
    printf("sumArraysOnHost Time elapsed %f sec\n", iElaps);

    // malloc device global memory
    float *d_A, *d_B, *d_C;
    CHECK(cudaMalloc((float**)&d_A, nBytes));
    CHECK(cudaMalloc((float**)&d_B, nBytes));
    CHECK(cudaMalloc((float**)&d_C, nBytes));

    // transfer data from host to device
    CHECK(cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_C, gpuRef, nBytes, cudaMemcpyHostToDevice));

    // invoke kernel at host side
    int iLen = 512;
    dim3 block (iLen);
    dim3 grid  ((nElem + block.x - 1) / block.x);

    iStart = seconds();
    sumArraysOnGPU<<<grid, block>>>(d_A, d_B, d_C, nElem);
    CHECK(cudaDeviceSynchronize());
    iElaps = seconds() - iStart;
    printf("sumArraysOnGPU <<<  %d, %d  >>>  Time elapsed %f sec\n", grid.x,
           block.x, iElaps);

    // check kernel error
    CHECK(cudaGetLastError()) ;

    // copy kernel result back to host side
    CHECK(cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost));

    // check device results
    checkResult(hostRef, gpuRef, nElem);

    // free device global memory
    CHECK(cudaFree(d_A));
    CHECK(cudaFree(d_B));
    CHECK(cudaFree(d_C));

    // free host memory
    free(h_A);
    free(h_B);
    free(hostRef);
    free(gpuRef);

    return(0);
}

nvprof

nvprof ./sumArraysOnGPU-timer

兔子牙丫丫
关注
专栏目录
cuda-tutorial:CUDA编程简短教程
05-18
本教程旨在为初学者提供一个快速入门CUDA编程的指导，帮助他们理解基本概念，并知道如何开始深入学习。 在GPU编程中，CUDA的核心是C/C++的编程模型，它扩展了标准C++，增加了对GPU编程的支持。CUDA编程主要包括以下...
【CUDA】新手的第一个CUDA C程序：“Hello World from GPU !“
TangPlusHPC的博客
06-11 1884
CUDA是一种通用的并行计算平台和编程模型，它利用NVIDIA GPU中的并行计算引擎能更有效地解决复杂的计算问题。CUDA C是标准ANSI C的一个扩展，它带有的少数语言扩展功能使异构编程成为可能，同时也能用过API来管理设备、内存和其他任务。一个CUDA程序包含了以下两部分的混合。NVIDIA 的CUDA nvcc编译器在编译过程中将设备代码从主机代码中分离出来。主机代码是标准的C代码，使用C编译器进行编译。设备代码，也就是核函数，是用扩展的带有标记数据并行函数关键字的CUDA C语言编写的。设备代码
CUDA【1】HelloWorld
椰子奶糖的博客
07-19 665
代码】CUDA【1】HelloWorld。
深入浅出之CUDA编程 Hello World
最新发布
浩瀚之水的专栏
10-15 781
值得注意的是，在cpp或h文件中调用cu或cuh中定义的函数，直接采用include的方式包含然后调用暂时是行不通的，可以采用extern的方式将cu或cuh中定义的CUDA函数设置为公共函数，然后在cpp或h中声明这个extern函数，最后在cpp或h文件中需要的地方调用即可。4、最后，打开项目的属性页，如下图所示，在“配置属性”——“链接器”——“输入”中的“附件依赖项”里面添加。：将C++代码和CUDA代码编译并链接在一起，生成可执行文件。下面是一个简单的例子，演示如何在C++中调用CUDA代码。
[CUDA] 快速入门CUDA(1)-基本了解和HelloWorld
weixin_52157994的博客
04-10 9100
CUDA基础 文章目录CUDA基础1 CUDA简介2 GPU和CPU架构的不同之处3 查看GPU硬件信息4 需要建立的基本概念5 总结 1 CUDA简介 CUDA的全程是Computer Unified Device Architecture，是由显卡头子NVIDIA发明的。有的人对于显卡的印象在于它可以玩游戏，效果十分逼真，但从背后而言，正是因为显卡强大的图形计算能力，才使得计算机可以运行这些大型的3D游戏，并且拥有较高的画质和帧数。 2 GPU和CPU架构的不同之处 CPU具有以下特点： 对单线程有优
Cuda实战-01 入门程序HelloWorld
星辰火之梦
03-01 379
不管是学习什么编程语言，一般入门的实战程序一般都是从Hello World开始的。这一节，我们就用Hello World程序进行Cuda实战，程序的内容主要是让CPU和GPU分别打印Hello World，在CPU中打印内容为，在GPU中打印内容为，具体实现请参考程序案例。
【CUDA Runtime】第一个“Hello World“程序
m0_62599305的博客
07-26 958
CUDA（Compute Unified Device Architecture）是 NVIDIA 提供的一个并行计算平台和编程模型，旨在通过利用 GPU 的强大计算能力来加速计算密集型应用。对于许多开发者而言，开始学习 CUDA 编程的第一步通常是编写一个简单的“Hello World”程序。尽管 CUDA 的实际应用远远超出一个简单的例子，但这个初步的示例能帮助开发者快速入门，理解基本的 CUDA 编程流程。
CUDA编程（三）：Hello world
博学而笃志，切问而近思。
04-02 1147
CUDA编程（三）：Hello world
CUDA C编程权威指南.pdf
07-08
### CUDA C编程权威指南知识点概览 #### 第1章 基于CUDA的异构并行计算 ...以上内容涵盖了《CUDA C编程权威指南》的主要知识点，旨在帮助读者全面了解CUDA编程的基本原理和高级技巧，为后续深入学习打下坚实的基础。
GPU并行计算与CUDA编程 - 副本.zip
09-18
学生可能会学习到如何创建第一个CUDA程序，即“Hello World”程序，以了解CUDA编程的基本结构。 第二课可能深入讲解CUDA编程的核心概念，如线程块和网格的概念，以及线程索引和内存模型。线程块和网格是CUDA编程中...
CUDA编程学习（六份pdf）
05-03
在CUDA编程学习1.pdf中，可能涵盖了CUDA环境的设置、CUDA SDK的安装与使用，以及第一个CUDA程序“Hello World”的编写。 第二部分：CUDA内存模型 CUDA的内存模型是其性能关键。GPU内存分为全局内存、共享内存、常量...
GPU高性能编程CUDA实战
03-27
GPU高性能编程CUDA实战GPU高性能编程CUDA实战GPU高性能编程CUDA实战GPU高性能编程CUDA实战
cuda编程入门-HelloWorld
@way
11-11 1011
为了提高并行运算速度，开始接触cuda编程，cuda 在安装的时候提供了很多示例，一般位于Home目录下，可以作为学习的参考，同时官网提供的入门教程，就是很好的学习资料。 cuda编程首先要转换与cpu 下编程的思想，在cpu 下，我们一般都是按照顺序进行执行，或者开几个线程处理几个任务，而在gpu 下，是分block 和 thread的，gpu由多个block（x,y ,z ）组成，每个bloc...
CUDA学习笔记1——线程组织、HelloWorld、数组相加、内存结构、核函数、设备函数
illyh的博客
12-10 947
一个SM中寄存器的数量比较有限，一旦核函数使用了超过硬件限制的寄存器数量，则会使用本地内存来代替多占用的寄存器，这种寄存器溢出的情况会带来性能上的不利影响，实际编程过程中我们应该避免这种情况。操作系统在逻辑层面将CPU内存分为两类：可分页内存（Pageable Memory）和页锁定内存（Page Lock Memory, 又称为Pinned Memory），可分页内存没有锁定特性，可能会被交换出去，比如传输到硬盘上，而页锁定内存具有锁定特性，不会被交换出去。共享内存具有仅次于寄存器的读写速度。
GPU（四）CUDA hello world
NUCEMLS的博客
03-09 979
CUDA GPU hello world
CUDA从入门到放弃（三）：Hello World 之第一个CUDA程序
mingo_敏
05-12 2829
vs2015中打开新建一个CUDA程序 拿个CUDA中的 hello world 示例程序过来跑一下吧！ 示例来源于：CUDA Samples 计算两个数组间的加法 #include "cuda_runtime.h" #include "device_launch_parameters.h" #include <stdio.h> cudaError_t addWithCuda(i...
《cuda c编程权威指南》01- 用gpu输出hello world
一千零一夜的博客
07-30 492
cuda 输出hello world
CUDA编程：Hello world
BigCowPeking
07-30 1797
第三章 CUDA C简介 输出hello world #include&lt;stdio.h&gt; __global__ void kernel() { printf("hello world"); } int main() { kernel&lt;&lt;&lt;1, 1&gt;&gt;&gt;(); return 0; } 这个程序和普通的C程序的区别值得...
CUDA学习笔记02：测试程序hello world
weixin_42130300的博客
03-04 686
一些测试的小笔记
评论 1
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值