《CUDA 编程入门》笔记

最新推荐文章于 2024-11-02 23:48:09 发布
最新推荐文章于 2024-11-02 23:48:09 发布
阅读量256 收藏
点赞数
分类专栏： CUDA 文章标签： 笔记 cuda
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/taifyang/article/details/131179029
版权

在这里插入图片描述

helloworld：vec-add

#include <stdio.h>
#include <cuda.h>

#include "_aux.h"

/* host, add */
void vec_add_host(float *x, float *y, float *z, int N);

/* device function */
__global__ void vec_add(float *x, float *y, float *z, int N)
{
    /* 1D block */
    int idx = get_tid();

    if (idx < N) z[idx] = z[idx] + y[idx] + x[idx];
}

void vec_add_host(float *x, float *y, float *z, int N)
{
    int i;

    for (i = 0; i < N; i++) z[i] = z[i] + y[i] + x[i];
}

int main()
{
    int N = 20000000;
    int nbytes = N * sizeof(float);

    /* 1D block */
    int bs = 256;

    /* 2D grid */
    int s = ceil(sqrt((N + bs - 1.) / bs));
    dim3 grid = dim3(s, s);

    float *dx = NULL, *hx = NULL;
    float *dy = NULL, *hy = NULL;
    float *dz = NULL, *hz = NULL;

    int itr = 30;
    int i;
    double th, td;

    /* allocate GPU mem */
    cudaMalloc((void **)&dx, nbytes);
    cudaMalloc((void **)&dy, nbytes);
    cudaMalloc((void **)&dz, nbytes);

    if (dx == NULL || dy == NULL || dz == NULL) {
        printf("couldn't allocate GPU memory\n");
        return -1;
    }

    printf("allocated %.2f MB on GPU\n", nbytes / (1024.f * 1024.f));

    /* alllocate CPU mem */
    hx = (float *) malloc(nbytes);
    hy = (float *) malloc(nbytes);
    hz = (float *) malloc(nbytes);

    if (hx == NULL || hy == NULL || hz == NULL) {
        printf("couldn't allocate CPU memory\n");
        return -2;
    }
    printf("allocated %.2f MB on CPU\n", nbytes / (1024.f * 1024.f));

    /* init */
    for (i = 0; i < N; i++) {
        hx[i] = 1;
        hy[i] = 1;
        hz[i] = 1;
    }

    /* copy data to GPU */
    cudaMemcpy(dx, hx, nbytes, cudaMemcpyHostToDevice);
    cudaMemcpy(dy, hy, nbytes, cudaMemcpyHostToDevice);
    cudaMemcpy(dz, hz, nbytes, cudaMemcpyHostToDevice);

    /* call GPU */
    cudaDeviceSynchronize();
    td = get_time();
    
    for (i = 0; i < itr; i++) vec_add<<<grid, bs>>>(dx, dy, dz, N);

    cudaDeviceSynchronize();
    td = get_time() - td;

    /* CPU */
    th = get_time();
    for (i = 0; i < itr; i++) vec_add_host(hx, hy, hz, N);
    th = get_time() - th;

    printf("GPU time: %e, CPU time: %e, speedup: %g\n", td, th, th / td);

    cudaFree(dx);
    cudaFree(dy);
    cudaFree(dz);

    free(hx);
    free(hy);
    free(hz);

    return 0;
}

内存管理

#include <stdio.h>
#include <cuda.h>


__global__ void sum(float *x)
{
    int tid = threadIdx.x;
    x[tid] += 1;
}

