能否在kernel里调用其他的kernel？

CUDA > CUDA编程
收藏 回复
[问题点数：20分]
更多帖子
私信 空间 博客
dushuchen
等级
本版专家分：0
结帖率 75%
有点复杂
不知道现在的CUDA是否支持
只看楼主 引用 举报 楼主 收起
展开阅读全文
私信 空间 博客
linxxx3
等级
本版专家分：2962
CUDA最新的版本可以支持，但是关键是显卡设备需要3.0以上计算能力，一些老显卡可以排除了。而且这种特性对性能和编程复杂度上，我没看出有明显好处。
只看TA 引用 举报 #1得分 0
更多帖子
私信 空间 博客
dushuchen
等级
本版专家分：0
引用 1 楼 linxxx3 的回复:
CUDA最新的版本可以支持，但是关键是显卡设备需要3.0以上计算能力，一些老显卡可以排除了。而且这种特性对性能和编程复杂度上，我没看出有明显好处。

谢谢

为什么没有明显好处呢？具体在哪些方面？
只看TA 引用 举报 #2得分 0
私信 空间 博客
yjoe61
等级
本版专家分：75
需要3.5或以上的计算能力。

好处还是有的，比如你需要host kernel的每一个thread再调用n个block和thread。这时你不需要把数据传回host再运行kernel，而是可以直接在kernel里运行新的kernel。
只看TA 引用 举报 #3得分 0
私信 空间 博客
linxxx3
等级
本版专家分：2962
引用 3 楼 yjoe61 的回复:
需要3.5或以上的计算能力。

好处还是有的，比如你需要host kernel的每一个thread再调用n个block和thread。这时你不需要把数据传回host再运行kernel，而是可以直接在kernel里运行新的kernel。


即使不用这个特性，也不需要将数据传回主机，再启动新kernel，因为数据本身就在显存里，全部算完了，最后传一次就行。
启动时间可能减少，是一个好处。但是要注意的是：1. 如果启动时间是一个问题，说明写的kernel粒度太细，需要调整；2. 这种新特性其实容易让人产生疑惑，编程学习的成本不是减少了，而是增加。

用CUDA强调的就是计算效率，想避开底层架构不需要了解不现实。这个特性出来，可以想象的是，主要的作用是方便已有的CPU程序移植到GPU上，对原程序的结构修改小，很少的时间就能得到一定效果。但是真的想要得到比较高的加速比，简单改改一般靠不住，这个特性就鸡肋了。
只看TA 引用 举报 #4得分 0
私信 空间 博客
yjoe61
等级
本版专家分：75
一个简单的例子是quicksort。如果不在kernel里调用kernel，那么必须在host做一次synchronization，将上一层的运算结果（比如排序的pivot序号）传回host，然后再在host执行下一层的kernel。如果多个kernel同时运算，这将大大增加host控制的复杂性。总的来说，任何分而治之的算法都有这个问题。

如果能在kernel直接执行kernel，并且同步的话，复杂度就大大降低了。
只看TA 引用 举报 #5得分 0
私信 空间 博客
linxxx3
等级
本版专家分：2962
引用 5 楼 yjoe61 的回复:
一个简单的例子是quicksort。如果不在kernel里调用kernel，那么必须在host做一次synchronization，将上一层的运算结果（比如排序的pivot序号）传回host，然后再在host执行下一层的kernel。如果多个kernel同时运算，这将大大增加host控制的复杂性。总的来说，任何分而治之的算法都有这个问题。

如果能在kernel直接执……


递归调用的例子确实是我没考虑过的，学习了。不过还是有些地方要探讨的，我也不百分比确定：
1. host上要使多个kernel同时启动，好像一定必须使用流(stream)？从kernel内启动kernel隐含的保证一个执行顺序，即调用的子函数必须完成之后，母函数才会继续，隐含了同步操作，这样确实能使编程容易。但是这样做，多个kernel能并行运算吗？能避免使用流吗？

