列出帖子

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


显示所有帖子 - 屠戮人神

页: [1] 2 3 ... 28
1
想把分配的local memory全部设置为某个值,需要一个一个单元赋值吗?有没有类似memset的函数?

PS. GPU的local memory缺省全是0吗?还是随机数?

谢谢!

(1)最快的方式是work-group集体清零。这也是推荐的方式。暂时没有专用于清零local memory的特定函数(as of OpenCL 1.2)。
(2)默认的值不是0的,如果你需要0值,你得显式的手工清零。
(3)OpenCL有async_work_group_copy函数,可以从global载入一堆0值给local memory, 但该函数除了编码时候方便一点点,并没有性能上的优势,不建议这样做。


2
CUDA / Re: 请问这个是内存分配的错误吗?
« 于: 今天 01:36:07 pm »
楼主你好,并不存在不能使用cudaDeviceSynchronize(), "一使用kernel就挂了“的情况。

同步的时候kernel挂掉,只能证明是你的kernel本身存在问题,而不是同步让它挂掉。

你的代码至少存在如下问题:
(1)wrapper函数传入的是host内存地址:
程序代码: [选择]
cudaError_t convertBinaryToDecimal(int* a,  unsigned int size)
(2)wrapper内部正确的分配了device显存缓冲区:
程序代码: [选择]
cudaStatus = cudaMalloc((void**)& dev_a, size * sizeof(int));
(3)以上两步都没问题,然而你的kernel并没有使用该分配出来的缓冲区(显存),而是试图直接使用内存:
程序代码: [选择]
convertBinaryToDecimalKernel <<<1, 1 >>> (a);

这里才是导致你的kernel挂掉的根源。楼主你未能理解为何kernel要使用显存,为何CPU使用内存的根本原理,只是照抄了大致的骨架(形式看上去是对的),而未能明白kernel给他传递的地址到底应当是什么。

建议的解决方案:
(1)将a改成dev_a。或者:
(2)阅读本论坛的Sisiy妹子的《阅读CUDA 100天》系列文章,学会使用其中的managed memory (unified memory), 后者可以全自动的规避你的这种问题。

3
CUDA / Re: 请问这个是内存分配的错误吗?
« 于: 昨天 02:38:43 pm »
Hello.
"an illegal memory access was encountered" error means that: Your kernel has failed due to invalid memory access patterns.

Please check the followings conditions according to your screen shots and partial code lines:

(1)Whether kernel argument "a" is a valid pointer. As error code 700 suggests, the address (the value of your pointer a) is invalid.

(2)if "a" is invalid, for example, it's NULL. then check your host code, and find why the allocations on "a" failed.

(3)Please note you have provided only a portion of code. We cannot analyze your host code for you in this situation.

mao@GPUWorld

As you have two kernels, which one really failed remains unknown to us at the moment. You have to check both of them. (Your pictures are mis-leading. In-consistent kernel names in your error logs)

The second kernel accessed only one element of buffer a. that's fine. just check the allocation on "a" itself.

The first kernel however accessed several elements of three buffers. You do have to check all three of them: their allocations, and their capacities. If you used more elements from these buffers, going beyong their capacities. Kernel would fail too.

You have to check all the conditions above.

4
CUDA / Re: 请问这个是内存分配的错误吗?
« 于: 昨天 02:28:55 pm »
请问这个是内存分配的错误吗?我是按照例子程序copy的,为什么会分配失败。。我的核函数是把2进制数转为10进制数

Hello.
"an illegal memory access was encountered" error means that: Your kernel has failed due to invalid memory access patterns.

Please check the followings conditions according to your screen shots and partial code lines:

(1)Whether kernel argument "a" is a valid pointer. As error code 700 suggests, the address (the value of your pointer a) is invalid.

(2)if "a" is invalid, for example, it's NULL. then check your host code, and find why the allocations on "a" failed.

(3)Please note you have provided only a portion of code. We cannot analyze your host code for you in this situation.

mao@GPUWorld

5
CUDA / Re: 這是怎麽回事
« 于: 昨天 02:20:18 pm »
如題

Hi. Lady. Your CUDA code lines returned 0x0b which means invalid value(s).

Solutions:
Examine the parameters you passed to API call cudaMemset(), and fix one or more invalid values you found. Then re-build your project and run it again.