int main()
{
    int N = 32;
    int nbytes = N * sizeof(float);

    float *dx = NULL, *hx = NULL;
    int i;

    /* allocate GPU mem */
    cudaMalloc((void **)&dx, nbytes);

    if (dx == NULL) {
        printf("couldn't allocate GPU memory\n");
        return -1;
    }

    /* allocate CPU host mem: memory copy is faster than malloc */
    //hx = (float *)malloc(nbytes);
    cudaMallocHost((void **)&hx, nbytes);

    if (hx == NULL) {
        printf("couldn't allocate CPU memory\n");
        return -2;
    }

    /* init */
    printf("hx original: \n");
    for (i = 0; i < N; i++) {
        hx[i] = i;
        printf("%g\n", hx[i]);
    }

    /* copy data to GPU */
    cudaMemcpy(dx, hx, nbytes, cudaMemcpyHostToDevice);

    /* call GPU */
    sum<<<1, N>>>(dx);

    /* let GPU finish */
    cudaDeviceSynchronize();

    /* copy data from GPU */
    cudaMemcpy(hx, dx, nbytes, cudaMemcpyDeviceToHost);

    printf("\nhx from GPU: \n");
    for (i = 0; i < N; i++) {
        printf("%g\n", hx[i]);
    }

    cudaFree(dx);
    //free(hx);
    cudaFreeHost(hx);

    return 0;
}

规约算法

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

volatile关键字能防止编译器优化循环展开，造成结果错误。
在这里插入图片描述
注意：下面的程序若去掉volatile关键字，则函数reduction_2和reduction_3的计算结果有误。

#include <stdio.h>
#include <cuda.h>


__global__ void reduction_1(const float* x, float* y)
{
	__shared__ float sdata[256];
	int tid = threadIdx.x;
	/*load data to shared mem*/
	sdata[tid] = x[tid];
	__syncthreads();

	/*reduction using shared mem*/
	if (tid < 128)sdata[tid] += sdata[tid + 128]; __syncthreads();
	if (tid < 64) sdata[tid] += sdata[tid + 64]; __syncthreads();
	if (tid < 32) sdata[tid] += sdata[tid + 32]; __syncthreads();
	if (tid < 16) sdata[tid] += sdata[tid + 16]; __syncthreads();
	if (tid < 8) sdata[tid] += sdata[tid + 8];__syncthreads();
	if (tid < 4) sdata[tid] += sdata[tid + 4]; __syncthreads();
	if (tid < 2) sdata[tid] += sdata[tid + 2]; __syncthreads();
	if (tid == 0) *y = sdata[0] + sdata[1];
}

__global__ void reduction_2(const float* x, float* y)
{
	__shared__ volatile float sdata[256];
	int tid = threadIdx.x;
	/*load data to shared mem*/
	sdata[tid] = x[tid];
	__syncthreads();
	/*reduction using shared mem*/
	if (tid < 128) sdata[tid] += sdata[tid + 128]; __syncthreads();
	if (tid < 64)  sdata[tid] += sdata[tid + 64]; __syncthreads();
	if (tid < 32) {
		sdata[tid] += sdata[tid + 32];
		sdata[tid] += sdata[tid + 16];
		sdata[tid] += sdata[tid + 8];
		sdata[tid] += sdata[tid + 4];
		sdata[tid] += sdata[tid + 2];
		sdata[tid] += sdata[tid + 1];
	}
	if (tid == 0) y[0] = sdata[0];
}

__device__ void warpReduce(volatile float* sdata, int tid)
{
	sdata[tid] += sdata[tid + 32]; 
	sdata[tid] += sdata[tid + 16]; 
	sdata[tid] += sdata[tid + 8]; 
	sdata[tid] += sdata[tid + 4]; 
	sdata[tid] += sdata[tid + 2]; 
	sdata[tid] += sdata[tid + 1];
}

__global__ void reduction_3(const float* x, float* y)
{
	__shared__ float sdata[256];
	int tid = threadIdx.x;
	/* load data to shared mem*/
	sdata[tid] = x[tid];
	__syncthreads();
	/*reduction using shared mem*/
	if (tid < 128)sdata[tid] += sdata[tid + 128]; __syncthreads();
	if (tid < 64) sdata[tid] += sdata[tid + 64]; __syncthreads();
	if (tid < 32) warpReduce(sdata, tid);
	if (tid == 0) y[0] = sdata[0];
}


