CUDA编程（六）进一步并行

CUDA编程（六）

进一步并行

在之前我们使用Thread完成了简单的并行加速，虽然我们的程序运行速度有了50甚至上百倍的提升，但是根据内存带宽来评估的话我们的程序还远远不够，在上一篇博客中给大家介绍了一个访存方面非常重要的优化，我们通过使用连续的内存存取模式，取得了令人满意的优化效果，最终内存带宽也达到了GB/s的级别。

之前也已经提到过了，CUDA不仅提供了Thread，还提供了Grid和Block以及Share Memory这些非常重要的机制，我的显卡的Thread极限是1024，但是通过block和Grid，线程的数量还能成倍增长，甚至用几万个线程。所以本篇博客我们将再次回到线程和并行的角度，进一步的并行加速我们的程序。

Thread AND Block AND Grid

第一篇博客的时候就给大家说明过thread-block-grid 结构了，这里我们再复习一下。

在 CUDA 架构下，显示芯片执行时的最小单位是thread。数个 thread 可以组成一个block。一个 block 中的 thread 能存取同一块共享的内存，而且可以快速进行同步的动作。

每一个 block 所能包含的 thread 数目是有限的。不过，执行相同程序的 block，可以组成grid。不同 block 中的 thread 无法存取同一个共享的内存，因此无法直接互通或进行同步。因此，不同 block 中的 thread 能合作的程度是比较低的。不过，利用这个模式，可以让程序不用担心显示芯片实际上能同时执行的 thread 数目限制。例如，一个具有很少量执行单元的显示芯片，可能会把各个 block 中的 thread 顺序执行，而非同时执行。不同的 grid 则可以执行不同的程序（即 kernel）。

每个 thread 都有自己的一份 register 和 local memory 的空间。同一个 block 中的每个thread 则有共享的一份 share memory。此外，所有的 thread（包括不同 block 的 thread）都共享一份 global memory、constant memory、和 texture memory。不同的 grid 则有各自的 global memory、constant memory 和 texture memory。

大家可能注意到不同block之间是无法进行同步工作的，不过，在我们的程序中，其实不太需要进行 thread 的同步动作，因此我们可以使用多个 block 来进一步增加thread 的数目。

通过多个block使用更多的线程

下面我们就开始继续修改我们的程序：

先贴一下之前的完整代码：

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

//CUDA RunTime API
#include <cuda_runtime.h>

//1M
#define DATA_SIZE 1048576

#define THREAD_NUM 1024

int data[DATA_SIZE];

//产生大量0-9之间的随机数
void GenerateNumbers(int *number, int size)
{
    for (int i = 0; i < size; i++) {
        number[i] = rand() % 10;
    }
}

//打印设备信息
void printDeviceProp(const cudaDeviceProp &prop)
{
    printf("Device Name : %s.\n", prop.name);
    printf("totalGlobalMem : %d.\n", prop.totalGlobalMem);
    printf("sharedMemPerBlock : %d.\n", prop.sharedMemPerBlock);
    printf("regsPerBlock : %d.\n", prop.regsPerBlock);
    printf("warpSize : %d.\n", prop.warpSize);
    printf("memPitch : %d.\n", prop.memPitch);
    printf("maxThreadsPerBlock : %d.\n", prop.maxThreadsPerBlock);
    printf("maxThreadsDim[0 - 2] : %d %d %d.\n", prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);
    printf("maxGridSize[0 - 2] : %d %d %d.\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
    printf("totalConstMem : %d.\n", prop.totalConstMem);
    printf("major.minor : %d.%d.\n", prop.major, prop.minor);
    printf("clockRate : %d.\n", prop.clockRate);
    printf("textureAlignment : %d.\n", prop.textureAlignment);
    printf("deviceOverlap : %d.\n", prop.deviceOverlap);
    printf("multiProcessorCount : %d.\n", prop.multiProcessorCount);
}

//CUDA 初始化
bool InitCUDA()
{
    int count;

    //取得支持Cuda的装置的数目
    cudaGetDeviceCount(&count);

    if (count == 0) {
        fprintf(stderr, "There is no device.\n");
        return false;
    }

    int i;

    for (i = 0; i < count; i++) {

        cudaDeviceProp prop;
        cudaGetDeviceProperties(&prop, i);
        //打印设备信息
        printDeviceProp(prop);

        if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) {
            if (prop.major >= 1) {
                break;
            }
        }
    }

    if (i == count) {
        fprintf(stderr, "There is no device supporting CUDA 1.x.\n");
        return false;
    }

    cudaSetDevice(i);

    return true;
}