6
CUDA / Re: CUDA版本和驱动的更新会带来负优化吗?
« 于: 十月 12, 2019, 01:55:47 pm »
老项目使用的sm52配CUDA7.5+驱动354,由于系统更新需要换更高版本的CUDA,更新了CUDA10.1+驱动432,结果发现运算效率和IO性能都有明显的下降。咨询以后得到的信息是由于CUDA版本的提升运算模型有发生变化,因此导致性能的下降。但是鄙人没能找到相关的资料文档,如果有更详细一点的说明就好了! :P

您好,我们也一直在关注此问题。但根据目前的情况,我们无法给出负面的或者正面的回答。请保持您自己的经验不要声张为好。

7
OpenCL / Re: 怎样查看GPU能支持多少个CommandQueue同时运行?
« 于: 九月 18, 2019, 05:06:34 pm »
非常感谢您的答复。其实我主要想了解GPU的最小工作单位是什么?各个work-item是否能独立工作,还是不管任务多小,一次都会占据一整个计算单元,甚至整个GPU?

我想投递多个queue的原因是担心kernel在某些参数下运行时间过长,导致同一批任务被其中一个拖后腿(比如平均一个任务3毫秒,千分之一的可能性需要2-3秒,跑百万个批次的任务就会造成极大延时),所以想分开跑,提高运行通量。可能想法幼稚了:P

这些都属于基本问题:
(1)请看一下N卡或者A卡的基本架构。手册均在网上公开可用。你得学会自学基本的常识。注意N卡的架构是用CUDA的术语来说的(例如线程,例如块,而不是你常见的work-item或者group之类的)。

(2)你可以认为一个kernel的启动规模较大(分别是CUDA或者OCL的grid或者ndrange)。但其中的最小调度单位(执行单位),往往是warp或者wavefront, 请自行搜索这两个是什么。而资源占用的单位往往是block或者group(也请自行搜索)。当某个kernel中有1个warp/wave迟迟不能结束的时候(这叫long-tail,也就是长尾/拖尾效应),最多只会有1个block或者group的资源被占用。而关于GPU如何具体blocks/groups上去执行的(这往往叫驻留,resident), 也请自行搜索。

(3)可以将多个kernel并发执行作为解决长尾效应的一种方式,这没错。但不能作为一种常规的并行手段(你并行的是work-items,而尽量不要是kernels)。

8
OpenCL / Re: 怎样查看GPU能支持多少个CommandQueue同时运行?
« 于: 九月 18, 2019, 12:53:45 pm »
新手求助解惑
RT, 假设1个GPU有16个计算单元,每个单元有32个核。

设计一个kernel只占用1个work item(1个核?)运行死循环,那么我开512(16*32)个CommandQueue,每个CommandQueue都提交这个kernel,最多能有多少个kernel同时运行呢?不同GPU是不是不一样?有没有命令可以查询这个数值?

谢谢!

具体一个GPU能同时在卡上执行的kernel数量,和具体的硬件有关的。

但你的假设的应用场景不存在,与其1次kernel启动只有1个work-item,启动512次;不如直接考虑启动1次512个work-items的规模,你说呢。

关于你的第三个问题,目前没有直接能查询到这个数值的命令或者说设备属性。但常见GPU有32个,128个(这两个是N家的),以及,64个(A家的,GCN 1.1/2). AMD稍微有点文档可以看ACE的数量,以及每个ACE支持的队列数量之类的信息。(但如同上一段说的,这没必要,你应当考虑一次性启动更大规模,而不是考虑卡能否支持多次启动的kernel在并发)

9
CUDA / Re: 多线程调用GPU问题
« 于: 九月 11, 2019, 09:32:01 pm »
下图里4个host线程,每个线程里创建一个stream,期望是每个线程的内存拷贝顺序执行,和kernel可以overlap。
但从nvvp图上看,各线程的memcpy和launchKernel相互影响,虽然stream1的拷贝完成了,但kernel并没能跟着执行。
我猜测是cuda的调用内部有锁机制,一个context里的cuda调用还是顺序执行的,这样的话多个host线程的意义不大。
[/size]测试总体时间基本是4倍的单线程时间。

你这缺少太多的信息了,例如:
(1)是否使用Windows?如果是,是否是专业卡驱动(TCC驱动?),普通WDDM驱动将限制可能的一次或者多次overlap

(2)传输是否使用了page-locked memory? 只有page-locked memory才会可能会kernel交叠。