int main()
{
	int N = 256;
	int nbytes = N * sizeof(float);

	float* dx = NULL, *hx = NULL;
	float* dy = NULL;
	int i;
	float as = 0;

	cudaMalloc((void**)& dx, nbytes);
	cudaMalloc((void**)& dy, sizeof(float));

	if (dx == NULL || dy == NULL) 
		return -1;

	hx = (float*)malloc(nbytes);
	if (hx == NULL) 
		return -2;

	for (i = 0; i < N; i++) 
		hx[i] = 1;

	cudaMemcpy(dx, hx, nbytes, cudaMemcpyHostToDevice);

	reduction_1 << <1, N >> > (dx, dy);
	//reduction_2 << <1, N >> > (dx, dy);
	//reduction_3 << <1, N >> > (dx, dy);

	cudaDeviceSynchronize();

	cudaMemcpy(&as, dy, sizeof(float), cudaMemcpyDeviceToHost);

	printf("answer: 256, calculated by GPU:%g\n", as);

	cudaFree(dx);
	cudaFree(dy);
	free(hx);

	return 0;
}

完整的三阶段规约算法
1.块大小, 256：数组长度降低256倍 ——大规模数组依旧很长，例如256万降低到1万
2.对部分和继续使用上一步的算法
3.使用一个块，将最后结果规约


/* asum: sum of all entries of a vector */

#include "aux.h"

typedef double FLOAT;

__device__ void warpReduce(volatile FLOAT *sdata, int tid)
{
    sdata[tid] += sdata[tid + 32];
    sdata[tid] += sdata[tid + 16];
    sdata[tid] += sdata[tid + 8];
    sdata[tid] += sdata[tid + 4];
    sdata[tid] += sdata[tid + 2];
    sdata[tid] += sdata[tid + 1];
}

/* sum all entries in x and asign to y
 * block dim must be 256 */
__global__ void asum_stg_1(const FLOAT *x, FLOAT *y, int N)
{
    __shared__ FLOAT sdata[256];
    int idx = get_tid();
    int tid = threadIdx.x;
    int bid = get_bid();

    /* load data to shared mem */
    if (idx < N) {
        sdata[tid] = x[idx];
    }
    else {
        sdata[tid] = 0;
    }

    __syncthreads();

    /* reduction using shared mem */
    if (tid < 128) sdata[tid] += sdata[tid + 128];
    __syncthreads();

    if (tid < 64) sdata[tid] += sdata[tid + 64];
    __syncthreads();

    if (tid < 32) warpReduce(sdata, tid);

    if (tid == 0) y[bid] = sdata[0];
}

__global__ void asum_stg_3(FLOAT *x, int N)
{
    __shared__ FLOAT sdata[128];
    int tid = threadIdx.x;
    int i;

    sdata[tid] = 0;

    /* load data to shared mem */
    for (i = 0; i < N; i += 128) {
        if (tid + i < N) sdata[tid] += x[i + tid];
    }

    __syncthreads();

    /* reduction using shared mem */
    if (tid < 64) sdata[tid] = sdata[tid] + sdata[tid + 64];
    __syncthreads();

    if (tid < 32) warpReduce(sdata, tid);
    __syncthreads();

    if (tid == 0) x[0] = sdata[0];
}

