列出帖子

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


只显示主题 - jinyer

页: [1]
1
CUDA / 多个线程进行判断修改同一个全局内存遇到的问题
« 于: 三月 30, 2023, 11:21:58 am »
程序代码: [选择]
int *d_flag;
checkCudaErrors(cudaMalloc((void**)&d_flag, 1*sizeof(int)));
cudaMemset(d_flag, 1, sizeof(int)*1);  //初始化为1
block = 96;
isIter<<<(M+block-1)/block,block>>>(d_flag, d_check);
isIter核函数想实现的功能是 ：flag[0]初始化是1，若check数组有一个数不为0，则将flag[0]置为0
程序代码: [选择]
__global__ void isIter(int *flag, int *checkResult) {
int idx = threadIdx.x + blockDim.x * blockIdx.x;
if(idx < M){
if(checkResult[idx] != 0){
flag[0] = 0;
}
}
}
可这样老是会输出flag[0]=16843009（也有部分线程进入if语句里边，输出flag[0]=0），可将flag[0]复制到CPU读取的时候是16843009，请问这是为什么啊？
flag存储在全局内存，不应该若有线程修改的话值就变为0，没有线程修改还是保持原值1吗？

2
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记录的核函数执行时间】

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

4
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达到这两种需求同时运行的效果吗？

5
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
，不知道要怎么做，请教大家

6
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的四分之一，这又是怎么做到的呢】

7
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();吗？

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

9
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没有这个选项，请教一下要怎么改正？

10
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);

​

11
CUDA / 使用cudaEvent计时，也使用了cudaEventSynchronize(stopAll);计时结果为0是为什么呢
« 于: 十二月 14, 2021, 11:43:55 am »
程序代码: [选择]
// 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]