首页 > TAG信息列表 > threadIdx
cuda 编 程（三） helloworld 打印 blockIdx和threadIdx.x threadIdx.y
#include <stdio.h> #include <iostream> using namespace std; __global__ void hello_from_gpu() { const int b = blockIdx.x; const int tx = threadIdx.x; const int ty = threadIdx.y; // cout<<b<<endl; printf("HelCUDA C编程（十四）合并的全局内存访问
使用共享内存也能帮助避免对未合并的全局内存的访问，矩阵转置就是一个典型的例子：读操作被自然合并，但写操作是按照交叉访问的。其中交叉访问是全局内存中最糟糕的访问模式，因为它浪费总线带宽，在共享内存的帮助下，可以先在共享内存中进行转置操作，然后再对全局内存进行合并写操CUDA快速入门——学习总结
**写在前面：**本文是作者总结的cuda基础部分重点内容，一些琐碎的重要细节部分没有提及，建议有需要的同学自行学习。 1. 基本概念 1.1 核函数、网格、线程块与线程 启动一个核函数，即调用一个网格（grid） kernal<<<block, thread>>> (args); grid中包含很多线程块（blocks），block中包含三种基于CUDA的归约实现
归约在并行计算中很常见，并且在实现上具有一定的套路。本文分别基于三种机制（Intrinsic，共享内存，atomic），实现三个版本的归约操作，完成一个warp（32）大小的整数数组的归约求和计算。 Intrinsic版本 基于Intrinsic函数 __shfl_down_sync 实现，使一个warp内的线程通过读取相邻线程寄存器test-case
#define FINAL_MASK 0xffffffff template <typename T> __inline__ __device__ T warpReduceSum(T val) { for(int mask = 16; mask > 0; mask >>= 1) val += __shfl_xor_sync(FINAL_MASK, val, mask, 32); return val; } template <typename T>CUDA学习2-编程部分
CUDA编程 函数声明 host：主机端，通常指CPU device：设备端，通常指GPU（数据可并行） kernel：数据并行处理函数，在主机端调用kernel可以在设备端创建大量轻量级线程 host 和 device拥有各自的存储器 CUDA编程包括主机端和设备端两部分代码 执行位置调用位置_device_float DeviceFunc()d如何使用TensorCores优化卷积
如何使用TensorCores优化卷积 本文将演示如何在TVM中使用TensorCores编写高性能的卷积计划。假设卷积的输入有大量数据。首先介绍如何在GPU上优化卷积。 TensorCore简介 每个Tensor核心都提供一个4x4x4的矩阵处理阵列，该阵列可以运行 ，其中A，B，C和D是4x4矩阵，如图所示。矩阵乘法输Nsight Compute内存访问常用Metrics含义理解
Nsight Compute 软件Source模块提供了精确到源代码行号的metrics参数，用于辅助性能调优，本篇基于访问共享内存的矩阵转置核函数的实现，记录一下对常用metrics含义的理解。 Metrics含义 Memory L1 Transcations Global:实际全局内存加载至L1缓存的内存交换次数，粒度128bytes Memor