- iOS/Objective-C 元类和类别
- objective-c - -1001 错误,当 NSURLSession 通过 httpproxy 和/etc/hosts
- java - 使用网络类获取 url 地址
- ios - 推送通知中不播放声音
我一直在网上搜索和搜索,但似乎无法找到我正在寻找的答案。我有一个特别的问题。
我正在编辑它以简化问题并希望它更具可读性和可理解性。
假设我有 5000 个 20x20 的对称密集矩阵。我想在 CUDA 中创建一个内核,让每个线程负责计算每个对称矩阵的特征值。
如果可能,CUDA 内核的示例代码会很棒。
如有任何帮助/建议,我们将不胜感激!
谢谢,
乔纳森
最佳答案
I would like to create a kernel in CUDA that will have each thread responsible for calculating the eigenvalues for each of the symmetric matrices.
我怀疑这是否是最快的方法,但它可能适用于非常小的矩阵。即使在这种情况下,也可能会进行一些数据存储优化(跨线程交错全局数据),但这会使事情复杂化。
如前所述,该请求可以映射到“令人尴尬的并行”算法中,其中每个线程处理一个完全独立的问题。我们只需要找到合适的单线程“捐赠者代码”。快速谷歌搜索后,我遇到了 this .修改该代码以以这种与线程无关的方式运行是相当简单的。我们只需要借用 3 个例程(jacobi_eigenvalue
、r8mat_diag_get_vector
和 r8mat_identity
),并用 __host__ __device__
适本地装饰这些例程在 GPU 上使用,同时没有其他更改。
有问题的代码似乎是由佛罗里达州立大学的 J Burkardt 许可的 GNU LGPL。因此,鉴于此,并关注conventional wisdom我没有在这个答案中包含任何大量的代码。但是您应该能够使用我提供的说明通过实验重建我的结果。
注意:我不确定使用此代码的法律后果是什么,它声称已获得 GNU LGPL 许可。您应该确保遵守 any necessary requirements如果您选择使用此代码或其中的一部分。我在这里使用它的主要目的是展示单线程问题解决器的相对微不足道的“令人尴尬的并行”扩展的概念。
通过 here 重构我的完整代码应该是微不足道的并将 3 个指示的函数复制粘贴到剩余代码骨架中指示的位置。但这不会改变前面提到的任何通知/免责声明。使用它需要您自担风险。
同样,从性能的角度来看,不进行其他更改可能不是最好的主意,但它会产生微不足道的工作量,并且可以作为一个可能有用的起点。一些可能的优化可能是:
new
和delete
函数,用固定分配代替(这很容易做到)无论如何,使用上面修饰过的捐赠者代码,我们只需要在它周围包裹一个简单的内核 (je
),以启动每个线程在单独的数据集(即矩阵)上运行,并且每个线程产生自己的一组特征值(和特征向量 - 对于这个特定的代码库)。
出于测试目的,我将其设计为仅使用 3 个线程和 3 个 4x4 矩阵,但将其扩展到您希望的任意多个矩阵/线程应该是微不足道的。
为了简洁起见,我省略了 the usual error checking ,但我建议您使用它,或者如果您进行了任何修改,至少使用 cuda-memcheck
运行您的代码。
我还构建了代码来向上调整设备堆大小以适应内核中的 new
操作,具体取决于矩阵(即线程)的数量和矩阵维度。如果您进行了上述第二项优化,您可能会删除它。
t1177.cu:
#include <stdio.h>
#include <iostream>
const int num_mat = 3; // total number of matrices = total number of threads
const int N = 4; // square symmetric matrix dimension
const int nTPB = 256; // threads per block
// test symmetric matrices
double a1[N*N] = {
4.0, -30.0, 60.0, -35.0,
-30.0, 300.0, -675.0, 420.0,
60.0, -675.0, 1620.0, -1050.0,
-35.0, 420.0, -1050.0, 700.0 };
double a2[N*N] = {
4.0, 0.0, 0.0, 0.0,
0.0, 1.0, 0.0, 0.0,
0.0, 0.0, 3.0, 0.0,
0.0, 0.0, 0.0, 2.0 };
double a3[N*N] = {
-2.0, 1.0, 0.0, 0.0,
1.0, -2.0, 1.0, 0.0,
0.0, 1.0, -2.0, 1.0,
0.0, 0.0, 1.0, -2.0 };
/* ---------------------------------------------------------------- */
//
// the following functions come from here:
//
// https://people.sc.fsu.edu/~jburkardt/cpp_src/jacobi_eigenvalue/jacobi_eigenvalue.cpp
//
// attributed to j. burkardt, FSU
// they are unmodified except to add __host__ __device__ decorations
//
//****************************************************************************80
__host__ __device__
void r8mat_diag_get_vector ( int n, double a[], double v[] )
/* PASTE IN THE CODE HERE, FROM THE ABOVE LINK, FOR THIS FUNCTION */
//****************************************************************************80
__host__ __device__
void r8mat_identity ( int n, double a[] )
/* PASTE IN THE CODE HERE, FROM THE ABOVE LINK, FOR THIS FUNCTION */
//****************************************************************************80
__host__ __device__
void jacobi_eigenvalue ( int n, double a[], int it_max, double v[],
double d[], int &it_num, int &rot_num )
/* PASTE IN THE CODE HERE, FROM THE ABOVE LINK, FOR THIS FUNCTION */
// end of FSU code
/* ---------------------------------------------------------------- */
__global__ void je(int num_matr, int n, double *a, int it_max, double *v, double *d){
int idx = threadIdx.x+blockDim.x*blockIdx.x;
int it_num;
int rot_num;
if (idx < num_matr){
jacobi_eigenvalue(n, a+(idx*n*n), it_max, v+(idx*n*n), d+(idx*n), it_num, rot_num);
}
}
void initialize_matrix(int mat_id, int n, double *mat, double *v){
for (int i = 0; i < n*n; i++) *(v+(mat_id*n*n)+i) = mat[i];
}
void print_vec(int vec_id, int n, double *d){
std::cout << "matrix " << vec_id << " eigenvalues: " << std::endl;
for (int i = 0; i < n; i++) std::cout << i << ": " << *(d+(n*vec_id)+i) << std::endl;
std::cout << std::endl;
}
int main(){
// make sure device heap has enough space for in-kernel new allocations
const int heapsize = num_mat*N*sizeof(double)*2;
const int chunks = heapsize/(8192*1024) + 1;
cudaError_t cudaStatus = cudaDeviceSetLimit(cudaLimitMallocHeapSize, (8192*1024) * chunks);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "set device heap limit failed!");
}
const int max_iter = 1000;
double *h_a, *d_a, *h_v, *d_v, *h_d, *d_d;
h_a = (double *)malloc(num_mat*N*N*sizeof(double));
h_v = (double *)malloc(num_mat*N*N*sizeof(double));
h_d = (double *)malloc(num_mat* N*sizeof(double));
cudaMalloc(&d_a, num_mat*N*N*sizeof(double));
cudaMalloc(&d_v, num_mat*N*N*sizeof(double));
cudaMalloc(&d_d, num_mat* N*sizeof(double));
memset(h_a, 0, num_mat*N*N*sizeof(double));
memset(h_v, 0, num_mat*N*N*sizeof(double));
memset(h_d, 0, num_mat* N*sizeof(double));
initialize_matrix(0, N, a1, h_a);
initialize_matrix(1, N, a2, h_a);
initialize_matrix(2, N, a3, h_a);
cudaMemcpy(d_a, h_a, num_mat*N*N*sizeof(double), cudaMemcpyHostToDevice);
cudaMemcpy(d_v, h_v, num_mat*N*N*sizeof(double), cudaMemcpyHostToDevice);
cudaMemcpy(d_d, h_d, num_mat* N*sizeof(double), cudaMemcpyHostToDevice);
je<<<(num_mat+nTPB-1)/nTPB, nTPB>>>(num_mat, N, d_a, max_iter, d_v, d_d);
cudaMemcpy(h_d, d_d, num_mat*N*sizeof(double), cudaMemcpyDeviceToHost);
print_vec(0, N, h_d);
print_vec(1, N, h_d);
print_vec(2, N, h_d);
return 0;
}
编译和示例运行:
$ nvcc -o t1177 t1177.cu
$ cuda-memcheck ./t1177
========= CUDA-MEMCHECK
matrix 0 eigenvalues:
0: 0.166643
1: 1.47805
2: 37.1015
3: 2585.25
matrix 1 eigenvalues:
0: 1
1: 2
2: 3
3: 4
matrix 2 eigenvalues:
0: -3.61803
1: -2.61803
2: -1.38197
3: -0.381966
========= ERROR SUMMARY: 0 errors
$
输出对我来说似乎是合理的,大部分匹配输出 here .
关于c++ - 使用 CUDA 并行特征值求解器,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/38112248/
这是我关于 Stack Overflow 的第一个问题,这是一个很长的问题。 tl;dr 版本是:我如何使用 thrust::device_vector如果我希望它存储不同类型的对象 DerivedC
我已使用 cudaMalloc 在设备上分配内存并将其传递给内核函数。是否可以在内核完成执行之前从主机访问该内存? 最佳答案 我能想到的在内核仍在执行时启动 memcpy 的唯一方法是在与内核不同的流
是否可以在同一节点上没有支持 CUDA 的设备的情况下编译 CUDA 程序,仅使用 NVIDIA CUDA Toolkit...? 最佳答案 你的问题的答案是肯定的。 nvcc编译器驱动程序与设备的物
我不知道 cuda 不支持引用参数。我的程序中有这两个函数: __global__ void ExtractDisparityKernel ( ExtractDisparity& es)
我正在使用 CUDA 5.0。我注意到编译器将允许我在内核中使用主机声明的 int 常量。但是,它拒绝编译任何使用主机声明的 float 常量的内核。有谁知道这种看似差异的原因? 例如,下面的代码可以
自从 CUDA 9 发布以来,显然可以将不同的线程和 block 分组到同一组中,以便您可以一起管理它们。这对我来说非常有用,因为我需要启动一个包含多个 block 的内核并等待所有 block 都同
我需要在 CUDA 中执行三线性插值。这是问题定义: 给定三个点向量:x[nx]、y[ny]、z[nz] 和一个函数值矩阵func[nx][ny][nz],我想在 x、y 范围之间的一些随机点处找到函
我认为由于 CUDA 可以执行 64 位 128 位加载/存储,因此它可能具有一些用于加/减/等的内在函数。像 float3 这样的向量类型,在像 SSE 这样更少的指令中。 CUDA 有这样的功能吗
我有一个问题,每个线程 block (一维)必须对共享内存内的一个数组进行扫描,并执行几个其他任务。 (该数组最多有 1024 个元素。) 有没有支持这种操作的好库? 我检查了 Thrust 和 Cu
我对线程的形成和执行方式有很多疑惑。 首先,文档将 GPU 线程描述为轻量级线程。假设我希望将两个 100*100 矩阵相乘。如果每个元素都由不同的线程计算,则这将需要 100*100 个线程。但是,
我正在尝试自己解决这个问题,但我不能。 所以我想听听你的建议。 我正在编写这样的内核代码。 VGA 是 GTX 580。 xxxx >> (... threadNum ...) (note. Shar
查看 CUDA Thrust 代码中的内核启动,似乎它们总是使用默认流。我可以让 Thrust 使用我选择的流吗?我在 API 中遗漏了什么吗? 最佳答案 我想在 Thrust 1.8 发布后更新 t
我想知道 CUDA 应用程序的扭曲调度顺序是否是确定性的。 具体来说,我想知道在同一设备上使用相同输入数据多次运行同一内核时,warp 执行的顺序是否会保持不变。如果没有,是否有任何东西可以强制对扭曲
一个 GPU 中可以有多少个 CUDA 网格? 两个网格可以同时存在于 GPU 中吗?还是一台 GPU 设备只有一个网格? Kernel1>(dst1, param1); Kernel1>(dst2,
如果我编译一个计算能力较低的 CUDA 程序,例如 1.3(nvcc 标志 sm_13),并在具有 Compute Capability 2.1 的设备上运行它,它是否会利用 Compute 2.1
固定内存应该可以提高从主机到设备的传输速率(api 引用)。但是我发现我不需要为内核调用 cuMemcpyHtoD 来访问这些值,也不需要为主机调用 cuMemcpyDtoA 来读取值。我不认为这会奏
我希望对 CUDA C 中负载平衡的最佳实践有一些一般性的建议和说明,特别是: 如果经纱中的 1 个线程比其他 31 个线程花费的时间长,它会阻止其他 31 个线程完成吗? 如果是这样,多余的处理能力
CUDA 中是否有像 opencl 一样的内置交叉和点积,所以 cuda 内核可以使用它? 到目前为止,我在规范中找不到任何内容。 最佳答案 您可以在 SDK 的 cutil_math.h 中找到这些
有一些与我要问的问题类似的问题,但我觉得它们都没有触及我真正要寻找的核心。我现在拥有的是一种 CUDA 方法,它需要将两个数组定义到共享内存中。现在,数组的大小由在执行开始后读入程序的变量给出。因此,
经线是 32 根线。 32 个线程是否在多处理器中并行执行? 如果 32 个线程没有并行执行,则扭曲中没有竞争条件。 在经历了一些例子后,我有了这个疑问。 最佳答案 在 CUDA 编程模型中,warp
我是一名优秀的程序员,十分优秀!