最新帖子

页: [1] 2 3 ... 10
1
CUDA / Re: 线程分化-以warp为分支单位
« 最后发表 作者 屠戮人神 六月 06, 2022, 04:54:32 pm »
感谢您的解答。
请问您说的异步载入是指这个吧?
程序代码: [选择]
__pipeline_memcpy_async(&shared[blockDim.x * i + threadIdx.x],
                            &global[blockDim.x * i + threadIdx.x], sizeof(T));
这个代码的效果是不是和我上述使用if~else效果相同,只是异步载入使用起来会更方便?

但是我又有一个疑问,如果使用异步载入的话,执行计算是一个核函数,进行异步载入又是一个核函数,可是在单流中,两个核函数只能串行进行,那就达不到计算和搬运数据同时进行的效果了。


首先说:
(1)和你用if...else, 在warp的边界上,不同的warp干不同的活(例如一个warp计算,一个warp载入)不同的。异步载入shared memory在载入命令发出后,warp可以继续往下做其他事情,而无需卡住。这样所有的warps都可以发出大量的进行中的load请求,也都可以计算一些不依赖的其他事情。和单独选出1个warp单独干载入(以前有篇文章专门讲这个),可能效果并不太一样。

其次:
(2)异步载入并不会产生一个新的kernel的,不会的。这点手册上已经说的很明白了,就不重复说了。你过度担忧了。
2
CUDA / Re: 线程分化-以warp为分支单位
« 最后发表 作者 jinyer 五月 25, 2022, 09:35:56 am »
感谢您的解答。
请问您说的异步载入是指这个吧?
程序代码: [选择]
__pipeline_memcpy_async(&shared[blockDim.x * i + threadIdx.x],
                            &global[blockDim.x * i + threadIdx.x], sizeof(T));
这个代码的效果是不是和我上述使用if~else效果相同,只是异步载入使用起来会更方便?
但是我又有一个疑问,如果使用异步载入的话,执行计算是一个核函数,进行异步载入又是一个核函数,可是在单流中,两个核函数只能串行进行,那就达不到计算和搬运数据同时进行的效果了。
3
CUDA / Re: 线程分化-以warp为分支单位
« 最后发表 作者 jinyer 五月 25, 2022, 09:09:00 am »
感谢您的解答。
请问您说的异步载入是指这个吧?
程序代码: [选择]
__pipeline_memcpy_async(&shared[blockDim.x * i + threadIdx.x],
                            &global[blockDim.x * i + threadIdx.x], sizeof(T));
这个代码的效果是不是和我上述使用if~else效果相同,只是异步载入使用起来会更方便?
4
CUDA / Re: cuda C和fortran 混编的实例代码
« 最后发表 作者 流云 五月 23, 2022, 03:54:55 pm »
学习一下 :)
5
CUDA / Re: 访问texture和global memory的区别
« 最后发表 作者 屠戮人神 五月 19, 2022, 07:16:40 pm »
测试程序kernel是参考官方例程写的为:
程序代码: [选择]
__global__ void LtsTBytesEquivL1sectormissPipeTexMemTextureKernel0(float *output, cudaTextureObject_t tex_obj,
                                                                   int width, int height, float theta) {
    // Calculate normalized texture coordinates
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

    float u = x / static_cast<float>(width);
    float v = y / static_cast<float>(height);


    // Read from texture and write to global memory
    output[y * width + x] = tex2D<float>(tex_obj, u, v);
}

