【高性能计算】CUDA编程之线程存储与原子操作（教程与代码-2）

最新推荐文章于 2024-08-12 01:01:56 发布
最新推荐文章于 2024-08-12 01:01:56 发布
阅读量895 收藏 2
点赞数 1
分类专栏： CUDA 文章标签： 人工智能 深度学习 c语言
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/heroybc/article/details/123130714
版权

3 线程、同步、存储器

3.1 线程与存储

  • tid=blockIdx.x（当前块的ID）*blockDim.x（当前块里面的线程数量）+threadIdx.x（当前线程在块中的ID）
  • gridDim.x*blockDim.x来计算，前者代表了本次启动的块的数量，而后者代表了每个块里面的线程数量，然后每次while循环，tid变量加上这个值，向后偏移以得到下个任务的索引
  • 所有线程都有一个寄存器堆，它是最快的。
  • 共享内存只能被块中的线程访问，但比全局内存块。
  • 全局内存是最慢的，但可以被所有的块访问。
  • 常量和纹理内存用于特殊用途
  • 所有通过cudaMalloc分配的存储器都是全局内存
  • 本地内存和寄存器堆对每个线程都是唯一的。
  • 寄存器是每个线程可用的最快存储器。
  • 当内核中使用的变量在寄存器堆中装不下的时候，将会使用本地内存存储它们，这叫寄存器溢出
  • “读取旧值-累加-回写新值”操作是不可被其他线程扰乱的原子性的整体完成的。
  • 使用atomicAdd进行原子累加的内核函数
  • 使用原子操作后程序具有更大的执行代价。可以通过使用共享内存来加速这些原子累加操作
    GPU卡从逻辑上对用户提供了64KB的常量内存空间，可以用来存储内核执行期间所需要的恒定数据
  • 常量内存有助于节省全局内存的访问带宽
  • warp整体进行一次常量内存的读取，结果广播给warp里的32个线程。同时，常量内存具有cache缓冲。当后续的在邻近位置上访问，将不会发生额外的从显存过来的传输。每个warp里的32个线程，进行一致性的相同常量内存位置读取的时候，这种广播效果和cache命中效果可以节省执行时间
  • 当程序进行具有很大程度上的空间邻近性的访存的时候，纹理变得非常高效。空间邻近性的意思是，每个线程的读取位置都和其他线程的读取位置邻近。
  • 请一定要确保纹理引用被定义成全局静态变量，同时还要确保它不能作为参数传递给任何其他函数。
  • 原子操作-求和
#include <stdio.h>
#define NUM_THREADS 10000
#define SIZE  10
#define BLOCK_WIDTH 100
__global__ void gpu_increment_without_atomic(int *d_a)
{
    // Calculate thread id for current thread
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    // each thread increments elements wrapping at SIZE variable
    tid = tid % SIZE;
    d_a[tid] += 1;
}
__global__ void gpu_increment_atomic(int *d_a)
{
    // Calculate thread id for current thread
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    // each thread increments elements wrapping at SIZE variable
    tid = tid % SIZE;
    atomicAdd(&d_a[tid], 1);
}
int main(int argc, char **argv)
{
    printf("%d total threads in %d blocks writing into %d array elements\n",
        NUM_THREADS, NUM_THREADS / BLOCK_WIDTH, SIZE);
    // declare and allocate host memory
    int h_a[SIZE];
    const int ARRAY_BYTES = SIZE * sizeof(int);
    // declare and allocate GPU memory
    int *d_a, *d_aA;
    cudaMalloc((void **)&d_a, ARRAY_BYTES);
    //Initialize GPU memory to zero
    cudaMemset((void *)d_a, 0, ARRAY_BYTES);
    gpu_increment_without_atomic << <NUM_THREADS / BLOCK_WIDTH, BLOCK_WIDTH >> >(d_a);
    // copy back the array to host memory
    cudaMemcpy(h_a, d_a, ARRAY_BYTES, cudaMemcpyDeviceToHost);
    printf("Number of times a particular Array index has been incremented without atomic add is: \n");
    for (int i = 0; i < SIZE; i++)
    {
        printf("index: %d --> %d times\n ", i, h_a[i]);
    }
    cudaFree(d_a);
    
    cudaMalloc((void **)&d_aA, ARRAY_BYTES);
    //Initialize GPU memory to zero
    cudaMemset((void *)d_aA, 0, ARRAY_BYTES);
    gpu_increment_atomic << <NUM_THREADS / BLOCK_WIDTH, BLOCK_WIDTH >> >(d_aA);
    // copy back the array to host memory
    cudaMemcpy(h_a, d_aA, ARRAY_BYTES, cudaMemcpyDeviceToHost);
    printf("Number of times a particular Array index has been incremented is: \n");
    for (int i = 0; i < SIZE; i++) 
    { 
        printf("index: %d --> %d times\n ", i, h_a[i]); 
    }
    cudaFree(d_aA);
    return 0;
}
  • 常量内存与纹理内存
