列出帖子

该操作将允许你查看该会员所有的帖子,注意你只能看到你有权限看到的板块的帖子。


只显示主题 - jinyer

页: [1]
1
CUDA / 使用了共享内存最终时间没有没有减少
« 于: 十月 19, 2022, 03:23:15 pm »
程序中 kernel 1得到u值,kernel 2和kernel 3都要使用这个u值,因此我申请了全局可见的共享内存用来存储u,以减少2、3步中读取u的延时,可是为什么使用共享内存优化之后runtime没减少呢?
如下图,可以明显看出访问全局内存的次数明显变少,但时间却从1.96ms变成了1.97ms【这是使用nsight compute记录的核函数执行时间】

2
CUDA / GPU寄存器使用
« 于: 九月 07, 2022, 03:52:21 pm »
手册上说若寄存器不足,会使用本地内存,而本地内存位于全局内存上,理论上不会有限制。
3090卡1个SM理论可以同时调度1536个线程,可是在实际使用中之恩那个调度1152个线程,通过ncu查看当前函数执行时一个线程使用寄存器56个【65536/1152约等于56】.这说明寄存器限制了SM可同时调度的线程数量。
请问为什么没有使用本地内存呢?

3
CUDA / 线程分化-以warp为分支单位
« 于: 五月 06, 2022, 07:32:19 pm »
对开发者而言,应该尽量避免在同一个Warp中有不同的执行路径。当必须写分支语句时,尽量让分支宽度大于WarpSize,即不以Thread为分支单位,而是以Warp为分支单位,这样就能保证同一个Warp内部不会出现分支。
请问如果分配线程块 block(256)
核函数中是
程序代码: [选择]
int warpIdx = threadIdx.x/32;
if(warpIdx < 2){
    执行语句1;
}else{
    执行语句2;
}
代码是以warp为分支单位,那这种情况下if和else可以同时运行吗?
即现在有一个需求,想让线程块中一部分线程来执行计算,另一部分线程来把数据从全局内存搬运到共享内存,请问我可以通过上述if-else达到这两种需求同时运行的效果吗?

4
CUDA / Nsight Compute没有权限问题(linux系统上)
« 于: 四月 13, 2022, 09:34:44 am »
遇到的问题:launch之后每个核函数都会出现ERR_NVGPUCTRPERM - The user does not have permission to profile on the target device这个错误
查阅解决方法大家一般都是两点:
1、使用管理员权限运行nsight compute,但这好像是针对windows系统,对linux不太适用;
2、https://developer.nvidia.com/nvidia-development-tools-solutions-err_nvgpuctrperm-permission-issue-performance-counters官网给出了启用该权限的方法,但我按步骤做的时候这一步会出错:modprobe -r nvidia_uvm nvidia_drm nvidia_modeset nvidia-vgpu-vfio nvidia[2] # Unload dependent modules
,不知道要怎么做,请教大家

5
CUDA / SM、block、warp
« 于: 四月 01, 2022, 03:56:39 pm »
一个Multi-Processor上面可以同时驻留多个warps,例如一个计算能力8.6的卡,它的1个SM里面,最多可能驻留1536个线程,也就是大约48个warps,在同时等待被调度执行。并且同时最多能上16个block。
请问一个SM是同时在执行16个block吗?
若是的话,SM一个时刻只能执行每个block里边的一个warp吗?即并行线程数为 82*16*32吗?【目前我的理解是这样,可3090卡的SP个数只有10496是82*16*32的四分之一,这又是怎么做到的呢】

6
CUDA / 同一个CUDA流操作
« 于: 三月 16, 2022, 11:37:08 am »
看到大家说:同一个CUDA流操作是严格顺序的。
但是我跑代码时却出现了不一样的结果
程序代码: [选择]
noisePdf1<<<gridNoiPdf1, blockNoiPdf1 >>>(noiPdf1, d_rx, sigma);
printf("After---------------noisePdf1\n");
convPy1<<<gridConvPy, blockConvPy >>>(convResult, u0, u1);
printf("After------------------convPy\n");
然后分别在两个核函数中有一个printf语句。
在默认流中先后调用这两个核函数,结果输出语句如下:
After---------------noisePdf1
After------------------convPy
run----noisePdf1
running------convPy
请问这是表示如果第二个核函数想用第一个核函数的结果,这两个核函数调用中间必须加cudaDeviceSynchronize();吗?

7
CUDA / CUDA中FAM指令问题
« 于: 一月 13, 2022, 11:01:31 am »
请教大家:CUDA中编译器会自动优化a*b+c为FMA单指令吗?还是只有使用fma(a,b,c)才会融合乘加为单指令