2. 如果有多个kernel同时运算，在host执行使用sync会造成什么结果？实际上我们需要的只是保证执行的先后顺序，不需要全局同步来控制，而应该使用阻塞的调用，比如使用事件（event）。对快排这个例子，在partition之后阻塞，直到partition完成之后，再调用子序列的quicksort就可以了。调用的两个子序列quicksort为了并行运算，可以启动两个新的流，不过可能遇到递归调用流的数目不够用。

3. 新特性多比较麻烦，一知半解危险，要用的话，至少又要读一篇新东西 "CUDA
Dynamic Parallelism programming guide"，我还没找到有中文的资料能把这点讲清楚的。看了也不保证用的好。打比方就是本来难走的路，有很多陷阱，它拿些树叶啥的给盖上，看起来是好走，等掉坑里的时候就想哭了。
只看TA 引用 举报 #6得分 0
私信 空间 博客
yjoe61
等级
本版专家分：75
我想了一下，貌似用device端的kernel还是可以使编程简单不少。如有不妥欢迎指教。

关于stream我是这样理解的：device启动kernel，也可以指定stream。如果不指定stream，就默认使用NULL stream。这个NULL stream是整个block共享的：也就是说这个block启动的所有子kernel都会进入这个stream。但是子kernel和母kernel还是可以异步运行的。

假设现在在host启动第一级的qsort，这样得到两个子序列，传回host。然后就需要在host再启动两个新的不相干的kernel。其中一个完成后又传回两个子序列。而这两个子序列和另一个kernel传回的子序列又是不相干的。。。这样貌似还是有很大的复杂性，差不多等同于需要和kernel同样多数量的host thread来控制每一级的kernel启动。而这些工作，如果有device级的kernel，就可以直接在kernel里完成。从易用性的角度来说，我感觉这样对程序员更具亲和性。

不知道那个Guide有没有人打算翻译一下？
只看TA 引用 举报 #7得分 0
私信 空间 博客
linxxx3
等级
本版专家分：2962
引用 7 楼 yjoe61 的回复:
我想了一下，貌似用device端的kernel还是可以使编程简单不少。如有不妥欢迎指教。

关于stream我是这样理解的：device启动kernel，也可以指定stream。如果不指定stream，就默认使用NULL stream。这个NULL stream是整个block共享的：也就是说这个block启动的所有子kernel都会进入这个stream。但是子ker……


从我看guide的结论，母kernel上启动子kernel，如果用默认stream，等子kernel执行完成，母kernel才继续执行，给你贴个图，如果你手头有这个文档，是在第4页：


不管怎么说，CUDA学习曲线挺陡的，多学一个特性其实也没啥，如果能用好的话就赚了，不熟悉的话，不用也能写程序。
只看TA 引用 举报 #8得分 0
更多帖子
私信 空间 博客
rainliuyu
等级
本版专家分：20
引用 8 楼 linxxx3 的回复:
引用 7 楼 yjoe61 的回复:
我想了一下，貌似用device端的kernel还是可以使编程简单不少。如有不妥欢迎指教。

关于stream我是这样理解的：device启动kernel，也可以指定stream。如果不指定stream，就默认使用NULL stream。这个NULL stream是整个block共享的：也就是说这个block启动的所有子kernel都会进入这个stream……



刚好碰到这个问题，想上来请教就看见这贴，好巧

我目前需要解决的是把一个串行算法调整为GPU算法
算法中有3层循环嵌套，每一层循环量都很大，而只在最里一层循环得出运算数据结果

设想的就是把每一层循环都做成一个device kernal，然后上层调用下层
如以上诸层所说，这样做是可行的？
只看TA 引用 举报 #9得分 0
收起
收藏 回复
dushuchen

等级：

关注 私信 TA的帖子
浅谈在linux kernel中打印函数调用的堆栈的方法

下面小编就为大家带来一篇浅谈在linux kernel中打印函数调用的堆栈的方法。小编觉得挺不错的。现在就分享给大家。也给大家做个参考。一起跟随小编过来看看吧