#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
//Defining two constants
__constant__ int constant_f;
__constant__ int constant_g;
#define N   5
//Kernel function for using constant memory
__global__ void gpu_constant_memory(float *d_in, float *d_out) {
    //Thread index for current kernel
    int tid = threadIdx.x;  
    d_out[tid] = constant_f*d_in[tid] + constant_g;
}
#define NUM_THREADS 5
texture <float, 1, cudaReadModeElementType> textureRef;
__global__ void gpu_texture_memory(int n, float *d_out)
{
    int idx = blockIdx.x*blockDim.x + threadIdx.x;
    if (idx < n) {
        float temp = tex1D(textureRef, float(idx));
        d_out[idx] = temp;
    }
}
int main(void) {
    //Defining Arrays for host
    float h_in[N], h_out[N];
    //Defining Pointers for device
    float *d_in, *d_out;
    // 常量内存
    int h_f = 2;
    int h_g = 20;
    // allocate the memory on the cpu
    cudaMalloc((void**)&d_in, N * sizeof(float));
    cudaMalloc((void**)&d_out, N * sizeof(float));
    //Initializing Array
    for (int i = 0; i < N; i++) {
        h_in[i] = i;
    }
    //Copy Array from host to device
    cudaMemcpy(d_in, h_in, N * sizeof(float), cudaMemcpyHostToDevice);
    //Copy constants to constant memory
    cudaMemcpyToSymbol(constant_f, &h_f, sizeof(int),0,cudaMemcpyHostToDevice);
    cudaMemcpyToSymbol(constant_g, &h_g, sizeof(int));
    //Calling kernel with one block and N threads per block
    gpu_constant_memory << <1, N >> >(d_in, d_out);
    //Coping result back to host from device memory
    cudaMemcpy(h_out, d_out, N * sizeof(float), cudaMemcpyDeviceToHost);
    //Printing result on console
    printf("Use of Constant memory on GPU \n");
    for (int i = 0; i < N; i++) {
        printf("The expression for input %f is %f\n", h_in[i], h_out[i]);
    }
    //Free up memory
    cudaFree(d_in);
    cudaFree(d_out);
    // 纹理内存
    //Calculate number of blocks to launch
    int num_blocks = N / NUM_THREADS + ((N % NUM_THREADS) ? 1 : 0);
    //Declare device pointer
    float *d_outM;
    // allocate space on the device for the result
    cudaMalloc((void**)&d_outM, sizeof(float) * N);
    // allocate space on the host for the results
    float *h_outM = (float*)malloc(sizeof(float)*N);
    //Declare and initialize host array
    float h_inM[N];
    for (int i = 0; i < N; i++) {
        h_inM[i] = float(i);
    }
    //Define CUDA Array
    cudaArray *cu_Array;
    cudaMallocArray(&cu_Array, &textureRef.channelDesc, N, 1);
    //Copy data to CUDA Array
    cudaMemcpyToArray(cu_Array, 0, 0, h_inM, sizeof(float)*N, cudaMemcpyHostToDevice);
    // bind a texture to the CUDA array
    cudaBindTextureToArray(textureRef, cu_Array);
    //Call Kernel   
    gpu_texture_memory << <num_blocks, NUM_THREADS >> >(N, d_outM);
    
    // copy result back to host
    cudaMemcpy(h_outM, d_outM, sizeof(float)*N, cudaMemcpyDeviceToHost);
    printf("Use of Texture memory on GPU: \n");
    for (int i = 0; i < N; i++) {
        printf("Texture element at %d is : %f\n",i, h_outM[i]);
    }
    free(h_outM);
    cudaFree(d_outM);
    cudaFreeArray(cu_Array);
    cudaUnbindTexture(textureRef);
    return 0;
}
  • 全局、局部、共享内存