8
CUDA / Linux系统:cuFFT使用错误 undefined reference to XXX
« 于: 十二月 23, 2021, 10:23:26 am »
linux系统在nsight edition中使用cufft库函数,已经#include <cufft.h>,还是出现以下错误:
     undefined reference to `cufftPlan1d'
查阅后发现windows下解决方法如下:
    项目——属性——平台选x64——链接器——输入——附加依赖项 在这个里面加入cufft.lib这个名字。
但是linux中的IDE没有这个选项,请教一下要怎么改正?

9
CUDA / CUDA多个核函数kernel之间的数据如何传递
« 于: 十二月 15, 2021, 10:22:52 am »
假设程序中有两个核函数:kernel1和kernel2

        kernel2中想要用kernel1 的结果,可kernel2使用的时候值为空,即一个kernel运行的结果没有传递给另外一个kernel使用,那么多个kernel之间的数据应该怎么做才能传递使用??

         下边程序中开辟的内存也没有释放,为什么kernel1:initV 时对u0、u1成功赋值之后kernel2:convPy 使用时值却为Nan呢?
程序代码: [选择]

// 先在device开辟内存
cudaMalloc((void**)&u0, mea * wr * sizeof(float));
cudaMalloc((void**)&u1, mea * wr * sizeof(float));

// 通过initV对u0、u1 赋值
initV <<<grid1, block1 >>> (u0, u1, k);

// 同步,确保上边u0、u1成功初始化
cudaDeviceSynchronize();

// convPy函数需要用到u0、u1的值来计算temp
convPy <<<grid22, block22 >>> (temp, u0, u1);

// 最后才释放内存
cudaFree(u0);
cudaFree(u1);


10
程序代码: [选择]
// Main.cu
#define  _CRT_SECURE_NO_WARNINGS
#define nFrame 2000

int main()
{
//译码程序计时
// struct timeval startAll,endAll;
// gettimeofday(&startAll, NULL);

cudaEvent_t startAll, stopAll;
float elapsedTime = 0.0;
cudaEventCreate(&startAll);
cudaEventCreate(&stopAll);

cudaEventRecord(startAll, 0);

genSource(src,mCol); //信源产生
rcmEncode(P,V, src, encOut,i,nFrame); //编码

//计算程序运行时间
cudaEventRecord(stopAll, 0);
cudaEventSynchronize(startAll);    //这句有没有结果都是0
cudaEventSynchronize(stopAll);

cudaEventElapsedTime(&elapsedTime, startAll, stopAll);
cudaEventDestroy(startAll);
cudaEventDestroy(stopAll);
printf("程序运行时间time = %f(ms)\n", elapsedTime);
printf("Processing Finish!!");
return 0;
}
// rcmEncode的实现

//kernel:编码,实质就是测量矩阵与权矢量相乘
__global__ void encode1(int *d_P,int* d_S,int* d_C)
{
unsigned int i = threadIdx.x + blockIdx.x * blockDim.x;
d_C[i] = -4 * d_S[d_P[i * wr + 0]] - 4 * d_S[d_P[i * wr + 1]] - 2 * d_S[d_P[i * wr + 2]] - d_S[d_P[i * wr + 3]]
     + d_S[d_P[i * wr + 4]] + 2 * d_S[d_P[i * wr + 5]] + 4 * d_S[d_P[i * wr + 6]] + 4 * d_S[d_P[i * wr + 7]];
}

void rcmEncode(int* P, int* V, int* src, int* encOut,int fi,int frame) {

//分配设备全局内存
int *d_S, *d_C,*d_P;

cudaMalloc((int**)&d_S, mCol*sizeof(int));//src
cudaMalloc((int**)&d_C, mea*sizeof(int));  //测量值
cudaMalloc((int**)&d_P, mea * wr * sizeof(int));//P
cudaMemset(d_C, 0, mea * sizeof(int));//初始化

//复制数据:从host --> device global,默认同步方式
cudaMemcpy(d_S, src, frameLen * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(d_P, P, mea * wr * sizeof(int), cudaMemcpyHostToDevice);

//分配线程和线程块:开启M个线程并行计算M个测量值,其中每个线程负责计算其中的一个测量值。
dim3 block(512);  //一个block的最大线程数为1024
dim3 grid((mea+block.x-1)/block.x);

//调用核函数
encode1 <<<grid, block >>>(d_P, d_S,d_C); //异步执行
//同步,编码完成后再进行数据的复制
cudaDeviceSynchronize();

//复制数据:从device --> host
cudaMemcpy(encOut, d_C, mea* sizeof(int), cudaMemcpyDeviceToHost);

//释放设备的全局内存(常量区不需要释放)
cudaFree(d_S);
cudaFree(d_C);
cudaFree(d_P);
}


页: [1]