// __global__ 函数 (GPU上执行) 计算立方和
__global__ static void sumOfSquares(int *num, int* result, clock_t* time)
{

    //表示目前的 thread 是第几个 thread（由 0 开始计算）
    const int tid = threadIdx.x;

    int sum = 0;

    int i;

    //记录运算开始的时间
    clock_t start;

    //只在 thread 0（即 threadIdx.x = 0 的时候）进行记录
    if (tid == 0) start = clock();

    for (i = tid; i < DATA_SIZE; i += THREAD_NUM) {

        sum += num[i] * num[i] * num[i];

    }

    result[tid] = sum;

    //计算时间的动作，只在 thread 0（即 threadIdx.x = 0 的时候）进行
    if (tid == 0) *time = clock() - start;

}

int main()
{

    //CUDA 初始化
    if (!InitCUDA()) {
        return 0;
    }

    //生成随机数
    GenerateNumbers(data, DATA_SIZE);

    /*把数据复制到显卡内存中*/
    int* gpudata, *result;

    clock_t* time;

    //cudaMalloc 取得一块显卡内存 ( 其中result用来存储计算结果，time用来存储运行时间 )
    cudaMalloc((void**)&gpudata, sizeof(int)* DATA_SIZE);
    cudaMalloc((void**)&result, sizeof(int)*THREAD_NUM);
    cudaMalloc((void**)&time, sizeof(clock_t));

    //cudaMemcpy 将产生的随机数复制到显卡内存中
    //cudaMemcpyHostToDevice - 从内存复制到显卡内存
    //cudaMemcpyDeviceToHost - 从显卡内存复制到内存
    cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);

    // 在CUDA 中执行函数 语法：函数名称<<<block 数目, thread 数目, shared memory 大小>>>(参数...);
    sumOfSquares << < 1, THREAD_NUM, 0 >> >(gpudata, result, time);

    /*把结果从显示芯片复制回主内存*/

    int sum[THREAD_NUM];

    clock_t time_use;

    //cudaMemcpy 将结果从显存中复制回内存
    cudaMemcpy(&sum, result, sizeof(int)* THREAD_NUM, cudaMemcpyDeviceToHost);
    cudaMemcpy(&time_use, time, sizeof(clock_t), cudaMemcpyDeviceToHost);

    //Free
    cudaFree(gpudata);
    cudaFree(result);
    cudaFree(time);

    int final_sum = 0;

    for (int i = 0; i < THREAD_NUM; i++) {

        final_sum += sum[i];

    }

    printf("GPUsum: %d  gputime: %d\n", final_sum, time_use);

    final_sum = 0;

    for (int i = 0; i < DATA_SIZE; i++) {

        final_sum += data[i] * data[i] * data[i];

    }

    printf("CPUsum: %d \n", final_sum);

    return 0;
}

我们要去加入多个block来继续增加我们的线程数量：

首先define一个block的数目

#define THREAD_NUM 256
#define BLOCK_NUM 32

我们准备建立 32 个 blocks，每个 blocks 有 256个 threads，也就是说总共有 32*256= 8192个threads，这里有一个问题，我们为什么不用极限的1024个线程呢？那样就是32*1024 = 32768 个线程，难道不是更好吗？其实并不是这样的，从线程运行的原理来看，线程数量达到一定大小后，我们再一味的增加线程也不会取得性能提升了，反而有可能会让性能下降，感兴趣的同学可以改一下数量试一下。另外我们的加和部分是在CPU上进行的，越多的线程意味着越多的结果，而这也意味着CPU上的运算压力会越来越大。

接着，我们需要修改kernel 部份,加入bid = blockIdx.x：