#include <stdio.h>
#define N 5
__global__ void gpu_global_memory(float *d_a)
{
    // "array" is a pointer into global memory on the device
    d_a[threadIdx.x] = threadIdx.x;
}
__global__ void gpu_local_memory(float d_in)
{
    int t_local;    
    t_local = d_in * threadIdx.x;     
    printf("Value of Local variable in current thread is: %d \n", t_local);
}
__global__ void gpu_shared_memory(float *d_a)
{
    // Defining local variables which are private to each thread
    int i, index = threadIdx.x;
    float average, sum = 0.0f;
    //Define shared memory
    __shared__ float sh_arr[5];
    sh_arr[index] = d_a[index];
    __syncthreads();    // This ensures all the writes to shared memory have completed
    for (i = 0; i<= index; i++) 
    { 
        sum += sh_arr[i]; 
    }
    average = sum / (index + 1.0f);
    d_a[index] = average; 
    sh_arr[index] = average;
}
int main(int argc, char **argv)
{
    // Define Host Array
    float h_a[N];
    //Define device pointer 
    float *d_a;       
    // 全局内存
    cudaMalloc((void **)&d_a, sizeof(float) *N);
    // now copy data from host memory to device memory 
    cudaMemcpy((void *)d_a, (void *)h_a, sizeof(float) *N, cudaMemcpyHostToDevice);
    // launch the kernel 
    gpu_global_memory << <1, N >> >(d_a);  
    // copy the modified array back to the host memory
    cudaMemcpy((void *)h_a, (void *)d_a, sizeof(float) *N, cudaMemcpyDeviceToHost);
    printf("Array in Global Memory is: \n");
    //Printing result on console
    for (int i = 0; i < N; i++) {
        printf("At Index: %d --> %f \n", i, h_a[i]);
    }
    // 本地内存
    printf("Use of Local Memory on GPU:\n");
    gpu_local_memory << <1, N >> >(5);  
    cudaDeviceSynchronize();
      
    // 共享内存 
    for (int i = 0; i < 5; i++) {
        h_a[i] = i;
    }
    // allocate global memory on the device
    cudaMalloc((void **)&d_a, sizeof(float) * 5);
    // now copy data from host memory  to device memory 
    cudaMemcpy((void *)d_a, (void *)h_a, sizeof(float) * 5, cudaMemcpyHostToDevice);
    
    gpu_shared_memory << <1, 5 >> >(d_a);
    // copy the modified array back to the host memory
    cudaMemcpy((void *)h_a, (void *)d_a, sizeof(float) * 5, cudaMemcpyDeviceToHost);
    printf("Use of Shared Memory on GPU:  \n");
    //Printing result on console
    for (int i = 0; i < 5; i++) {
        printf("The running average after %d element is %f \n", i, h_a[i]);
    }
    
    return 0;
}
  • 多线程