void LtsTBytesEquivL1sectormissPipeTexMemTexture() {
    const int test_num = 1;
    const int loop_num = 1;
    int grid_num[test_num]{32};
    int thread_num[test_num]{32};
    float *h_src[test_num] = {nullptr};
    float *d_res[test_num] = {nullptr};

    float angle = 0.5;
    cudaChannelFormatDesc channel_desc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
    cudaArray_t cu_array[test_num];

    std::srand(time(nullptr));  // use current time as seed for random generator
    CUDA_CHECK(cudaSetDevice(0));
    for (int i = 0; i < loop_num; i++) {
        dim3 block(thread_num[i]);
        dim3 grid(grid_num[i]);

        // alloc cpu memory
        h_src[i] = new float[thread_num[i] * grid_num[i]];
        for (int j = 0; j < grid_num[i] * thread_num[i]; j++) {
            h_src[i][j] = (std::rand() % 255);
        }

        // Allocate CUDA array in device memory
        CUDA_CHECK(cudaMallocArray(&cu_array[i], &channel_desc, thread_num[i], grid_num[i]));

        // Set pitch of the source (the width in memory in bytes of the 2D array pointed
        // to by src, including padding), we dont have any padding
        const size_t spitch = thread_num[i] * sizeof(float);

        // Copy data located at address h_data in host memory to device memory
        CUDA_CHECK(cudaMemcpy2DToArray(cu_array[i], 0, 0, h_src[i], spitch, thread_num[i] * sizeof(float), grid_num[i],
                                       cudaMemcpyHostToDevice));

        // Specify texture
        struct cudaResourceDesc res_desc;
        memset(&res_desc, 0, sizeof(res_desc));
        res_desc.resType = cudaResourceTypeArray;
        res_desc.res.array.array = cu_array[i];

        // Specify texture object parameters
        struct cudaTextureDesc tex_desc;
        memset(&tex_desc, 0, sizeof(tex_desc));
        tex_desc.addressMode[0] = cudaAddressModeWrap;
        tex_desc.addressMode[1] = cudaAddressModeWrap;
        tex_desc.filterMode = cudaFilterModeLinear;
        tex_desc.readMode = cudaReadModeElementType;
        tex_desc.normalizedCoords = 1;

        // Create texture object
        cudaTextureObject_t tex_obj = 0;
        cudaCreateTextureObject(&tex_obj, &res_desc, &tex_desc, NULL);

        // Allocate result of transformation in device memory
        CUDA_CHECK(cudaMalloc(&d_res[i], grid_num[i] * thread_num[i] * sizeof(float)));

        // kernel
        switch (i) {
            case 0: {
                LtsTBytesEquivL1sectormissPipeTexMemTextureKernel0<<<grid, block>>>(d_res[i], tex_obj, thread_num[i],
                                                                                    grid_num[i], angle);
                break;
            }
            default:
                break;
        }

        CUDA_CHECK(cudaStreamSynchronize(0));
        CUDA_CHECK(cudaGetLastError());

        // Destroy texture object
        cudaDestroyTextureObject(tex_obj);
        cudaFreeArray(cu_array[i]);

        delete[] h_src[i];
        CUDA_FREE(d_res[i]);
    }
}

程序实现的功能为:从texture中读取32*32的数据写入到global memory中,然后测试了lts__t_bytes_equiv_l1sectormiss_pipe_tex、lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_texture、lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_texture_op_ld和lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_texture_op_tex这四个metrics,测试结果为:
程序代码: [选择]
  LtsTBytesEquivL1sectormissPipeTexMemTextureKernel0(float *, unsigned long long, int, int, float), 2022-Apr-15 10:38:35, Context 1, Stream 7
    Section: Command line profiler metrics
    ---------------------------------------------------------------------- --------------- ------------------------------
    lts__t_bytes_equiv_l1sectormiss_pipe_tex.avg                                      byte                            512
    lts__t_bytes_equiv_l1sectormiss_pipe_tex.max                                      byte                            nan
    lts__t_bytes_equiv_l1sectormiss_pipe_tex.min                                      byte                            nan
    lts__t_bytes_equiv_l1sectormiss_pipe_tex.sum                                     Kbyte                          16.38
    lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_texture.avg                          byte                            512
    lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_texture.max                          byte                            nan
    lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_texture.min                          byte                            nan
    lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_texture.sum                         Kbyte                          16.38
    lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_texture_op_ld.avg                    byte                              0
    lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_texture_op_ld.max                    byte                            nan
    lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_texture_op_ld.min                    byte                            nan
    lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_texture_op_ld.sum                    byte                              0
    lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_texture_op_tex.avg                   byte                            512
    lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_texture_op_tex.max                   byte                            nan
    lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_texture_op_tex.min                   byte                            nan
    lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_texture_op_tex.sum                  Kbyte                          16.38
    ---------------------------------------------------------------------- --------------- ------------------------------

以lts__t_bytes_equiv_l1sectormiss_pipe_tex_mem_texture_op_tex为例,该metrics的解释为:
# of bytes requested for TEX instructions
使用sass代码的确是有这个指令,但是我比较疑惑的是为什么数据量是16384byte,
因为我的数据总量为32*32*4=4096byte,请问这是为什么呢?

可能存在插值时候的读取放大效应,例如你可能在2D的归一化坐标的线性滤波的时候,某个坐标需要读取临近的4个点;1D的时候,可能需要读取临近的2个点(或者精确的说,纹元,texel,texture elment)。

放大的读取也可能被CUDA Array的不透明的数据重排列方式+cache的读取而导致的。关于这点,2013年的<CUDA Handbook>里面解释了这个不透明的数据安排方式(2D下4个点一组)。可能这也会贡献了读取放大的一部分可能因素。

欢迎继续回帖,我回复的太晚了。
6
CUDA / Re: SM、block、warp
« 最后发表 作者 屠戮人神 五月 19, 2022, 07:11:48 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的四分之一,这又是怎么做到的呢】

