- c - 在位数组中找到第一个零
- linux - Unix 显示有关匹配两种模式之一的文件的信息
- 正则表达式替换多个文件
- linux - 隐藏来自 xtrace 的命令
在这里我尝试使用一些伪代码 self 解释 CUDA 启动参数模型(或执行配置模型),但我不知道是否有一些大错误,所以希望有人帮助审查它,并给我一些建议。感谢先进。
这里是:
/*
normally, we write kernel function like this.
note, __global__ means this function will be called from host codes,
and executed on device. and a __global__ function could only return void.
if there's any parameter passed into __global__ function, it should be stored
in shared memory on device. so, kernel function is so different from the *normal*
C/C++ functions. if I was the CUDA authore, I should make the kernel function more
different from a normal C function.
*/
__global__ void
kernel(float *arr_on_device, int n) {
int idx = blockIdx.x * blockDIm.x + threadIdx.x;
if (idx < n) {
arr_on_device[idx] = arr_on_device[idx] * arr_on_device[idx];
}
}
/*
after this definition, we could call this kernel function in our normal C/C++ codes !!
do you feel something wired ? un-consistant ?
normally, when I write C codes, I will think a lot about the execution process down to
the metal in my mind, and this one...it's like some fragile codes. break the sequential
thinking process in my mind.
in order to make things normal, I found a way to explain: I expand the *__global__ * function
to some pseudo codes:
*/
#define __foreach(var, start, end) for (var = start, var < end; ++var)
__device__ int
__indexing() {
const int blockId = blockIdx.x * gridDim.x + gridDim.x * gridDim.y * blockIdx.z;
return
blockId * (blockDim.x * blockDim.y * blockDim.z) +
threadIdx.z * (blockDim.x * blockDim.y) +
threadIdx.x;
}
global_config =:
{
/*
global configuration.
note the default values are all 1, so in the kernel codes,
we could just ignore those dimensions.
*/
gridDim.x = gridDim.y = gridDim.z = 1;
blockDim.x = blockDim.y = blockDim.z = 1;
};
kernel =:
{
/*
I thought CUDA did some bad evil-detail-covering things here.
it's said that CUDA C is an extension of C, but in my mind,
CUDA C is more like C++, and the *<<<>>>* part is too tricky.
for example:
kernel<<<10, 32>>>(); means kernel will execute in 10 blocks each have 32 threads.
dim3 dimG(10, 1, 1);
dim3 dimB(32, 1, 1);
kernel<<<dimG, dimB>>>(); this is exactly the same thing with above.
it's not C style, and C++ style ? at first, I thought this could be done by
C++'s constructor stuff, but I checked structure *dim3*, there's no proper
constructor for this. this just brroke the semantics of both C and C++. I thought
force user to use *kernel<<<dim3, dim3>>>* would be better. So I'd like to keep
this rule in my future codes.
*/
gridDim = dimG;
blockDim = dimB;
__foreach(blockIdx.z, 0, gridDim.z)
__foreach(blockIdx.y, 0, gridDim.y)
__foreach(blockIdx.x, 0, gridDim.x)
__foreach(threadIdx.z, 0, blockDim.z)
__foreach(threadIdx.y, 0, blockDim.y)
__foreach(threadIdx.x, 0, blockDim.x)
{
const int idx = __indexing();
if (idx < n) {
arr_on_device[idx] = arr_on_device[idx] * arr_on_device[idx];
}
}
};
/*
so, for me, gridDim & blockDim is like some boundaries.
e.g. gridDim.x is the upper bound of blockIdx.x, this is not that obvious for people like me.
*/
/* the declaration of dim3 from vector_types.h of CUDA/include */
struct __device_builtin__ dim3
{
unsigned int x, y, z;
#if defined(__cplusplus)
__host__ __device__ dim3(unsigned int vx = 1, unsigned int vy = 1, unsigned int vz = 1) : x(vx), y(vy), z(vz) {}
__host__ __device__ dim3(uint3 v) : x(v.x), y(v.y), z(v.z) {}
__host__ __device__ operator uint3(void) { uint3 t; t.x = x; t.y = y; t.z = z; return t; }
#endif /* __cplusplus */
};
typedef __device_builtin__ struct dim3 dim3;
最佳答案
CUDA 驱动程序 API
CUDA 驱动程序 API v4.0 及更高版本使用以下函数来控制内核启动:
cuFuncSetCacheConfig
cuFuncSetSharedMemConfig
cuLaunchKernel
在 v4.0 中引入 cuLaunchKernel 之前使用了以下 CUDA Driver API 函数。
cuFuncSetBlockShape()
cuFuncSetSharedSize()
cuParamSet{Size,i,fv}()
cuLaunch
cuLaunchGrid
关于这些函数的更多信息可以在 cuda.h 中找到。
CUresult CUDAAPI cuLaunchKernel(CUfunction f,
unsigned int gridDimX,
unsigned int gridDimY,
unsigned int gridDimZ,
unsigned int blockDimX,
unsigned int blockDimY,
unsigned int blockDimZ,
unsigned int sharedMemBytes,
CUstream hStream,
void **kernelParams,
void **extra);
cuLaunchKernel 将整个启动配置作为参数。
参见 NVIDIA Driver API[执行控制] 1了解更多详情。
CUDA 内核发布
cuLaunchKernel 将1.验证启动参数2.更改共享内存配置3.更改本地内存分配4. 将流同步 token 推送到命令缓冲区,以确保流中的两个命令不重叠4.将启动参数推送到命令缓冲区5.将启动命令推送到命令缓冲区6. 将命令缓冲区提交给设备(在 wddm 驱动程序上,此步骤可能会延迟)7. 在wddm上,内核驱动程序将对设备内存中所需的所有内存进行分页
GPU 将1.验证命令2. 将命令发送到计算工作分配器3. 将启动配置和线程 block 分发给 SM
当所有线程 block 都完成后,工作分配器将刷新缓存以遵循 CUDA 内存模型,并将内核标记为已完成,以便流中的下一个项目可以向前推进。
调度线程 block 的顺序因架构而异。
计算能力 1.x 设备将内核参数存储在共享内存中。计算能力 2.0-3.5 设备将 kenrel 参数存储在常量内存中。
CUDA 运行时 API
CUDA 运行时是一个 C++ 软件库,是在 CUDA Driver API 之上构建工具链。 CUDA 运行时使用以下函数来控制内核启动:
cuda配置调用cudaFuncSetCacheConfigcudaFuncSetSharedMemConfigcuda启动cudaSetup参数
参见 NVIDIA Runtime API[执行控制] 2
<<<>>> CUDA 语言扩展是用于启动内核的最常用方法。
在编译过程中,nvcc 将为使用 <<<>>> 调用的每个内核函数创建一个新的 CPU stub 函数,并将用对 stub 函数的调用替换 <<<>>>。
例如
__global__ void kernel(float* buf, int j)
{
// ...
}
kernel<<<blocks,threads,0,myStream>>>(d_buf,j);
产生
void __device_stub__Z6kernelPfi(float *__par0, int __par1){__cudaSetupArgSimple(__par0, 0U);__cudaSetupArgSimple(__par1, 4U);__cudaLaunch(((char *)((void ( *)(float *, int))kernel)));}
您可以通过将 --keep 添加到您的 nvcc 命令行来检查生成的文件。
cudaLaunch 调用 cuLaunchKernel。
CUDA 动态并行
CUDA CDP 的工作方式类似于上述 CUDA Runtime API。
关于CUDA 内核启动参数解释对了吗?,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/19240658/
我在一个项目中工作,该项目需要 SQL 结果的最佳性能,并且希望优化查询,但经过反复试验后,我在 IN 方面遇到了一些问题。 -- THIS RETURNS NO RESULTS AT ALL. SE
在尝试创建一个实际上非常简单的 SQL 语句时,我发现自己迷失了方向。 我有一个包含 3 个表的数据库: 食谱 - 存储一些用于 cooking 的食谱名称 配料食谱 - 将配料与食谱链接 成分 -
我正在尝试理解 PHP 中的 Hebrev 函数。 https://php.net/manual/en/function.hebrevc.php 它说:“将逻辑希伯来语文本转换为视觉文本”。但我不明白
嗨,我在 Grid view 的 android 文档中发现了一段代码对于以下代码。 gridview.setOnItemClickListener(new OnItemClickListener()
谁能解释一下 InfiniBand 是什么?与以太网相比的主要区别是什么,这些差异如何使其比以太网更快? 在官方description从 mellanox 写到 Introduce InfiniBan
这个问题已经有答案了: How are java increment statements evaluated in complex expressions (1 个回答) 已关闭 8 年前。 我知道
我正在阅读 MySQL 教程,我遇到了这个: SELECT /*! SQL_NO_CACHE */ user FROM users; 为什么优化提示 SQL_NO_CACHE 包含在: /*!
我无法理解$(this),我做了一个剪刀石头布的版本,并应用了 jQuery 让用户在计算机上选择按钮选项。我希望有人能解释一下 $(this) 指的是什么,它是 btn-primary 吗?该函数在
我不是很确定 while(choice == 1 || choice ==2);谁能解释一下。我明白这一点 if(choice ==1) displayMonthly(rainfall); e
let flyRight = CABasicAnimation(keyPath: "position.x") flyRight.toValue = view.bounds.size.width/2 f
目录 解释:int型默认值为0 但我们尝试发现并不能通过: 原因: int的默认值为0,而Integer的默认值为null
我正在处理一个查询,自从一个 SSRS 服务器传输到另一个服务器后,它似乎没有按预期执行,并且 where 语句的一部分中出现了以下行 找出不同之处,或者至少从我能找到的地方来看。 where COA
我正在制作一个退回检测程序,读取退回邮件。我们的设置是发送电子邮件,在发送的邮件中添加一个 noreply@domain.tl。一些收件人不再存在,因此我们想要读取退回邮件,并检测它发送给谁。我已经崩
我有一个关于公式通过控制点弯曲的问题。 如您所知,HTML Canvas 有 quadraticCurveTo(x1, y1, x2, y2)与 x1 and x2作为控制点。 但是,当您尝试使用它绘
我有一个 Emakefile看起来像: %% -- %% %% -- {'/Users/user/projects/custom_test/trunk/*', [debug_info, {out
我有一个非常简单的问题。这不仅适用于 spray-json,而且我已经阅读了 argonaut 和 circe 的类似声明。所以请赐教。 在 spray-json 中,我遇到了 There is no
我正在为视频添加水印。我试图让水印与视频尺寸成比例。我已经使用 scale2ref 看到了十几个不同的答案,但没有解释实际发生了什么,所以我发现很难知道如何实现/更改配置以适应我的情况。 当前覆盖命令
因为我正在学习语言,所以我在玩 Haskell,我只是发现了一些我不理解的东西,我找不到解释。如果我尝试运行此代码: map (`div` 0) [1,2,3,4] 我得到一个除以 0 的异常,这是预
我正在寻找解决错误对象引用未设置到对象实例的步骤/指南。以及问题发生原因的解释。 我正在寻找更一般的解释,所以如果我收到错误,我应该采取什么步骤来查找问题。我经常看到有人提供特定代码段的帖子,而其他人
我最近想升级我的知识React ,所以我从组件生命周期方法开始。让我好奇的第一件事是这个componentWillReceiveProps .所以,文档说当组件接收新的(不一定是更新的) Prop 时
我是一名优秀的程序员,十分优秀!