#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
//Defining number of elements in Array
#define N   50000
//Defining Kernel function for vector addition
__global__ void gpuAdd(int *d_a, int *d_b, int *d_c) {
    //Getting block index of current kernel
    int tid = threadIdx.x + blockIdx.x * blockDim.x;    
    while (tid < N)
    {
        d_c[tid] = d_a[tid] + d_b[tid];
        tid += blockDim.x * gridDim.x;
    }
        
}
int main(void) {
    //Defining host arrays
    int h_a[N], h_b[N], h_c[N];
    //Defining device pointers
    int *d_a, *d_b, *d_c;
    // allocate the memory
    cudaMalloc((void**)&d_a, N * sizeof(int));
    cudaMalloc((void**)&d_b, N * sizeof(int));
    cudaMalloc((void**)&d_c, N * sizeof(int));
    //Initializing Arrays
    for (int i = 0; i < N; i++) {
        h_a[i] = 2 * i*i;
        h_b[i] = i;
    }
    // Copy input arrays from host to device memory
    cudaMemcpy(d_a, h_a, N * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, N * sizeof(int), cudaMemcpyHostToDevice);
    //Calling kernels with N blocks and one thread per block, passing device pointers as parameters
    gpuAdd << <512, 512 >> >(d_a, d_b, d_c);
    //Copy result back to host memory from device memory
    cudaMemcpy(h_c, d_c, N * sizeof(int), cudaMemcpyDeviceToHost);
    cudaDeviceSynchronize();
    int Correct = 1;
    printf("Vector addition on GPU \n");
    //Printing result on console
    for (int i = 0; i < N; i++) {
        if ((h_a[i] + h_b[i] != h_c[i]))
        {
            Correct = 0;
        }
        
    }
    if (Correct == 1)
    {
        printf("GPU has computed Sum Correctly\n");
    }
    else
    {
        printf("There is an Error in GPU Computation\n");
    }
    //Free up memory
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    return 0;
}

3.2 向量点乘与矩阵乘法例子

  • 向量点乘
  • 矩阵乘法