// __global__ 函数 (GPU上执行) 计算立方和
__global__ static void sumOfSquares(int *num, int* result, clock_t* time)
{

    //表示目前的 thread 是第几个 thread（由 0 开始计算）
    const int tid = threadIdx.x;

    //表示目前的 thread 属于第几个 block（由 0 开始计算）
    const int bid = blockIdx.x;

    int sum = 0;

    int i;

    //记录运算开始的时间
    clock_t start;

    //只在 thread 0（即 threadIdx.x = 0 的时候）进行记录，每个 block 都会记录开始时间及结束时间
    if (tid == 0) time[bid]= clock();

    //thread需要同时通过tid和bid来确定，同时不要忘记保证内存连续性
    for (i =  bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM) {

        sum += num[i] * num[i] * num[i];

    }

    //Result的数量随之增加
    result[bid * THREAD_NUM + tid] = sum;

    //计算时间的动作，只在 thread 0（即 threadIdx.x = 0 的时候）进行，每个 block 都会记录开始时间及结束时间
    if (tid == 0) time[bid + BLOCK_NUM] = clock();

}

关于修改注释已经写得很清楚了。

blockIdx.x 和 threadIdx.x 一样是 CUDA 内建的变量，它表示的是目前的 block 编号。

另外我们把计算时间的方式改成每个 block 都会记录开始时间及结束时间。

因此我们result和time变量的长度需要进行更改：

//cudaMalloc 取得一块显卡内存 ( 其中result用来存储计算结果，time用来存储运行时间 )
cudaMalloc((void**) &result, sizeof(int) * THREAD_NUM * BLOCK_NUM);
cudaMalloc((void**) &time, sizeof(clock_t) * BLOCK_NUM * 2);

然后在调用核函数的时候，把控制block数量的的参数改成我们的block数：


sumOfSquares << < BLOCK_NUM, THREAD_NUM, 0 >> >(gpudata, result, time);

注意从显存复制回内存的部分也需要修改(由于result和time长度的改变)：


    /*把结果从显示芯片复制回主内存*/

    int sum[THREAD_NUM*BLOCK_NUM];

    clock_t time_use[BLOCK_NUM * 2];

    //cudaMemcpy 将结果从显存中复制回内存
    cudaMemcpy(&sum, result, sizeof(int)* THREAD_NUM*BLOCK_NUM, cudaMemcpyDeviceToHost);
    cudaMemcpy(&time_use, time, sizeof(clock_t)* BLOCK_NUM * 2, cudaMemcpyDeviceToHost);

    //Free
    cudaFree(gpudata);
    cudaFree(result);
    cudaFree(time);

    int final_sum = 0;

    for (int i = 0; i < THREAD_NUM*BLOCK_NUM; i++) {

        final_sum += sum[i];

    }

此外，由于涉及到block，我们需要采取不同的计时方式，即把每个 block 最早的开始时间，和最晚的结束时间相减，取得总运行时间。

    //采取新的计时策略 把每个 block 最早的开始时间，和最晚的结束时间相减，取得总运行时间
    clock_t min_start, max_end;

    min_start = time_use[0];

    max_end = time_use[BLOCK_NUM];

    for (int i = 1; i < BLOCK_NUM; i++) {
        if (min_start > time_use[i])
            min_start = time_use[i];
        if (max_end < time_use[i + BLOCK_NUM])
            max_end = time_use[i + BLOCK_NUM];
    }

    printf("GPUsum: %d  gputime: %d\n", final_sum, max_end - min_start);

完整程序：

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

//CUDA RunTime API
#include <cuda_runtime.h>

//1M
#define DATA_SIZE 1048576

#define THREAD_NUM 256

#define BLOCK_NUM 32

int data[DATA_SIZE];

//产生大量0-9之间的随机数
void GenerateNumbers(int *number, int size)
{
    for (int i = 0; i < size; i++) {
        number[i] = rand() % 10;
    }
}