/* dy and dz serve as cache: result stores in dz[0] */
void asum(FLOAT *dx, FLOAT *dy, FLOAT *dz, int N)
{
    /* 1D block */
    int bs = 256;

    /* 2D grid */
    int s = ceil(sqrt((N + bs - 1.) / bs));
    dim3 grid = dim3(s, s);
    int gs = 0;

    /* stage 1 */
    asum_stg_1<<<grid, bs>>>(dx, dy, N);

    /* stage 2 */
    {
        /* 1D grid */
        int N2 = (N + bs - 1) / bs;

        int s2 = ceil(sqrt((N2 + bs - 1.) / bs));
        dim3 grid2 = dim3(s2, s2);

        asum_stg_1<<<grid2, bs>>>(dy, dz, N2);

        /* record gs */
        gs = (N2 + bs - 1.) / bs;
    }

    /* stage 3 */
    asum_stg_3<<<1, 128>>>(dz, gs);
}

FLOAT asum_host(FLOAT *x, int N);

FLOAT asum_host(FLOAT *x, int N)
{
    int i;
    FLOAT t = 0;

    for (i = 0; i < N; i++) t += x[i];

    return t;
}

int main(int argc, char **argv)
{
    int N = 10000070;
    int nbytes = N * sizeof(FLOAT);

    FLOAT *dx = NULL, *hx = NULL;
    FLOAT *dy = NULL, *dz;
    int i, itr = 20;
    FLOAT asd = 0, ash;
    double td, th;

    if (argc == 2) {
        int an;

        an = atoi(argv[1]);
        if (an > 0) N = an;
    }

    /* allocate GPU mem */
    cudaMalloc((void **)&dx, nbytes);
    cudaMalloc((void **)&dy, sizeof(FLOAT) * ((N + 255) / 256));
    cudaMalloc((void **)&dz, sizeof(FLOAT) * ((N + 255) / 256));

    if (dx == NULL || dy == NULL || dz == NULL) {
        printf("couldn't allocate GPU memory\n");
        return -1;
    }

    printf("allocated %e MB on GPU\n", nbytes / (1024.f * 1024.f));

    /* alllocate CPU mem */
    hx = (FLOAT *) malloc(nbytes);

    if (hx == NULL) {
        printf("couldn't allocate CPU memory\n");
        return -2;
    }
    printf("allocated %e MB on CPU\n", nbytes / (1024.f * 1024.f));

    /* init */
    for (i = 0; i < N; i++) {
        hx[i] = 1;
    }

    /* copy data to GPU */
    cudaMemcpy(dx, hx, nbytes, cudaMemcpyHostToDevice);

    /* let dust fall */
    cudaDeviceSynchronize();
    td = get_time();

    /* call GPU */
    for (i = 0; i < itr; i++) asum(dx, dy, dz, N);

    /* let GPU finish */
    cudaDeviceSynchronize();
    td = get_time() - td;

    th = get_time();
    for (i = 0; i < itr; i++) ash = asum_host(hx, N);
    th = get_time() - th;

    /* copy data from GPU */
    cudaMemcpy(&asd, dz, sizeof(FLOAT), cudaMemcpyDeviceToHost);

    printf("asum, answer: %d, calculated by GPU:%f, calculated by CPU:%f\n", N, asd, ash);
    printf("GPU time: %e, CPU time: %e, speedup: %g\n", td, th, th / td);

    cudaFree(dx);
    cudaFree(dy);
    cudaFree(dz);
    free(hx);

    return 0;
}

#include "_aux.h"
#include <assert.h>


/* host, add */
float dot_host(float *x, float *y, int N)
{
    int i;
    float t = 0;

    assert(x != NULL);
    assert(y != NULL);

    for (i = 0; i < N; i++) t += x[i] * y[i];

    return t;
}

__device__ void warpReduce(volatile float *sdata, int tid)
{
    sdata[tid] += sdata[tid + 32];
    sdata[tid] += sdata[tid + 16];
    sdata[tid] += sdata[tid + 8];
    sdata[tid] += sdata[tid + 4];
    sdata[tid] += sdata[tid + 2];
    sdata[tid] += sdata[tid + 1];
}