#include "stdio.h"
#include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>
#include <math.h>
#define TILE_SIZE 2
#define N 1024
#define threadsPerBlock 512
__global__ void gpu_dot(float *d_a, float *d_b, float *d_c) {
    //Declare shared memory
    __shared__ float partial_sum[threadsPerBlock];
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    //Calculate index for shared memory 
    int index = threadIdx.x;
    //Calculate Partial Sum
    float sum = 0;
    while (tid < N) 
    {
        sum += d_a[tid] * d_b[tid];
        tid += blockDim.x * gridDim.x; // 前者代表了本次启动的块的数量，而后者代表了每个块里面的线程数量
    }
    // Store partial sum in shared memory
    partial_sum[index] = sum;
    // synchronize threads 
    __syncthreads();
    // Calculating partial sum for whole block in reduce operation  
    int i = blockDim.x / 2;
    while (i != 0) {
        if (index < i)
            partial_sum[index] += partial_sum[index + i];
        __syncthreads();
        i /= 2;
    }
    //Store block partial sum in global memory
    if (index == 0)
        d_c[blockIdx.x] = partial_sum[0];
}
//Matrix multiplication using non shared kernel
__global__ void gpu_Matrix_Mul_nonshared(float *d_a, float *d_b, float *d_c, const int size)
{
    int row, col;
    col = TILE_SIZE * blockIdx.x + threadIdx.x;
    row = TILE_SIZE * blockIdx.y + threadIdx.y;
    for (int k = 0; k< size; k++)
    {
        d_c[row*size + col] += d_a[row * size + k] * d_b[k * size + col];
    }
}
// Matrix multiplication using shared kernel
__global__ void gpu_Matrix_Mul_shared(float *d_a, float *d_b, float *d_c, const int size)
{
    int row, col;
    //Defining Shared Memory
    __shared__ float shared_a[TILE_SIZE][TILE_SIZE];
    __shared__ float shared_b[TILE_SIZE][TILE_SIZE];
    col = TILE_SIZE * blockIdx.x + threadIdx.x;
    row = TILE_SIZE * blockIdx.y + threadIdx.y;
    for (int i = 0; i< size / TILE_SIZE; i++) 
    {
        shared_a[threadIdx.y][threadIdx.x] = d_a[row* size + (i*TILE_SIZE + threadIdx.x)];
        shared_b[threadIdx.y][threadIdx.x] = d_b[(i*TILE_SIZE + threadIdx.y) * size + col];
        __syncthreads(); 
        for (int j = 0; j<TILE_SIZE; j++)
            d_c[row*size + col] += shared_a[threadIdx.y][j] * shared_b[j][threadIdx.x];
        __syncthreads(); 
    }
}
int main_dot(void) {
    //Declare Host Array
    float *h_a, *h_b, h_c, *partial_sum;
    //Declare device Array
    float *d_a, *d_b, *d_partial_sum;
    //Calculate total number of blocks per grid
    int block_calc = (N + threadsPerBlock - 1) / threadsPerBlock;
    int blocksPerGrid = (32 < block_calc ? 32 : block_calc);
    // allocate memory on the host side
    h_a = (float*)malloc(N * sizeof(float));
    h_b = (float*)malloc(N * sizeof(float));
    partial_sum = (float*)malloc(blocksPerGrid * sizeof(float));
    // allocate the memory on the device
    cudaMalloc((void**)&d_a, N * sizeof(float));
    cudaMalloc((void**)&d_b, N * sizeof(float));
    cudaMalloc((void**)&d_partial_sum, blocksPerGrid * sizeof(float));
    // fill the host array with data
    for (int i = 0; i<N; i++) {
        h_a[i] = i;
        h_b[i] = 2;
    }
    cudaMemcpy(d_a, h_a, N * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, N * sizeof(float), cudaMemcpyHostToDevice);
    //Call kernel 
    gpu_dot << <blocksPerGrid, threadsPerBlock >> >(d_a, d_b, d_partial_sum);
    // copy the array back to host memory
    cudaMemcpy(partial_sum, d_partial_sum, blocksPerGrid * sizeof(float), cudaMemcpyDeviceToHost);
    // Calculate final dot product on host
    h_c = 0;
    for (int i = 0; i<blocksPerGrid; i++) {
        h_c += partial_sum[i];
    }
    printf("The computed dot product is: %f\n", h_c);
    #define cpu_sum(x) (x*(x+1))
    if (h_c == cpu_sum((float)(N - 1)))
    {
        printf("The dot product computed by GPU is correct\n");
    }
    else
    {
        printf("Error in dot product computation");
    }
    
    // free memory on host and device
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_partial_sum);
    free(h_a);
    free(h_b);
    free(partial_sum);
    return 0;
}
int main_matrix()
{
    const int size = 4;
    //Define Host Array
    float h_a[size][size], h_b[size][size],h_result[size][size];
    //Defining device Array
    float *d_a, *d_b, *d_result; 
    //Initialize host Array
    for (int i = 0; i<size; i++)
    {
        for (int j = 0; j<size; j++)
        {
            h_a[i][j] = i;
            h_b[i][j] = j;
        }
    }
    cudaMalloc((void **)&d_a, size*size*sizeof(int));
    cudaMalloc((void **)&d_b, size*size * sizeof(int));
    cudaMalloc((void **)&d_result, size*size* sizeof(int));
    //copy host array to device array
    cudaMemcpy(d_a, h_a, size*size* sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, size*size* sizeof(int), cudaMemcpyHostToDevice);
    
    //Define grid and block dimensions
    dim3 dimGrid(size / TILE_SIZE, size / TILE_SIZE, 1);
    dim3 dimBlock(TILE_SIZE, TILE_SIZE, 1);
    //gpu_Matrix_Mul_nonshared << <dimGrid, dimBlock >> > (d_a, d_b, d_result, size);
    gpu_Matrix_Mul_shared << <dimGrid, dimBlock >> > (d_a, d_b, d_result, size);
    cudaMemcpy(h_result, d_result, size*size * sizeof(int), cudaMemcpyDeviceToHost);
    printf("The result of Matrix multiplication is: \n");
    
    for (int i = 0; i< size; i++)
    {
        for (int j = 0; j < size; j++)
        {
            printf("%f   ", h_result[i][j]);
        }
        printf("\n");
    }
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_result);
    return 0;
}
int main(){
    main_matrix();
}
Carlos橘子柠檬
关注
专栏目录
第八章 CUDA内存应用与性能优化篇(上篇)
weixin_38252409的博客
08-13 2375
学习我的教程专栏，你将绝对能实现CUDA工程化，实现环境安装、index计算、kernel核函数编程、内存优化与steam性能优化、原子操作、nms的cuda算子、yolov5的cuda部署等内容，并开源教程源码。以上章节中，我们已经比较熟练掌握如何使用cuda编写自己想要的计算逻辑，已能成功编写cuda代码了。那么，另外一个重要问题值得我们关注，如何优化其性能，使其工程部署能加速运行了。而这种性能优化与cuda内存密切相关。为此，我们在本节中介绍cuda内存相关内容，并附其源码。
GPU 高性能编程 CUDA ：原子性
xxiaozr的博客
05-18 723
NVIDIA GPU 的计算功能集：NVIDIA 将 GPU 支持的各种功能统称为计算功能集nvcc -arch=sm_11 //告诉编译器需要1.1版本或者更高 nvcc -arch=sm_12 //告诉编译器需要1.2版本或者更高原子操作简介：在执行过程中不能分解为更小的部分，将满足这种条件限制的操作称为原子操作，当有数千个线程在内存访问上发生竞争时，这些操作可以保证在内存上实现安全的操作...
阿里云高性能计算（HPC）使用教程
weixin_40050195的博客
10-25 1922
关于阿里云高性能计算（HPC）的详细了解：阿里云高性能计算（HPC）使用教程 （高性能计算 (Alibaba Cloud HPC) 提供一种性能卓越、稳定、安全、便捷的计算服务，帮助您快速构建处理能力出色的应用，解放计算给服务带来的压力，使您的产品在计算效率上具有非凡竞争力。） 产品概述       高性能计算 (Alibaba Cloud HPC) 提供一种性能卓越、稳定、安全、便捷的计算...
高性能计算-devcloud平台使用简易教程与体验
hhhny的博客
04-22 1156
Intel® DevCloud平台为每个注册用户提供独立的沙箱运行环境，并构建了完整的软件栈，开发者登录后不需要对环境做任何配置，不用担心软件之间的兼容问题，即可使用浏览器随时随地访问Intel® DevCloud上的算力资源、官方提供的丰富教程和参考实例。“开箱即用”的用户体验，想开发者所想！Intel® DevCloud容器化应用部署平台已经上线，开发者可以在丰富的边缘节点上运行容器化应用，完成功能测试和基准测试。
cuda原子操作
IChocolateKapa
02-07 2784
C.1  数学函数C.1.1  atomicAdd() int atomicAdd(int* address, int val); unsigned int atomicAdd(unsigned int* address,                            unsigned int val); unsigned long long int atomicAdd(unsig
CUDA编程06 - 性能优化指南
最新发布
GPU全栈博主-程序哥
08-12 1198
并行程序的执行速度在很大程度上取决于程序的资源需求与硬件的资源限制。在几乎所有并行编程模型中，管控并行代码与硬件资源约束之间的相互影响对于实现高性能非常重要的。这是一种实用的技能，需要对硬件体系结构有深刻理解，并需要在为高性能设计的并行编程模型下不断练习。到目前为止，我们已经了解了GPU架构的各个方面及其对性能的影响。在前面的CUDA编程04 - GPU计算架构和线程调度。
CUDA编程01- 并行编程介绍
GPU全栈博主-程序哥
07-17 2225
自计算机诞生以来，许多高端应用程序对执行速度和资源的需求超出了计算机所能提供的能力范围。早期的应用依赖于处理器速度、存储速度、存储容量的提升来增强应用程序的能力，例如，天气预报的及时性、工程结构分析的准确性、计算机生成图形的逼真度、每秒处理的机票预订数量，以及每秒处理的资金转账数量。最近，深度学习等新应用对执行速度和资源的要求甚至超过了最好的计算能力。这些应用需求在过去的五十年里推动了计算设备能力的快速发展，并将在可预见的未来继续如此。
CUDA编程实战：GPU高性能优化指南
"《GPU高性能编程CUDA实战中文版》是一本详细介绍CUDA编程技术的书籍，旨在帮助读者理解和利用CUDA架构实现高性能计算。CUDA是NVIDIA公司推出的一种编程模型，允许开发者利用GPU的强大并行计算能力来加速计算密集型...
如何学习高性能计算
haimianjie2012的专栏
03-18 1417
High performance computing 高性能计算 定义 高性能计算（HPC）是利用先进的计算机技术，如超级计算机和并行处理算法，来解决复杂的计算问题。主要用于科学、军事和学术领域。现在也被利用在数据存储和事务处理领域，以及为统计目的挖掘或使用这些信息。 "深蓝"是一种超级计算机，它通过与人类大师对弈来展示人工智能。 高性能计算的主要目标是以精确、快速、高效的方式解决计算问题，并行处理是高性能计算成功的关键。并行处理使用多个中央处理器（CPU）来计算一个问题，在技术上可以使计算速度更快。其
cuda原子操作实例(图像直方图统计、图像像素求和)
热门推荐
thecentury的博客
12-14 10万+
1.cuda直方图统计 /////代码一 #include <iostream> #include "cuda_runtime.h" #include "time.h" using namespace std; #define num (256 * 1024 * 1024) // 核函数 // 注意，为了方便验证GPU的统计结果，这里采用了"逆直方图"， // 即每发现一个数字，...
CUDA-原子操作及优化
qq_52230866的博客
07-12 484
1、为了避免不同线程同时对数据进行读取/运算/写回导致结果错误，原子操作可以避免这种错误，当前面线程完成这三步操作后，后面的线程才会运行。返回旧值，原子操作非常耗时，因为是串行的,，新版本cuda很聪明会优化，所以有时候如下这么写也没问题。原子的判断是否相等，相等则写入：atomicCAS(dst, cmp, src)加法：atomicAdd(dst, src) : *dst += src;原子写入并读取旧值：atomicExch()按位异或：atomicSubXor()按位异或：atomicMin()
高性能计算学习资料简单汇总
qq_30182357的博客
03-31 441
高性能计算学习资料视频课简单汇总 视频课简单汇总 新竹清华大学：并行计算与并行编程课程 链接：https://www.bilibili.com/video/BV1Yt411W7td?p=15&spm_id_from=pageDriver GPU并行计算与CUDA编程（2017） 链接：https://www.bilibili.com/video/BV1z4411a7Ny?p=2&spm_id_from=pageDriver NVIDIA CUDA初级教程视频 链接：https:/
CUDA学习笔记（五）——原子性操作
学习笔记
03-31 1262
说明 原子性操作不是所有的显卡都支持的，只有算力在1.1及以上的才支持全局内存的原子操作，只有1.2及以上算力的才支持共享内存上的原子操作。 由于显卡算力高的是低的一个超集，类似可理解为于高算力版本可以向下兼容。同时，在编译代码的时候，可以告诉nvcc编译器需要什么样的版本，nvcc -arch=sm_1.2即需要1.2算力版本的支持。 如果不支持就需要考虑升级硬件了。 原子操作 原子操作的概念是由多部操作构成的一个操作集合，如果执行该操作集合，就必须全部执行，不能执行部分操作。 如： x++; 定义上述
cuda-使用cuda并行加速实现之FastAtomicAdd.zip
04-24
CUDA（Compute Unified Device Architecture）是NVIDIA推出的一种编程模型，用于利用GPU（Graphics Processing Unit）的强大计算能力来处理科学计算、机器学习、图像处理等高性能计算任务。在CUDA中，FastAtomicAdd...
VS2017 CUDA编程学习12：CUDA流
DU_YULIN的博客
11-12 3243
文章目录前言1. CUDA流的理解2. C++ 实现CUDA流3. 执行结果总结学习资料 VS2017 CUDA编程学习1：CUDA编程两变量加法运算 VS2017 CUDA编程学习2：在GPU上执行线程 VS2017 CUDA编程学习3：CUDA获取设备上属性信息 VS2017 CUDA编程学习4：CUDA并行处理初探 - 向量加法实现 VS2017 CUDA编程学习5：CUDA并行执行-线程 VS2017 CUDA编程学习6: GPU存储器架构 VS2017 CUDA编程学习7：线程同步-共享内存 VS
CUDA编程模型系列八(原子操作 / 规约 / 向量元素求和)
专注于人工智能领域的小何尚
07-01 1337
本期我介绍了cuda 当中规约算法的一种情况, 也是小何尚职业生涯中的第一道面试题, 计算数组中所有元素的和.本系列视频目的是帮助开发者们一步步地学会利用CUDA编程模型加速GPU应用, 我们的口号是: 让GPU飞起来。CUDA编程模型系列八(原子操作 / 规约 / 向量元素求和)
高性能计算课程 ppt
05-06
高性能计算课程
CUDA原子操作【参加CUDA线上训练营】
Mr_8j的博客
02-09 403
CUDA原子操作及求和计算【参加CUDA线上训练营】
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值