//打印设备信息
void printDeviceProp(const cudaDeviceProp &prop)
{
    printf("Device Name : %s.\n", prop.name);
    printf("totalGlobalMem : %d.\n", prop.totalGlobalMem);
    printf("sharedMemPerBlock : %d.\n", prop.sharedMemPerBlock);
    printf("regsPerBlock : %d.\n", prop.regsPerBlock);
    printf("warpSize : %d.\n", prop.warpSize);
    printf("memPitch : %d.\n", prop.memPitch);
    printf("maxThreadsPerBlock : %d.\n", prop.maxThreadsPerBlock);
    printf("maxThreadsDim[0 - 2] : %d %d %d.\n", prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);
    printf("maxGridSize[0 - 2] : %d %d %d.\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
    printf("totalConstMem : %d.\n", prop.totalConstMem);
    printf("major.minor : %d.%d.\n", prop.major, prop.minor);
    printf("clockRate : %d.\n", prop.clockRate);
    printf("textureAlignment : %d.\n", prop.textureAlignment);
    printf("deviceOverlap : %d.\n", prop.deviceOverlap);
    printf("multiProcessorCount : %d.\n", prop.multiProcessorCount);
}

//CUDA 初始化
bool InitCUDA()
{
    int count;

    //取得支持Cuda的装置的数目
    cudaGetDeviceCount(&count);

    if (count == 0) {
        fprintf(stderr, "There is no device.\n");
        return false;
    }

    int i;

    for (i = 0; i < count; i++) {

        cudaDeviceProp prop;
        cudaGetDeviceProperties(&prop, i);
        //打印设备信息
        printDeviceProp(prop);

        if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) {
            if (prop.major >= 1) {
                break;
            }
        }
    }

    if (i == count) {
        fprintf(stderr, "There is no device supporting CUDA 1.x.\n");
        return false;
    }

    cudaSetDevice(i);

    return true;
}

// __global__ 函数 (GPU上执行) 计算立方和
// __global__ 函数 (GPU上执行) 计算立方和
__global__ static void sumOfSquares(int *num, int* result, clock_t* time)
{

    //表示目前的 thread 是第几个 thread（由 0 开始计算）
    const int tid = threadIdx.x;

    //表示目前的 thread 属于第几个 block（由 0 开始计算）
    const int bid = blockIdx.x;

    int sum = 0;

    int i;

    //记录运算开始的时间
    clock_t start;

    //只在 thread 0（即 threadIdx.x = 0 的时候）进行记录，每个 block 都会记录开始时间及结束时间
    if (tid == 0) time[bid] = clock();

    //thread需要同时通过tid和bid来确定，同时不要忘记保证内存连续性
    for (i = bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM) {

        sum += num[i] * num[i] * num[i];

    }

    //Result的数量随之增加
    result[bid * THREAD_NUM + tid] = sum;

    //计算时间的动作，只在 thread 0（即 threadIdx.x = 0 的时候）进行，每个 block 都会记录开始时间及结束时间
    if (tid == 0) time[bid + BLOCK_NUM] = clock();

}