请问一个SM是同时在执行16个block吗
--这个得看你的”同时“的定义,如果是常见的类似CPU风格的”超线程“方式的理解的话,那么这多个blocks同时驻留在1个SM上(手册有关于”驻留“的概念的解释),随时可能它们中间的warps被执行,则可以认为是这么多的blocks/warps同时在执行。
--如果从SM内部结构的角度看,例如从手册的计算能力章节的scheduler的情况来说,那么最多可能同时只有0-4个warps在真正同时被执行(任意周期,和warps的就绪情况,以及计算能力有关)。详情可以看手册。
7
CUDA / Re: 线程分化-以warp为分支单位
« 最后发表 作者 屠戮人神 五月 19, 2022, 07:08:25 pm »
对开发者而言,应该尽量避免在同一个Warp中有不同的执行路径。当必须写分支语句时,尽量让分支宽度大于WarpSize,即不以Thread为分支单位,而是以Warp为分支单位,这样就能保证同一个Warp内部不会出现分支。
请问如果分配线程块 block(256)
核函数中是
程序代码: [选择]
int warpIdx = threadIdx.x/32;
if(warpIdx < 2){
    执行语句1;
}else{
    执行语句2;
}
代码是以warp为分支单位,那这种情况下if和else可以同时运行吗?
即现在有一个需求,想让线程块中一部分线程来执行计算,另一部分线程来把数据从全局内存搬运到共享内存,请问我可以通过上述if-else达到这两种需求同时运行的效果吗?

关于你的几个问题:
(1)“当必须写分支语句时,尽量让分支宽度大于WarpSize”---我猜测你的正确意思应该是分支宽度是warpSize的倍数。
(2)“那这种情况下if和else可以同时运行吗”--如果编译器生成了常规的跳转指令,同时2个warp执行到了不同分支的位置,则在SM有多个scheduler的情况下,的确2个warp可能在同时执行这2个不同的代码路径。
(3)“即现在有一个需求,想让线程块中一部分线程来执行计算,另一部分线程来把数据从全局内存搬运到共享内存,请问我可以通过上述if-else达到这两种需求同时运行的效果吗?”---这个的确可以,前人有一篇老文章,讲述了通过"warp specilization来执行特殊的载入任务,一个warp为其他warp服务“(大致意思,原文记不清了,你可以搜搜对应的英文原文”。

不过考虑到现在已经有了异步载入指令(例如你的8.6计算能力的RTX30卡),和最新9.0计算能力的内嵌在SM内部的微型DMA引擎,都可以完成你设想的载入global->shared的任务,你的独立使用1个warp来进行的主意,可能没有太大的用途。我建议你使用异步载入,手册有介绍(针对8.6的,9.X的你还得等)。
8
CUDA / Re: 线程分化-以warp为分支单位
« 最后发表 作者 jinyer 五月 09, 2022, 11:18:22 am »
或者怎么看作用于if的和作用于else的不同warp是否同时执行了呢
9
CUDA / 线程分化-以warp为分支单位
« 最后发表 作者 jinyer 五月 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达到这两种需求同时运行的效果吗?
10
CUDA / Re: Nsight Compute没有权限问题(linux系统上)
« 最后发表 作者 jinyer 四月 18, 2022, 10:04:58 am »
直接输入 ncu-ui 可以打开compute,就是launch时核函数会报没有权限的错误;
sudo ncu-ui 会出现说没有qt平台导致不能start the app【而上边直接ncu-ui软件是可以使用的,这是为什么呢】
命令执行如下:
dell@dell-Precision-5820-Tower-X-Series:~$ ncu-ui
********BP-RCM----nFrame:1, FrameLen=4096, mea=8192, iteration=1  [cudaEvent]***********
dell@dell-Precision-5820-Tower-X-Series:~$ sudo ncu-ui
No protocol specified
qt.qpa.xcb: could not connect to display :10.0
qt.qpa.plugin: Could not load the Qt platform plugin "xcb" in "" even though it was found.
This application failed to start because no Qt platform plugin could be initialized. Reinstalling the application may fix this problem.

Available platform plugins are: xcb.

/usr/local/cuda-11.2/bin/../nsight-compute-2020.3.0/host/linux-desktop-glibc_2_11_3-x64/ncu-ui:行 16: 3025862 已放弃               (核心已转储) "$NV_AGORA_PATH/CrashReporter" "NVIDIA Nsight Compute" "NVIDIA Nsight Compute" "2020.3.0.0 (build 29307467) (public-release)" "$NV_AGORA_PATH/ncu-ui.bin" "$@"
已解决,感谢回答
页: [1] 2 3 ... 10