Linux Kernel系统调用分析

学号118 原创作品转载请注明出处 + https://github.com/mengning/linuxkernel/ 实验过程 内核编译 下载Linux Kernel 5.0.1 source code...然后在menuconfig中选择编译内核时添加调试信息 然后执行 qemu-system-i38...

CUDA C中核函数（kernel）的编写及调用

在此，我们以HelloWorld程序为例，分别讲述kernel函数的编写和调用。1、kernel函数的编写：__global__ void kernel_hello(void) { printf("Hello Wrld from GPU ! \n"); }修饰符__global__表示该函数从...

在linux kernel中打印函数调用的堆栈的方法

在linux kernel中打印函数调用的堆栈的方法

在kernel space调用系统调用的例子

一般来说系统调用都是需要在user space调用的，但是其实在kernel space也可以调用原本应该在user space调用的系统调用. 例如在init/do_mounts.c中就会调用sys_open void __init change_floppy(char *fmt, ...)

kernel 系统调用----system call

Init在trap_init中对SYSCALL_VECTOR（编号0x80）的向量进行初始化。 808 set_system_trap_gate(SYSCALL_VECTOR, &system_call);将system call初始化为trap门，加入到IDT table中，发生中断以后，会跳转到对应system_...

OpenCL中kernel的循环调用

kernel的循环调用主要是涉及缓冲区的创建和主机端命令同步

如何给Linux kernel 5添加一个系统调用

如何给Linux Kernel 5添加一个系统调用 本篇博客从英文博客 Adding a Hello World System Call to Linux Kernel 全部翻译而来。原文链接为： ...

系统调用从User进入Kernel

#include"stdio.h" int main() { int pid; asm volatile( "movl $0x14,%eax\n\t" "int $0x80\n\t" "mov %%eax,%0\n\t" : "=m" (pid) ... printf("the pid is :%d\n",pid);...}这是今天看代码的时候看到

在linux kernel层调用用户层的方法

首先调用 call_usermodehelper_setup 来创建已初始化的 subprocess_info 结构。 然后使用 call_usermodehelper_exec 函数执行相关的命令,，这里执行 rm /home/xxf/test 命令 static int kernel_...

Jupyter安装R Kernel和调用R

安装和删除Python Kernel 安装Python Kernel by Stackoverflow @cle: 创建虚拟环境，然后手动添加ipykernel， conda create -n py27 python=2.7 conda activate py27 conda install notebook ipykernel ipython ...

Android kernel 打印函数调用栈

1 添加头文件 #include <asm/ptrace.h> 2 在需要打印调用栈的地方添加：dump_stack()

Linux Kernel - open系统调用

系列目录传送门Linux Kernel - 文件系统系列open系统调用简介open系统调用相信大家都不会陌生, ‘Everything is a file - 一切皆是文件’, 是Unix的一个设计特点, 和Unix有着密切关系的Linux也不例外. 设备是文件, ...

.net core 调用c dll_永远被误判的KernelBase.dll

今天维修一个APP故障，在某些电脑中爆发，而另外的电脑却正常运行但这个故障，在国内绝大多数的网站搜索当中都会说：替换KernelBase.dll或者重装系统如果是所有程序都报错的话（连windows自带的程序）KernelBase.dll...

ARM64内核系统调用详解（基于kernel-4.9）

首先来看异常向量表的配置，内核在arch/arm64/kernel/entry.S汇编代码中设置了异常向量表。 /* * Exception vectors. */ .pushsection ".entry.text", "ax" .align 11 ENTRY(vectors) ...

OpenCL调用kernel

OpenCL是一个API，就是应用程序的编程接口，OpenCL和OpenGL很像，这种API可以调用API最底层的数据，通过程序开发调用各种各样的函数，实现各种各样的功能。 对于API来说一般它对硬件设备有比较完整的访问权，以访问...

