最新帖子

页: [1] 2 3 ... 10
1
CUDA / 如何提高全局内存的写入速率?
« 最后发表 作者 皮皮虾and皮皮猪 八月 16, 2022, 04:25:52 pm »
代码如下所示,d_m_rgbbuffer指向的是全局内存,该函数的作用是YUV转RGB 注释部分(*lmbmp++部分)就是RGB数据写入全局内存中,这一部分所花时间占整个时间的70%,(一个线程处理192个数据)请问应该如何进行优化?
程序代码: [选择]
__device__ void StoreBuffer(short * QtZzMCUBuffer, short dx, short dy, unsigned char* d_m_rgbbuffer)
{
short i, j;
unsigned char* lpbmp;
unsigned char R, G, B;
int y, u, v, rr, gg, bb;
//unsigned char RGB[63 * 3];//得到该8*8像素的RGB
for (i = 0; i <8; i++)
{
if ((dy + i) < d_m_height)
{
lpbmp = ((unsigned char*)d_m_rgbbuffer + (dy + i) * 3 * d_m_width + dx * 3);
for (j = 0; j <  8; j++)
{
if ((dx + j) < d_m_width)
{
y = QtZzMCUBuffer[i * 8 + j];
u = QtZzMCUBuffer[i * 8 + j + 64];
v = QtZzMCUBuffer[i * 8 + j + 128];
rr = ((y << 8) + 18 * u + 367 * v) >> 8;
gg = ((y << 8) - 159 * u - 220 * v) >> 8;
bb = ((y << 8) + 411 * u - 29 * v) >> 8;
R = (unsigned char)rr;
G = (unsigned char)gg;
B = (unsigned char)bb;
if (rr & 0xffffff00) if (rr > 255) R = 255; else if (rr < 0) R = 0;
if (gg & 0xffffff00) if (gg > 255) G = 255; else if (gg < 0) G = 0;
if (bb & 0xffffff00) if (bb > 255) B = 255; else if (bb < 0) B = 0;
//*lpbmp++ = B;
//*lpbmp++ = G;
//*lpbmp++ = R;
}
else  break;
}
}
else break;
}
}
2
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的,不会的。这点手册上已经说的很明白了,就不重复说了。你过度担忧了。
3
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效果相同,只是异步载入使用起来会更方便?
但是我又有一个疑问,如果使用异步载入的话,执行计算是一个核函数,进行异步载入又是一个核函数,可是在单流中,两个核函数只能串行进行,那就达不到计算和搬运数据同时进行的效果了。
4
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效果相同,只是异步载入使用起来会更方便?
5
CUDA / Re: cuda C和fortran 混编的实例代码
« 最后发表 作者 流云 五月 23, 2022, 03:54:55 pm »
学习一下 :)
6
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个点一组)。可能这也会贡献了读取放大的一部分可能因素。

欢迎继续回帖,我回复的太晚了。
7
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的就绪情况,以及计算能力有关)。详情可以看手册。
8
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的你还得等)。
9
CUDA / Re: 线程分化-以warp为分支单位
« 最后发表 作者 jinyer 五月 09, 2022, 11:18:22 am »
或者怎么看作用于if的和作用于else的不同warp是否同时执行了呢
10
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达到这两种需求同时运行的效果吗?
页: [1] 2 3 ... 10