int main()
{

    //CUDA 初始化
    if (!InitCUDA()) {
        return 0;
    }

    //生成随机数
    GenerateNumbers(data, DATA_SIZE);

    /*把数据复制到显卡内存中*/
    int* gpudata, *result;

    clock_t* time;

    //cudaMalloc 取得一块显卡内存 ( 其中result用来存储计算结果，time用来存储运行时间 )
    cudaMalloc((void**)&gpudata, sizeof(int)* DATA_SIZE);
    cudaMalloc((void**)&result, sizeof(int)*THREAD_NUM* BLOCK_NUM);
    cudaMalloc((void**)&time, sizeof(clock_t)* BLOCK_NUM * 2);

    //cudaMemcpy 将产生的随机数复制到显卡内存中
    //cudaMemcpyHostToDevice - 从内存复制到显卡内存
    //cudaMemcpyDeviceToHost - 从显卡内存复制到内存
    cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);

    // 在CUDA 中执行函数 语法：函数名称<<<block 数目, thread 数目, shared memory 大小>>>(参数...);
    sumOfSquares << < BLOCK_NUM, THREAD_NUM, 0 >> >(gpudata, result, time);

    /*把结果从显示芯片复制回主内存*/

    int sum[THREAD_NUM*BLOCK_NUM];

    clock_t time_use[BLOCK_NUM * 2];

    //cudaMemcpy 将结果从显存中复制回内存
    cudaMemcpy(&sum, result, sizeof(int)* THREAD_NUM*BLOCK_NUM, cudaMemcpyDeviceToHost);
    cudaMemcpy(&time_use, time, sizeof(clock_t)* BLOCK_NUM * 2, cudaMemcpyDeviceToHost);

    //Free
    cudaFree(gpudata);
    cudaFree(result);
    cudaFree(time);

    int final_sum = 0;

    for (int i = 0; i < THREAD_NUM*BLOCK_NUM; i++) {

        final_sum += sum[i];

    }

    //采取新的计时策略 把每个 block 最早的开始时间，和最晚的结束时间相减，取得总运行时间
    clock_t min_start, max_end;

    min_start = time_use[0];

    max_end = time_use[BLOCK_NUM];

    for (int i = 1; i < BLOCK_NUM; i++) {
        if (min_start > time_use[i])
            min_start = time_use[i];
        if (max_end < time_use[i + BLOCK_NUM])
            max_end = time_use[i + BLOCK_NUM];
    }

    printf("GPUsum: %d  gputime: %d\n", final_sum, max_end - min_start);

    final_sum = 0;

    for (int i = 0; i < DATA_SIZE; i++) {

        final_sum += data[i] * data[i] * data[i];

    }

    printf("CPUsum: %d \n", final_sum);

    return 0;
}

运行结果：

为了对比我们把block改成1再运行一次：

我们看到32block 256 thread 连续存取的情况下运行用了133133个时钟周期

而在 1block 256 thread 连续存取的情况下运行用了3488971个时钟周期

3488971/133133= 26.21倍

可以看到我们的速度整整提升了26倍，这个版本的程序，执行的时间减少很多。

我们还是从内存带宽的角度来进行一下评估：

首先计算一下使用的时间：

133133/ (797000 * 1000) = 1.67e-4S

然后计算使用的带宽：

数据量仍然没有变 DATA_SIZE 1048576，也就是1024*1024 也就是 1M

1M 个 32 bits 数字的数据量是 4MB。

因此，这个程序实际上使用的内存带宽约为：

4MB / 1.67e-4S = 23945.9788MB/s = 23.38GB/s

这对于我这块640，频率仅有797000，已经是一个很不错的效果了，不过，这个程序虽然在GPU上节省了时间，但是在 CPU 上执行的部份，需要的时间加长了（因为 CPU 现在需要加总 8192 个数字）。为了避免这个问题，下一步我们可以让每个 block 把自己的每个 thread 的计算结果进行加总。

关于更多线程的小实验，越多线程越好？

之前中间提过我们为什么不用更多的线程，比如一个block 1024个，或者更多的block。这是因为从线程运行的原理来看，线程数量达到一定大小后，我们再一味的增加线程也不会取得性能提升了，反而有可能会让性能下降。我们可以试验一下：

1024Thread *128block = 101372 个 Thread 够多了吧，我们看下运行结果：

我们看到最终用了153292个时钟周期，劲爆的10万个线程真的变慢了。

为什么会这样呢？下面我们从GPU的原理上来讲解这个问题。

从GPU结构理解线程：

之前关于为什么线程不能这么多的问题，说的还是不是很清楚，其实从硬件角度分析，支持CUDA的NVIDIA 显卡，都是由多个multiprocessors 组成。每个 multiprocessor 里包含了8个stream processors，其组成是四个四个一组，也就是两组4D的处理器。

每个 multiprocessor 还具有 很多个（比如8192个）寄存器，一定的（比如16KB） share memory，以及 texture cache 和 constant cache

在 CUDA 中，大部份基本的运算动作，都可以由 stream processor 进行。每个 stream processor 都包含一个 FMA（fused-multiply-add）单元，可以进行一个乘法和一个加法。比较复杂的运算则会需要比较长的时间。

