cuda - 将两个 CUDA 内核合并为一个
问题描述
我正在使用 CUDA 来计算out = C(b(A(in)))
,其中函数A
和C
是卷积,并且b
是逐元素函数。一个玩具示例是:
#define N 1000
__device__ float b(float d_in){return min(d_in + 10.0f, 100.0f);}
__global__ void bA(float *d_in, float *d_out){
const int x = threadIdx.x + blockIdx.x * blockDim.x;
if (x >= N) return;
// replicate boundary
int x_left = max(x-1, 0); int x_right = min(x+1, N-1);
d_out[x] = b( d_in[x_left] + d_in[x] + d_in[x_right] );
}
__global__ void C(float *d_in, float *d_out){
const int x = threadIdx.x + blockIdx.x * blockDim.x;
if (x >= N) return;
// replicate boundary
int x_left = max(x-1, 0); int x_right = min(x+1, N-1);
d_out[x] = d_in[x_left] + d_in[x] + d_in[x_right];
}
void myfunc(float *d_data, float *d_temp){
dim3 threads(256);
dim3 blocks( (N + threads.x - 1) / threads.x ); // divide up
// kernels that I would like to merge into one:
bA<<<blocks, threads>>>(d_data, d_temp);
C <<<blocks, threads>>>(d_temp, d_data);
}
像这样的计算需要一个额外的变量d_temp
,这是我不想要的。所以我想将这些内核合并为一个,即一个用于计算的内核C(b(A(in)))
。
一个困难是,我怎样才能保存 的临时结果b(A(in))
,然后执行卷积函数C()
?我曾尝试使用共享内存,但不知道如何将临时结果加载b(A(in))
到共享内存。例如:
#define BLOCK_SIZE 32
__global__ void CbA(float *d_in, float *d_out){
const int x = threadIdx.x + blockIdx.x * blockDim.x;
if (x >= N) return;
// replicate boundary
int x_left = max(x-1, 0); int x_right = min(x+1, N-1);
// temp result for b(A(in))
float temp = b( d_in[x_left] + d_in[x] + d_in[x_right] );
// shared memory for convolution (stencil size of 3)
__shared__ float shmem[BLOCK_SIZE+2];
// load center part to shared memory
shmem[threadIdx.x+1] = temp;
// but how to load boundary parts from temp to shmem?
// ...
__syncthreads();
// perform function C()
// ...
}
非常感谢任何建议或提示。
解决方案
首先评论一下
// load center part to shared memory
shmem[threadIdx.x+1] = temp;
我会称之为保存到共享内存......
除此之外还有一些想法:
使用块中的第一个和最后一个线程仅计算b(A(in))
当然,您必须在计算x
( const int x = threadIdx.x + blockIdx.x * (blockDim.x-2);
) 时考虑这一点,并使用更多线程/块调用内核。
然后,当您执行时,每个块将有两个线程空闲C()
。但这不应该有很大的影响。
这是内核。如果您尝试可视化计算流程,则更容易理解。
__global__ void CbA(float *d_in, float *d_out)
{
const int x = threadIdx.x + blockIdx.x * (blockDim.x - 2);
if (x >= N) return;
int x_left = max(x-1, 0); int x_right = min(x+1, N-1);
float temp = b( d_in[x_left] + d_in[x] + d_in[x_right] );
__shared__ float shmem[BLOCK_SIZE]; // = 256
shmem[threadIdx.x] = temp;
__syncthreads();
if (threadIdx.x > 0 && threadIdx.x < blockDim.x-1)
d_out[x-1] = shmem[threadIdx.x-1] + d_in[threadIdx.x] + d_in[threadIdx.x+1];
}
让块中的一个线程也执行块b(A())
的“边界部分”
但是,对于每个块,您只会使用 32 个线程中的 1 个来进行该计算。最坏的情况是整个 SM 在额外计算时的比率为 1/32。
...
// but how to load boundary parts from temp to shmem?
if (threadIdx.x == 0)
{
{
const int x = 0 + blockIdx.x * blockDim.x;
int x_left = max(x-1, 0); int x_right = min(x+1, N-1);
float temp = b( d_in[x_left] + d_in[x] + d_in[x_right] );
shmem[0] = temp;
}
{
const int x = blockDim.x-1 + blockIdx.x * blockDim.x;
int x_left = max(x-1, 0); int x_right = min(x+1, N-1);
float temp = b( d_in[x_left] + d_in[x] + d_in[x_right] );
shmem[blockDim.x-1] = temp;
}
}
// perform function C()
...
避免使用共享内存
(至少在您的简化示例中) 的值temp
是非常简单计算的结果。也许最好C()
在该线程中本地计算您需要在该线程中执行的所有值。
__global__ void CbA(float *d_in, float *d_out)
{
const int x = threadIdx.x + blockIdx.x * blockDim.x;
if (x >= N) return;
float temp[3];
for (int i(0); i < 3; ++i)
{
int x_left = max(x-1-1+i, 0); int x_right = min(x+1-1+i, N-1);
temp[i] = b( d_in[x_left] + d_in[x-1+i] + d_in[x_right] );
}
// perform function C()
...
}
推荐阅读
- python - super() 将如何影响 MRO 中的订单?
- ssl - ClientHello 后 SSL 握手失败
- javascript - 鼠标离开时在锚标签上反向悬停
- python - 多次调用函数时出现 RecursionError
- string - 从 GNU Octave 中的字符串数组中访问字符串
- regex - 如果我有逗号分隔符,如何在谷歌工作表中拆分信息,但同时我在文本中有逗号?
- javascript - 如何在反应组件上设置Timout
- python-3.x - Python 3 需要帮助计算 tKinter 中的所有 3 个函数
- python - python 中的 *args 和 **kwargs
- android - recyclerview在android中滚动不顺畅