如何有效地将向量重复到库达的矩阵中？
如何有效地将向量重复到库达的矩阵中？
提问于 2014-08-22 17:27:34
我想重复一个向量，以形成一个矩阵在库达，避免太多的模件。在GPU上分配向量和矩阵。
例如：
我有一个向量：
a = [1 2 3 4]将其扩展到一个矩阵：
b = [1 2 3 4;
1 2 3 4;
.......
1 2 3 4]我尝试的是分配b的每个元素，但是这涉及到大量的GPU内存到GPU内存副本。
我知道这在matlab中很容易(使用爬虫垫)，但是如何在cuda中高效地实现呢？我没有在库布拉斯找到任何例行程序。
回答 1
Stack Overflow用户
回答已采纳
发布于 2014-08-22 18:32:53
基于注释的编辑，我已经将代码更新为将处理行主或列主底层存储的版本。
像这样的事情应该相当快：
// 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 %d\n",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/s\n", 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/s\n", 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等级-1更新功能实现，但要比上面的方法慢得多。
页面原文内容由Stack Overflow提供。腾讯云小微IT领域专用引擎提供翻译支持
原文链接：
https://stackoverflow.com/questions/25452519
复制相关文章
相似问题
腾讯云开发者