在执行 CUDA 程序的时候，每个 stream processor 就是对应一个 thread。每个 multiprocessor 则对应一个 block。但是我们一个block往往有很大量的线程，题主用到了640个和1024个，远超一个 multiprocessor 所有的8个 stream processor 。

实际上，虽然一个 multiprocessor 只有八个 stream processor，但是由于 stream processor 进行各种运算都有 latency，更不用提内存存取的 latency，因此 CUDA 在执行程序的时候，是以warp 为单位。

比如一个 warp 里面有 32 个 threads，分成两组 16 threads 的 half-warp。由于 stream processor 的运算至少有 4 cycles 的 latency，因此对一个 4D 的stream processors 来说，一次至少执行 16 个 threads（即 half-warp）才能有效隐藏各种运算的 latency。也因此，线程数达到隐藏各种latency的程度后，之后数量的提升就没有太大的作用了。

还有一个重要的原因是，由于 multiprocessor 中并没有太多别的内存，因此每个 thread 的状态都是直接保存在multiprocessor 的寄存器中。所以，如果一个 multiprocessor 同时有愈多的 thread 要执行，就会需要愈多的寄存器空间。例如，假设一个 block 里面有 256 个 threads，每个 thread 用到20 个寄存器，那么总共就需要 256x20 = 5,120 个寄存器才能保存每个 thread 的状态。

而一般每个 multiprocessor 只有 8,192 个寄存器，因此，如果每个 thread 使用到16 个寄存器，那就表示一个 multiprocessor 的寄存器同时最多只能维持 512 个 thread 的执行。如果同时进行的 thread 数目超过这个数字，那么就会需要把一部份的数据储存在显卡内存中，就会降低执行的效率了。

总结：

这篇博客主要使用block进行了进一步增大了线程数，进行进一步的并行，最终的结果还是比较令人满意的，至少对于我这块显卡来说已经很不错了，因为我的显卡主频比较低，如果用一块1.5Ghz的显卡，用的时间就会是我的一半，而这时候内存带宽也就基本达到45GB/s左右了。

同时也回答了很多人都会有的疑问，即为什么我们不搞几万个线程。

但是我们也看到了新的问题，我们在CPU端的加和压力变得很大，那么我们能不能从GPU上直接完成这个工作呢？我们知道每个block内部的Thread之间是可以同步和通讯的，下一步我们将让每个block把每个thread的计算结果进行加和。

希望我的博客能帮助到大家~

参考资料：《深入浅出谈CUDA》

时间： 2024-10-04 11:08:11

CUDA编程（六）进一步并行的相关文章

CUDA C 编程指导（二）：CUDA编程模型详解

CUDA编程模型详解 本文以vectorAdd为例,通过描述C在CUDA中的使用(vectorAdd这个例子可以在CUDA sample中找到.)来介绍CUDA编程模型的主要概念.CUDA C的进一步描述可以参考<Programming Interface>. 主要内容包括: 1.Kernels(核函数) 2.Thread Hierarchy(线程结构) 3.Memory Hierarchy(存储结构) 4.Heterogeneous Programming(异构编程) 5.Compute C

CUDA编程（四）并行化我们的程序

CUDA编程(四) CUDA编程(四)并行化我们的程序 上一篇博客主要讲解了怎么去获取核函数执行的准确时间,以及如何去根据这个时间评估CUDA程序的表现,也就是推算所谓的内存带宽,博客的最后我们计算了在GPU上单线程计算立方和的程序的内存带宽,发现其内存带宽的表现是十分糟糕的,其所使用的内存带宽大概只有 5M/s,而像GeForce 8800GTX这样比较老的显卡,也具有超过50GB/s 的内存带宽 . 面对我们首先需要解决的内存带宽问题,我们首先来分析这个问题,然后我们将使用并行化来大大改善这

CUDA编程（二） CUDA初始化与核函数

CUDA编程(二) CUDA初始化与核函数 CUDA初始化 在上一次中已经说过了,CUDA安装成功之后,新建一个工程还是十分简单的,直接在新建项目的时候选择NVIDIA CUDA项目就可以了,我们先新建一个MyCudaTest 工程,删掉自带的示例kernel.cu,然后新建项,新建一个CUDA C/C++ File ,我们首先看一下如何初始化CUDA,因此我命名为InitCuda.cu 首先我们要使用CUDA的RunTime API 所以 我们需要include cuda_runtime.h

