我是Thrust的新手.我看到所有Thrust演示文稿和示例仅显示主机代码.
我想知道我是否可以将device_vector传递给我自己的内核？怎么样？如果是,内核/设备代码中允许的操作是什么？
如何有效地规范化CUDA中的矩阵列？
我的矩阵存储在column-major中,典型大小为2000x200.
该操作可以用以下matlab代码表示.
A = rand(2000,200);
A = exp(A);
A = A./repmat(sum(A,1), [size(A,1) 1]);
这可以通过Thrust,cuBLAS和/或cuNPP有效地完成吗？
包括4个内核的快速实现如下所示.
想知道这些是否可以在1或2个内核中完成以提高性能,尤其是对于由cublasDgemv()实现的列求和步骤.
#include <cuda.h>
#include <curand.h>
#include <cublas_v2.h>
#include <thrust/device_vector.h>
#include <thrust/device_ptr.h>
#include <thrust/transform.h>
#include <thrust/iterator/constant_iterator.h>
#include <math.h>
struct Exp
{
__host__ __device__ void operator()(double& x)
{
x = exp(x);
}
};
struct Inv
{
__host__ __device__ void operator()(double& x)
{
x = (double) 1.0 / x;
}
};
int main()
{
cudaDeviceSetCacheConfig(cudaFuncCachePreferShared);
cublasHandle_t hd;
curandGenerator_t rng;
cublasCreate(&hd);
curandCreateGenerator(&rng, CURAND_RNG_PSEUDO_DEFAULT);
const size_t m = 2000, …我有一个相当简单的问题,但我无法找到一个优雅的解决方案.
我有一个Thrust代码,它生成c包含值的相同大小的向量.假设这些c向量中的每一个都有一个索引.我想为每个向量位置获取c值为最低的向量的索引:
例:
C0 = (0,10,20,3,40)
C1 = (1,2 ,3 ,5,10)
我会得到一个包含C具有最低值的向量索引的向量:
result = (0,1 ,1 ,0,1)
我已经考虑过使用推力zip迭代器来做这件事,但是已经遇到了问题:我可以压缩所有c向量并实现任意转换,它接受一个元组并返回其最低值的索引,但是:
10元素,并且可以存在比10 c矢量更多的元素.然后我考虑这样做:不是使用c单独的向量,而是将它们全部附加到单个向量中C,然后生成引用位置的键并按键执行稳定排序,这将从同一位置重新组合向量条目.在示例中,将给出:
C = (0,10,20,3,40,1,2,3,5,10)
keys = (0,1 ,2 ,3,4 ,0,1,2,3,4 )
after stable sort by key:
output = (0,1,10,2,20,3,3,5,40,10)
keys = (0,0,1 ,1,2 ,2,3,3,4 ,4 )
然后使用向量中的位置生成键,使用向量的索引压缩输出,c然后使用自定义函数执行按键缩减,对于每个缩减,输出具有最低值的索引.在示例中:
input = (0,1,10,2,20,3,3,5,40,10)
indexes= (0,1,0 ,1,0 ,1,0,1,0 ,1)
keys = (0,0,1 ,1,2 …我正在使用CUDA cuBLAS来执行矩阵运算.
我需要对矩阵的行(或列)求和.目前我是通过将矩阵与一个向量相乘来实现的,但这似乎并不那么有效.
有没有更好的方法？找不到任何东西cuBLAS.
谢谢.
在GPU上的一些计算中,我需要缩放矩阵中的行,以便给定行中的所有元素总和为1.
| a1,1 a1,2 ... a1,N | | alpha1*a1,1 alpha1*a1,2 ... alpha1*a1,N | | a2,1 a2,2 ... a2,N | => | alpha2*a2,1 alpha2*a2,2 ... alpha2*a2,N | | . . | | . . | | aN,1 aN,2 ... aN,N | | alphaN*aN,1 alphaN*aN,2 ... alphaN*a …
假设我有两个device_vector <byte>数组,d_keys并且d_data.
d_data例如,如果是扁平的2D 3x5阵列(例如{1,2,3,4,5,6,7,8,9,8,7,6,5,4,3})并且d_keys是1D阵列大小为5(例如{1,0,0,1,1}),我如何进行减少,这样如果相应的d_keys值为1 ,我最终只会按行添加值(例如,结束结果为{10,23,14})？
该sum_rows.cu例如允许我加入的每一个值d_data,但是这并不完全正确.
或者,我可以在每行的基础上使用a zip_iterator并一次合并d_keys一行d_data,然后执行a transform_reduce,仅在键值为1时添加,但是我必须循环遍历d_data数组.
我真正需要的是某种transform_reduce_by_key不是内置的功能,但肯定必须有一种方法来实现它!
我正在尝试沿着2D矩阵的行方向实现缩减.我从stackoverflow上找到的代码开始(非常感谢Robert!)
thrust :: max_element比较cublasIsamax慢 - 更有效的实现？
上面的链接显示了一个在单行上执行缩减的自定义内核.它将输入行分为多行,每行有1024个线程.效果很好.
对于2D情况,一切都是相同的,除了现在有一个网格尺寸.所以每个块的y维度仍然是1.问题是当我尝试将数据写入每个块内的共享内存(在代码中的"max_idx_kernel_reduction_within_block"内核中)时,需要很长时间(超过(行数)*(在1行上执行减少所需的时间.我宁愿运行for循环).我知道我有很多元素,但我期待比这更快的东西.
我不认为内存访问模式是一个问题,但我听说TOTAL共享内存量可能是限制？:CUDA:合并全局内存访问速度比共享内存快吗？另外,分配大型共享内存阵列会减慢程序的速度吗？
任何使我的代码更快的建议(第一个内核是瓶颈)？非常感谢,非常感谢!!
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <iostream>
#include <cuda_runtime.h>
#define NCOLS 163317 // number of columns
#define NROWS 8 // number of rows
#define nTPB 1024 // Threads per Block. nTPB should be a power-of-2
#define MAX_BLOCKS_X ((NCOLS/nTPB)+1) // # of blocks I will launch
#define MIN(a,b) ((a>b)?b:a)
#define FLOAT_MIN -1.0f // lowest anticipated number of the data. Values in array will be compared with this and updated …