/* partial dot product */
__global__ void dot_stg_1(const float *x, float *y, float *z, int N)
{
    __shared__ float sdata[256];
    int idx = get_tid();
    int tid = threadIdx.x;
    int bid = get_bid();

    /* load data to shared mem */
    if (idx < N) {
        sdata[tid] = x[idx] * y[idx];
    }
    else {
        sdata[tid] = 0;
    }

    __syncthreads();

    /* reduction using shared mem */
    if (tid < 128) sdata[tid] += sdata[tid + 128];
    __syncthreads();

    if (tid < 64) sdata[tid] += sdata[tid + 64];
    __syncthreads();

    if (tid < 32) warpReduce(sdata, tid);

    if (tid == 0) z[bid] = sdata[0];
}

/* sum all entries in x and asign to y
 * block dim must be 256 */
__global__ void dot_stg_2(const float *x, float *y, int N)
{
    __shared__ float sdata[256];
    int idx = get_tid();
    int tid = threadIdx.x;
    int bid = get_bid();

    /* load data to shared mem */
    if (idx < N) {
        sdata[tid] = x[idx];
    }
    else {
        sdata[tid] = 0;
    }

    __syncthreads();

    /* reduction using shared mem */
    if (tid < 128) sdata[tid] += sdata[tid + 128];
    __syncthreads();

    if (tid < 64) sdata[tid] += sdata[tid + 64];
    __syncthreads();

    if (tid < 32) warpReduce(sdata, tid);

    if (tid == 0) y[bid] = sdata[0];
}

__global__ void dot_stg_3(float *x, int N)
{
    __shared__ float sdata[128];
    int tid = threadIdx.x;
    int i;

    sdata[tid] = 0;

    /* load data to shared mem */
    for (i = 0; i < N; i += 128) {
        if (tid + i < N) sdata[tid] += x[i + tid];
    }

    __syncthreads();

    /* reduction using shared mem */
    if (tid < 64) sdata[tid] = sdata[tid] + sdata[tid + 64];
    __syncthreads();

    if (tid < 32) warpReduce(sdata, tid);

    if (tid == 0) x[0] = sdata[0];
}

/* dz and d serve as cache: result stores in d[0] */
void dot_device(float *dx, float *dy, float *dz, float *d, int N)
{
    /* 1D block */
    int bs = 256;

    /* 2D grid */
    int s = ceil(sqrt((N + bs - 1.) / bs));
    dim3 grid = dim3(s, s);
    int gs = 0;

    /* stage 1 */
    dot_stg_1<<<grid, bs>>>(dx, dy, dz, N);

    /* stage 2 */
    {
        /* 1D grid */
        int N2 = (N + bs - 1) / bs;

        int s2 = ceil(sqrt((N2 + bs - 1.) / bs));
        dim3 grid2 = dim3(s2, s2);

        dot_stg_2<<<grid2, bs>>>(dz, d, N2);

        /* record gs */
        gs = (N2 + bs - 1.) / bs;
    }

    /* stage 3 */
    dot_stg_3<<<1, 128>>>(d, gs);
}