CUDA编程（三）评估CUDA程序的表现

CUDA编程(三) 评估CUDA程序的表现 上一篇博客我们基本上搭建起来了CUDA程序的骨架,但是其中并没有涉及到我们之前不断提到的并行加速,毕竟只有当我们的程序高并行的运行在GPU上才能大大缩短运行时间.不过在加速之前我们还有一件非常重要的事情需要考虑,那就是我们的程序到底有没有一个好的表现,也就是我们要准确计算程序的运行时间,这对之后的程序优化也有至关重要的作用,所以值得我们去仔细研究一下~ 这里所谓的计算运行时间也不是单纯意义上的看运行时间,更重要的是我们要通过核函数的运行时间去计算程序实

CUDA编程（五）关注内存的存取模式

CUDA编程(五) 关注内存的存取模式 上一篇博客我们使用Thread完成了简单的并行加速,虽然我们的程序运行速度有了50甚至上百倍的提升,但是根据内存带宽来评估的话我们的程序还远远不够, 除了通过Block继续提高线程数量来优化性能,这次想给大家先介绍一个访存方面非常重要的优化,同样可以大幅提高程序的性能~ 什么样的存取模式是高效的? 大家知道一般显卡上的内存是 DRAM,因此最有效率的存取方式,是以连续的方式存取,单纯说连续存取可能比较抽象,我们还是通过例子来看这个问题. 之前的程序,大家可

CUDA加速为什么没有并行的效果

问题描述 CUDA加速为什么没有并行的效果 我对一幅图像做处理,640X480,对每个像素我分配了一个thread,即我每个block分配了640个thread,分配了480 个block.然后我对另一幅2048X1536的图像操作,我每个block分配了1024个thread,共分配了1536X2个block.用 同样的算法,问题规模增加了10倍,我不理解的是为什么运行时间也增加了10倍.不是只会增加一点读取数据的时间吗. 对CUDA的运行原理不是很了解,忘大神指教. 解决方案 CUDA编程的

CUDA编程（一）第一个CUDA程序

CUDA编程(一) 第一个CUDA程序 Kernel.cu CUDA是什么? CUDA(Compute Unified Device Architecture),是显卡厂商NVIDIA推出的运算平台.是一种通用并行计算架构,该架构使GPU能够解决复杂的计算问题.说白了就是我们可以使用GPU来并行完成像神经网络.图像处理算法这些在CPU上跑起来比较吃力的程序.通过GPU和高并行,我们可以大大提高这些算法的运行速度. 有的同学可能知道,在CPU和GPU上跑同一个神经网络,由于其大量的浮点数权重计算以

CUDA入门（二）cuda编程的基本知识与第一个cuda程序

多首先,先来了解一下GPU与CPU的区别,如图 可以看到CPU(Central Processing Unit,中央处理单元),由Control(控制台),ALU(Arithmetic Logic Unit,逻辑计算单元),Cache(高速缓存),而GPU(Graphic Processing Unit,图形处理单元)也是由相同的部件组成,但GPU的计算单元远比CPU多,这就决定了GPU适合大量简单,精度要求低的计算,CPU则适合复杂的,精度要求高的计算.(如果还不了解,可以回学校恶补一下微机原

《CUDA C编程权威指南》——2.1节CUDA编程模型概述

2.1 CUDA编程模型概述 CUDA编程模型提供了一个计算机架构抽象作为应用程序和其可用硬件之间的桥梁.图2-1说明了程序和编程模型实现之间的抽象结构的重要.通信抽象是程序与编程模型实现之间的分界线,它通过专业的硬件原语和操作系统的编译器或库来实现.利用编程模型所编写的程序指定了程序的各组成部分是如何共享信息及相互协作的.编程模型从逻辑上提供了一个特定的计算机架构,通常它体现在编程语言或编程环境中. 除了与其他并行编程模型共有的抽象外,CUDA编程模型还利用GPU架构的计算能力提供了以下几个特