Linux Kernel - read系统调用

系列目录传送门Linux Kernel - 文件系统系列read系统调用文件系统作为连接用户数据和磁盘等块存储设备的中间层, 必须提供数据的读取和写入接口, 而read系统调用就是用来完成数据读取功能的. 在上一篇Linux Kernel - ...

深度学习CUDA编程干货-kernel的编写和调用

title: 深度学习CUDA编程干货-kernel的编写和调用 date: 2020-06-09 12:08:54 category: 默认分类 本文介绍 深度学习CUDA编程干货-kernel的编写和调用 深度学习CUDA编程干货-kernel的编写和调用 本文由林大佬原创...

start_kernel之前的调用流程（head.s）

本文为本人笔记图，描述汇编start到start_kernel简要过程。     visio图： start_kernel之前.vsd

python ctypes调用kernel32调试游戏进程（二）：枚举线程

python ctypes调用kernel32调试游戏进程（一）：附加和停止附加进程 from ctypes import * import ctypes #这个文件是定义的结构体和一些常量 from debugger_defines import * kernel = ctypes.windll....

机器学习里的kernel是指什么？

1.先给个定义：核函数K（kernel function）就是指K(x, y) = ，其中x和y是n维的输入值，f(·) 是从n维到m维的映射（通常而言，m&...光看这一段还是不明白kernel是什么，用来干什么的...对吧？不要急。...

关于cuda kernel能否调用printf（）函数

笔者最近想编译cuda kernel的代码，不知道何种原因，cuda-gdb总是进入不了cuda 的...可是笔者有疑惑了，cuda __device__ 打头的函数不是不支持调用host端的函数吗，为何printf能调？ 我试着在kernel代码中加上print

java使用JNA调用系统Kernel32.dll文件示例

有时候需要调用系统的API，比如调用GetLocalTime(),GetSystemTime(),...用过查询知道这些方法封装在kernel32.dll文件中，而Java调用dll文件非常麻烦，因此记录下自己整理的一些知识，希望对大家有一点帮助。 1.

Kernel 中调用其它驱动程序的几种方法

1.ZwCreateFile UNICODE_STRING DeviceName; RtlInitUnicodeString( &DeviceName, L"\\Device\\MyDDKDA" ); OBJECT_ATTRIBUTES objectAttribut

linux kernel panic 调用过程

mdm_modem_initialize_dataàmdev->mdm_panic_blk.notifier_call = mdm_panic_prep;  àatomic_notifier_chain_register(&panic_notifier_list,  &mdev->mdm

Android 在 kernel 中重启系统的接口

2019-06-14 关键字：kernel 重启函数、kernel reboot、kernel restart Android 重启系统的方法可太多了。...对于一些特定的场景，比如在 Android 的 kernel 进程里，想要直接发起一个重启命令要如何做呢...

奥特曼大全及关系明细.pdf

此文档有详细奥特曼大全及关系明细

抢茅台脚本以及使用方法

抢茅台的方法，里面有脚本文件和python的安装包，小白可以学习使用，大佬绕行吧，哈哈

初级玩转Linux+Ubuntu(嵌入式开发基础课程)

课程主要面向嵌入式Linux初学者、工程师、学生

主要从一下几方面进行讲解：

1.linux学习路线、基本命令、高级命令

2.shell、vi及vim入门讲解

3.软件安装下载、NFS、Samba、FTP等服务器配置及使用



<

2020美赛题目的ADE.zip

2020年2月份美赛发布的题目ADE，是原始文件，包含有原题目给的数据等。 pdf格式全英，需要的朋友们直接下载就可以啦，里面的东西没有动过，需要翻译可以看我的博客/

相关热词 c# 打印 等比缩放 c#弹出右键菜单 c# 系统托盘图标 c#结构体定义 c# 根据网络定位 c# 清除html标签 c# uwp最小化 c#和java的区别 64位转换 c# 十进制 c# 点击保存草稿的实现