如何在cuda中有效地将向量重复到矩阵
How to efficiently repeat a vector to a matrix in cuda?
我想在cuda中重复一个向量来形成一个矩阵,避免太多的内存拷贝。向量和矩阵都在GPU上分配。
例如:我有一个向量:a = [1 2 3 4]
展开成矩阵:
b = [1 2 3 4;
1 2 3 4;
.......
1 2 3 4]
我所尝试的是分配b的每个元素,但这涉及到大量GPU内存到GPU内存复制。
我知道这在matlab中很容易(使用repmat),但如何在cuda中有效地做到这一点?
EDIT根据注释,我已经将代码更新为可以处理行为主或列为主的底层存储的版本。
这样的代码应该相当快:
// for row_major, blocks*threads should be a multiple of vlen
// for column_major, blocks should be equal to vlen
template <typename T>
__global__ void expand_kernel(const T* vector, const unsigned vlen, T* matrix, const unsigned mdim, const unsigned col_major=0){
if (col_major){
int idx = threadIdx.x+blockIdx.x*mdim;
T myval = vector[blockIdx.x];
while (idx < ((blockIdx.x+1)*mdim)){
matrix[idx] = myval;
idx += blockDim.x;
}
}
else{
int idx = threadIdx.x + blockDim.x * blockIdx.x;
T myval = vector[idx%vlen];
while (idx < mdim*vlen){
matrix[idx] = myval;
idx += gridDim.x*blockDim.x;
}
}
}
这假设您的矩阵的维度是mdim
行x vlen
列(似乎是您在问题中概述的)
您可以调整网格和块的尺寸,以找出在您的特定GPU上最快的工作方式。对于行为主的情况,从每个块256或512个线程开始,并将块的数量设置为等于或大于GPU中SMs数量的4倍。选择网格和块尺寸的乘积等于矢量长度vlen
的整数倍。如果这是困难的,选择一个任意的,但"大"的线程块大小,如250或500,应该不会导致太多的效率损失。
对于列为主的情况,选择每个块256或512个线程,并选择与向量长度vlen
相等的块数。如果vlen
> 65535,您将需要为计算能力3.0或更高版本编译此代码。如果vlen
很小,可能小于32,则该方法的效率可能会显著降低。如果您将每个块的线程数增加到GPU的最大值(512或1024),则会发现一些缓解。可能还有其他"扩展"实现,它们可能更适合列为主的"窄"矩阵情况。例如,对列主代码的直接修改将允许每个向量元素两个块,或者每个向量元素四个块,然后启动的总块将是2* vlen
或4* vlen
,例如。
bandwidthTest
所指示的吞吐量的约90%:
$ cat t546.cu
#include <stdio.h>
#define W 512
#define H (512*1024)
// for row_major, blocks*threads should be a multiple of vlen
// for column_major, blocks should be equal to vlen
template <typename T>
__global__ void expand_kernel(const T* vector, const unsigned vlen, T* matrix, const unsigned mdim, const unsigned col_major=0){
if (col_major){
int idx = threadIdx.x+blockIdx.x*mdim;
T myval = vector[blockIdx.x];
while (idx < ((blockIdx.x+1)*mdim)){
matrix[idx] = myval;
idx += blockDim.x;
}
}
else{
int idx = threadIdx.x + blockDim.x * blockIdx.x;
T myval = vector[idx%vlen];
while (idx < mdim*vlen){
matrix[idx] = myval;
idx += gridDim.x*blockDim.x;
}
}
}
template <typename T>
__global__ void check_kernel(const T* vector, const unsigned vlen, T* matrix, const unsigned mdim, const unsigned col_major=0){
unsigned i = 0;
while (i<(vlen*mdim)){
unsigned idx = (col_major)?(i/mdim):(i%vlen);
if (matrix[i] != vector[idx]) {printf("mismatch at offset %dn",i); return;}
i++;}
}
int main(){
int *v, *m;
cudaMalloc(&v, W*sizeof(int));
cudaMalloc(&m, W*H*sizeof(int));
int *h_v = (int *)malloc(W*sizeof(int));
for (int i = 0; i < W; i++)
h_v[i] = i;
cudaMemcpy(v, h_v, W*sizeof(int), cudaMemcpyHostToDevice);
// test row-major
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);
expand_kernel<<<44, W>>>(v, W, m, H);
cudaEventRecord(stop);
float et;
cudaEventSynchronize(stop);
cudaEventElapsedTime(&et, start, stop);
printf("row-majortime: %fms, bandwidth: %.0fMB/sn", et, W*H*sizeof(int)/(1024*et));
check_kernel<<<1,1>>>(v, W, m, H);
cudaDeviceSynchronize();
// test col-major
cudaEventRecord(start);
expand_kernel<<<W, 256>>>(v, W, m, H, 1);
cudaEventRecord(stop);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&et, start, stop);
printf("col-majortime: %fms, bandwidth: %.0fMB/sn", et, W*H*sizeof(int)/(1024*et));
check_kernel<<<1,1>>>(v, W, m, H, 1);
cudaDeviceSynchronize();
return 0;
}
$ nvcc -arch=sm_20 -o t546 t546.cu
$ ./t546
row-majortime: 13.066944ms, bandwidth: 80246MB/s
col-majortime: 12.806720ms, bandwidth: 81877MB/s
$ /usr/local/cuda/samples/bin/x86_64/linux/release/bandwidthTest
[CUDA Bandwidth Test] - Starting...
Running on...
Device 0: Quadro 5000
Quick Mode
Host to Device Bandwidth, 1 Device(s)
PINNED Memory Transfers
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 5864.2
Device to Host Bandwidth, 1 Device(s)
PINNED Memory Transfers
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 6333.1
Device to Device Bandwidth, 1 Device(s)
PINNED Memory Transfers
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 88178.6
Result = PASS
$
CUDA 6.5, RHEL 5.5
这也可以使用CUBLAS Rank-1更新函数来实现,但它将比上述方法慢得多。
- 写入向量<向量<bool>>
- 函数向量_指针有不同的原型,我可以构建一个吗
- 编译时未启用intel oneApi CUDA支持
- 在统一内存 CUDA C/C++ 中分配 2D 向量
- 如何在 CUDA 中(有效地)将大量向量相互比较
- CUDA:复杂标量 *双稀疏矩阵 *双向量
- 如何将密集的向量转换为CUDA中的稀疏向量
- CUDA矢量减少以处理长度小于512的向量
- Cuda 内核返回向量
- 如何在CUDA内核中添加向量元素
- CUDA 推力:类的:d向量 |错误
- CUDA内核自动调用内核来完成向量加法.为什么?
- 在 cuda 内核中创建向量
- CUDA - STL 向量在传递给使用 nvcc 编译的类时损坏
- 两个单精度浮点向量的点积在 CUDA 内核中产生的结果与在主机上的结果不同
- 获取CUDA推力::transform运算符()函数内向量的索引
- 如何在cuda中有效地将向量重复到矩阵
- 如何从两个数组中生成一对向量,然后使用CUDA/Thrust按对的第一个元素排序
- 在 CUDA 中相互减去向量的元素
- CUDA 中的稀疏矩阵向量乘法