CUDA:使用 tex2D() 的问题
CUDA: Problems Using tex2D()
我在尝试从 2D CUDA 数组中检索矩阵元素 A(m,n) 时遇到了很多麻烦。有趣的是,当 m = n 时,我获得了正确的元素;即元素沿对角线。否则,我会遇到一些意外的行为:例如,如果我想获取元素 A(13,12),并且我尝试使用 tex2D(tex, row + 0.5f, col + 0.5f) 检索它,我会得到 A(14,11)。据我所知,我正在尽我所能,所以我真的很想知道我哪里出错了。
内核如下。错误发生在前两个tex2D
调用之后,因此其余部分并不真正相关。
texture<float, 2, cudaReadModeElementType> tex_a;
texture<float, 2, cudaReadModeElementType> tex_b;
// Assume that BinaryFunc is multiplication, and AccumulationFunc is addition.
// Then this kernel computes the standard matrix product, and uses prefetching
// with tile sizes given by the template parameter TileSize.
template <unsigned TileSize, class T, class SizeType, class BinaryFunc,
class AccumulationFunc>
__global__ void
matrix_prod_tex_prefetch(T* c, const SizeType dim, BinaryFunc binary_func,
AccumulationFunc accum_func)
{
__shared__ T as[TileSize][TileSize];
__shared__ T bs[TileSize][TileSize];
SizeType row = blockIdx.y * TileSize + threadIdx.y;
SizeType col = blockIdx.x * TileSize + threadIdx.x;
T p = 0;
T l = tex2D(tex_a, row + 0.5f, threadIdx.x + 0.5f);
T m = tex2D(tex_b, threadIdx.y + 0.5f, col + 0.5f);
__syncthreads();
for (SizeType i = 1; i != dim / TileSize; ++i) {
as[threadIdx.y][threadIdx.x] = l;
bs[threadIdx.y][threadIdx.x] = m;
__syncthreads();
l = tex2D(tex_a, row + 0.5f, i * TileSize + threadIdx.x + 0.5f);
m = tex2D(tex_b, i * TileSize + threadIdx.y + 0.5f, col + 0.5f);
for (SizeType k = 0; k != TileSize; ++k) {
p = accum_func(p, binary_func(
as[threadIdx.y][k],
bs[k][threadIdx.x]
));
}
__syncthreads();
}
as[threadIdx.y][threadIdx.x] = l;
bs[threadIdx.y][threadIdx.x] = m;
__syncthreads();
for (SizeType k = 0; k != TileSize; ++k) {
p = accum_func(p, binary_func(
as[threadIdx.y][k],
bs[k][threadIdx.x]
));
}
c[dim * row + col] = p;
}
答:将threadIdx.x与threadIdx.y交换。最终,它归结为语义问题:纹理使用索引作为沿我们都熟悉的 x 轴和 y 轴的偏移量。矩阵使用索引来引用行和列的索引。本质上,基向量是交换的。
请注意,虽然将 threadIdx.x 和 threadIdx.y 的使用换成 1D 内存布局可能会产生等效的结果,但在此过程中可能会丢失合并的内存访问模式。
相关文章:
- 警告处理为错误这里有什么问题
- 最小硬币更换问题(自上而下方法)
- 为"adjacent"变量赋值时出现问题
- 我的神经网络不起作用 [XOR 问题]
- 在Ubuntu 16.04上安装Cilk时出现问题
- C++我的数学有什么问题,为什么我的代码不能正确循环
- 编译包含字符串的代码时遇到问题
- Project Euler问题4的错误解决方案
- 问题:什么是QAbstractItemView::NoEditTriggers的反面
- 在编译C++代码(具有dlib和opencv)到WASM时面临问题
- 在进程中对同一管道进行读取和写入时C++管道出现问题
- 静态数据成员的问题-修复链接错误会导致编译器错误
- C++ 雷神库 - 使用资源加载器类时出现问题(不命名类型)
- 一个关于在C++中重载布尔运算符的问题
- 首要问题的答案让值班员搞错了
- setlocale的C++土耳其字符串问题
- 如何重构类层次结构以避免菱形问题
- 基于boost的程序的静态链接——zlib问题
- C++格式化输出问题
- CUDA:使用 tex2D() 的问题