等等。请补充。

10
CUDA / Re: 多线程调用GPU问题
« 于: 九月 10, 2019, 10:21:43 pm »
目的是将GPU上的实现的功能封装成库对外调用,用多线程时发现时间反而变长,更像是顺序执行的,用nvvp查看,多线程调用的cuda函数耗时比较长。感觉一个context上的调用都是顺序执行。
先请教前辈这方面的经验,是不是不支持多线程调用。

没有必要一个host线程对应一个CUDA Context的,你完全可以多个Host线程共享1个CUDA Context, 然后用Stream隔离开前后依赖关系:

(1)这样多线程调用的时候并不会引入context切换的开销。
(2)是否特定的卡可以多个context并发的执行,目前无资料。

注意多个host线程共享1个CUDA Context,是目前使用Runtime API情况时候的默认处理状态。

11
CUDA / Re: CUDA kernel里面调用kernel
« 于: 八月 27, 2019, 01:23:40 pm »
CUDA里面有没有opencl 2.0里面的那种 设备端入队的操作,可以kernel里面调用kernel。

请直接在手册或者本论坛搜索CUDA Dynamic Parallism (CUDA动态并行)。

12
阅读CUDA英文手册100天 / Re: DAY37:
« 于: 八月 27, 2019, 01:22:14 pm »
关于是否仅限同一个文件内可见,这点还是不对的。可以跨文件使用的。请参考手册中的分步编译章节,以及,搜索-rdc=true选项。

跨文件使用__device__有时候还是重要的。

13
阅读CUDA英文手册100天 / Re: DAY37:
« 于: 八月 26, 2019, 01:39:44 pm »
目前尚不能支持:

__device__ type var[你需要的动态大小]格式的。

也就是说,目前只能是你说的静态的。

好处主要有两点:
(1)减少了kernel的参数传递,没有它,你必须传递一次额外的参数。
(2)很多形如__device__ float bruce[H][W];的形式,推导实际地址的时候(即: bruce首地址 + 4 * (y * W + x))的时候,可以让编译器方便的考虑编译时刻的优化(因为H,W都是编译时刻已知的)。不过第二点也可以通过形如float (*p)[H][W]的参数进行(其中H可以省略,即参数float (*p)[][W])。

关于你说的命名污染、对程序内聚力的破坏,和不符合软件工程学的相关问题,这些我不懂,暂时无法回答。

14
阅读CUDA英文手册100天 / Re: DAY38:阅读春存储器修饰符
« 于: 八月 26, 2019, 01:33:23 pm »
这里的指针推导/指针变换指的是什么?可不可以举例说明要避免的做法?

不建议:
type *p1 = f1(p0);
type *p2 = f2(p1);
type *p3 = f3(p2);
....

建议:
直接使用p0[offset1 + ....], p0[offset2 + ...], p0[offset3 + ...]


15
CUDA / Re: Volta架构中规约计算的同步问题
« 于: 八月 23, 2019, 03:18:33 pm »
引用
for (int offset = blockDim.x / 2; offset > 0; offset /= 2)
{
    if (tid < offset)
    {
        shmem[tid] += shmem[tid + offset]; // 有问题吗?为什么?
    }
    __syncthreads();
}

同时让我们来看看你的第三个代码,这代码其实和之前的并不同,这里多了这么一行:
if (tid < offset)   <---- 注意这里的约束条件
{[tid] = [tid] + [tid + offset];}
__syncthreads();

此代码在多了这个约束条件, 在此条件下,对于同1个block内部的线程来说(假设block是1D的),
不可能在该约束条件下,存在任意2个线程,使得读取的[tid]或者[tid + offset]和另外一个线程的写入的[tid']元素相同。

这个容易证明:
如果相同,
要么(1)下标tid = tid' (tid来自于threadIdx.x, 刚才强调了是1D的,因此不可能),这是不可能的。
要么(2)下标tid + offset = tid', 即存在2个线程,他们的下标差为offset。但是刚才约束了所有的下表范围均为[0, offset), 在此范围内是找不到任意两个坐标,能坐标差为offset的(显然!)

结合(1)(2),这种是安全的。

所以你看,我们分析的时候只需要线程间的读写。而无视了同一个线程对同一个下标位置的读取,然后写入(上面说了,同一个线程这样做是安全的)。

你以后可以用这种方式分析任意类似代码。

页: [1] 2 3 ... 28