int main(int argc, char **argv)
{
    int N = 10000070;
    int nbytes = N * sizeof(float);

    float *hx = NULL, *hy = NULL;
    float *dx = NULL, *dy = NULL, *dz = NULL, *d = NULL;
    int i, itr = 20;
    float asd = 0, ash;
    double td, th;

    if (argc == 2) {
        int an;

        an = atoi(argv[1]);
        if (an > 0) N = an;
    }

    /* allocate GPU mem */
    cudaMalloc((void **)&dx, nbytes);
    cudaMalloc((void **)&dy, nbytes);

    cudaMalloc((void **)&dz, sizeof(float) * ((N + 255) / 256));
    cudaMalloc((void **)&d, sizeof(float) * ((N + 255) / 256));

    if (dx == NULL || dy == NULL || dz == NULL || d == NULL) {
        printf("couldn't allocate GPU memory\n");
        return -1;
    }

    printf("allocated %e MB on GPU\n", nbytes / (1024.f * 1024.f));

    /* alllocate CPU mem */
    hx = (float *) malloc(nbytes);
    hy = (float *) malloc(nbytes);

    if (hx == NULL || hy == NULL) {
        printf("couldn't allocate CPU memory\n");
        return -2;
    }
    printf("allocated %e MB on CPU\n", nbytes / (1024.f * 1024.f));

    /* init */
    for (i = 0; i < N; i++) {
        hx[i] = 1;
        hy[i] = 2;
    }

    /* copy data to GPU */
    cudaMemcpy(dx, hx, nbytes, cudaMemcpyHostToDevice);
    cudaMemcpy(dy, hy, nbytes, cudaMemcpyHostToDevice);

    /* let dust fall */
    cudaDeviceSynchronize();
    td = get_time();

    /* call GPU */
    for (i = 0; i < itr; i++) dot_device(dx, dy, dz, d, N);

    /* let GPU finish */
    cudaDeviceSynchronize();
    td = get_time() - td;

    th = get_time();
    for (i = 0; i < itr; i++) ash = dot_host(hx, hy, N);
    th = get_time() - th;

    /* copy data from GPU */
    cudaMemcpy(&asd, d, sizeof(float), cudaMemcpyDeviceToHost);

    printf("dot, answer: %d, calculated by GPU:%f, calculated by CPU:%f\n", 2 * N, asd, ash);
    printf("GPU time: %e, CPU time: %e, speedup: %g\n", td, th, th / td);

    cudaFree(dx);
    cudaFree(dy);
    cudaFree(dz);
    cudaFree(d);

    free(hx);
    free(hy);

    return 0;
}
给算法爸爸上香
关注
专栏目录
cuda编码入门学习笔记
xiaomu_347的博客
06-27 876
首先，确保你对CUDA编程的基本概念有所了解：- CUDA是NVIDIA提供的用于并行计算的平台和编程模型。- CUDA允许你利用NVIDIA GPU的并行处理能力来加速计算任务。
CUDA C编程1 - 入门介绍
DU_YULIN的博客
01-13 1016
系列文章目录 文章目录系列文章目录前言1. CUDA C程序组成2. CUDA C编程范式总结参考资料 前言 今天开始又一次开始学习CUDA C编程，今后也会陆续更新学习笔记以及代码，希望加深记忆的同时，能够跟大家分享学习所得。 1. CUDA C程序组成 CUDA C程序主要包括两部分： （1）主机代码 （2）设备代码 CUDA程序保存为.cu文件，CUDA平台使用编译器nvcc编译器来编译设备代码，运行在GPU上；使用C语言编译器来编译主机代码，运行在CPU上。 2. CUDA C编程范式
cuda编程入门笔记
薛定谔的猫头咕咕咕
11-23 752
基本操作 一个Grid中含有多个Block，一个Block中含有多个thread gridDim.x表示网格的块数量 blockIdx.x表示当前块的索引 blockDim.x表示一个块中的线程数量 threadIdx.x表示当前块中线程的索引 <<<NUMBER_OF_BLOCK, NUMBER_OF_THREAD_PER_BLOCK>>> 启动核函数时，核函数代码由每个已配置的线程块中的每个线程执行； #include <stdio.h> // cud
CUDA编程入门笔记
mumusan2016的博客
11-02 346
1、线程块（block）是独立执行的，在执行的过程中线程块之间互不干扰，因此它们的执行顺序是随机的 2、同一线程块中的线程可以通过访问共享内存（shared memory）或者通过同步函数__syncthreads()来协调合作。 3、 4、
CUDA编程入门笔记1
【黑键】
06-05 331
要开始给项目中的程序做速度上的优化 由于CPU的计算速度比较慢 所以想用GPU来进行大量相同的计算 CUDA编程入门 CUDA C基础 CUDA C是对C/C++语言进行拓展后形成的变种，兼容C/C++语法，文件类型为”.cu”文件，编译器为”nvcc”，相比传统的C/C++，主要添加了以下几个方面： 函数类型限定符 执行配置运算符 五个内置变量 变量类型限定符 其...
【CUDA编程系列】CUDA编程基本入门学习笔记
sazass的博客
08-20 3379
参考博客： https://www.cnblogs.com/skyfsm/p/9673960.html 1.GPU特点 核心很多； 规则数据结构，高度统一； 相互无依赖的大规模数据； 不需要被打断的纯净的计算环境； 可预测存储模式； 在并行化计算表现极佳； 非常适合处理图像领域的计算，图像在计算机中的表现形式是矩阵，而很多矩阵运算可以实现并行计算； 一般，串行部分跑CPU，并行部分跑GPU，异构计算； 2.CUDA线程模型 从小到大： **1.Thread：**线程，并行的基本单位 **2.Thread B
【CUDA编程】CUDA入门笔记
非晚非晚的博客
01-23 1921
显卡内部，有三级结构：网格（grid）、块（block）、线程（thread）。每个显卡只有很少的网格，一个核函数只能运行在一个网格中，而一个网格里有多个块，每个块包含了若干线程。kernel在device上执行时实际上是启动很多线程，一个kernel所启动的所有线程称为一个网格（grid），同一个网格上的线程共享相同的全局内存空间，grid是线程结构的第一层次，而网格又可以分为很多线程块（block），一个线程块里面包含很多线程，这是第二个层次。线程两层组织结构如下图所示，这是一个gird和block均为
linux cuda 编程指南,CUDA编程指南阅读笔记
weixin_39575047的博客
05-16 843
3.3 内存层次(Memory Hierarchy)在GPU上CUDA线程可以访问到的存储资源有很多，每个CUDA线程拥有独立的本地内存(local Memory)；每一个线程块(block)都有其独立的共享内存(shared memory)，共享内存对于线程块中的每个线程都是可见的，它与线程块具有相同的生存时间；同时，还有一片称为全局内存(global memory)的区域对所有的CUDA线程都...
笔记-cuda编程:基础与实践
qq_43647367的博客
06-24 372
cuda程序框架一般为主机数据拷贝到设备，执行核函数，设备数据拷贝到主机，由于cpu是不需要数据拷贝的，所以这整体时间在于执行任务函数，而cuda正是可以加速这一点；但是cpu与gpu内存间的PCIe带宽远小于gpu的显存带宽(cpu与gpu数据传输很慢)，所以显而易见的加速gpu(相同cpu运行时间下，怎么让gpu运行时间更短)的方法有减少数据传输的时间，增加核函数的算术强度，增加并行程度；而上述都是从"宏观"或者软件层面上说的，而限于gpu硬件实现，也有一些有关速度的值得注意的地方。
CUDA编程之入门
USTC暖暖的博客
11-12 577
CUDA代码中的函数 CUDA程序文件后缀为.cu。一个.cu文件内既包含CPU程序（称为主机程序），也包含GPU程序（称为设备程序）。凡是挂有“__global__”或者“__device__”前缀的函数，都是在GPU上运行的设备程序，不同的是__global__设备程序可被主机程序调用，而__device__设备程序则只能被设备程序调用。没有挂任何前缀的函数，都是主机程序。主机程序显示声...
CUDA从入门到精通.docx
08-16
### CUDA从入门到精通 #### 一、简介与背景 CUDA (Compute Unified Device Architecture) 是由 NVIDIA 开发的一种并行计算平台和编程模型。通过利用 NVIDIA 的 GPU 进行通用计算，CUDA 将并行计算从传统的大型集群...
【CUDA编程笔记】（3）opencv4 + CUDA 并行图像处理实战入门：图像均值滤波和图像反色
qq_43799400的博客
05-15 1507
目录一、环境配置与测试二、cuda与OpenCV结合方法三、代码实例 相关阅读：【CUDA编程笔记】（1）CUDA C 编程入门 一、环境配置与测试 因为学校的机房vs2017已经配置好 cuda10.0 了，但没有配置opencv，所以这里只考虑配置opencv的事情，opencv4的配置参考视频 VS_2017(Visual Studio)配置OpenCV教程 跟着上面教程配置好opencv之后，新建一个cuda10.0 runtime的项目，在kernel.cu文件中运行以下测试代码： 测试代码来源于
CUDA入门笔记（二）CUDA编程基础——CPU、GPU的关系
AuntieLee的博客
09-07 1565
参考： 优达学城：https://classroom.udacity.com/courses/cs344/lessons/55120467/concepts/670611900923 CUDA Zone：https://www.nvidia.cn/object/cuda_education_cn_old.html   一、CPU与GPU的关系 1.计算机有两个不同的处理器，CPU（hos...
嵌入式通信协议：UART简明学习笔记
weixin_73867577的博客
10-31 737
学习UART串口的简明学习笔记
[自用，更新自day5]瑞吉外卖代码及笔记
lapiii的博客
11-01 615
当使用ThreadLocal维护变量时，ThreadLocal为每个使用该变量的线程提供独立的变量副本，所以每一个线程都可以独立地改变自己的副本，而不会影响其它线程所对应的副本。Mybatis Plus公共字段自动填充，也就是在插入或者更新的时候为指定字段赋予指定的值，使用它的好处就是可以统一对这些字段进行处理，避免了重复代码。因为在分页查询的Dish的records(菜品记录中)，只有这个菜品所属的categoryId，但是我们需要分页的时候展示的是菜品名字。
嵌入式通信协议：MODBUS简明学习笔记
weixin_73867577的博客
11-02 158
学习Modbus的简明学习笔记
点云学习笔记3——读取点云文件、进行统计滤波/直通滤波后可视化
qq_64095888的博客
10-29 196
【代码】点云学习笔记3——读取点云文件、进行统计滤波后可视化。
diffusion model 学习笔记
wubing's 博客
11-01 817
郎之万公式: 描述了粒子做随机布朗运动 (粒子位置随时间变化的关系), 是一种 SDE, 描述了 由梯度力(即.即我们在无条件模型 DDPM 估计的噪声中添加一个微小的扰动(. 基本上,在采样一半时,基本上预测的。输出的是噪声的估计, 这个估计的噪声。通过上面的公式，我们可以知道，当知道。),就可以作为条件模型的噪声估计.首先,假设我们已经有了一个训练好的。: 每一步中保留的原始信号的比例。的计算有各种不同的方法，). 已经接近于真实的。)影响的系统的时间演化.也就是,我们训练的网络。
Spring学习笔记_21——循环依赖
最新发布
LuckyLay的博客
11-02 553
在Spring中的循环依赖就是指一个或者多个Bean之间存在着互相依赖的关系，并且形成了循环调用。例如：在Spring中，Bean-A依赖Bean-B，Bean-B又依赖Bean-A，Bean-A和Bean-B之间就形成了相互依赖的关系。在创建A对象时，发现A对象依赖了B对象，此时先去创建B对象；创建B对象时，发现B对象又依赖了A对象，此时又区创建A对象……形成死循环。同理，在Spring中多个对象之间有可能存在循环依赖。
CUDA编程入门与接口详解
"《GPU编程指南》是一本针对CUDA编程的实用教程，由译者风辰翻译，特别适合初学者入门。CUDA是NVIDIA公司推出的一种通用并行计算架构，旨在利用GPU的强大并行处理能力进行高性能计算。本书主要讲解了CUDA编程的基本...
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

给算法爸爸上香

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付：¥1
获取中
扫码支付

您的余额不足，请更换扫码支付或充值

打赏作者

实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

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

余额充值