- html - 出于某种原因,IE8 对我的 Sass 文件中继承的 html5 CSS 不友好?
- JMeter 在响应断言中使用 span 标签的问题
- html - 在 :hover and :active? 上具有不同效果的 CSS 动画
- html - 相对于居中的 html 内容固定的 CSS 重复背景?
我在类作业中遇到了一个简单的 CUDA 问题,但教授添加了一个可选任务来使用共享内存来实现相同的算法。我无法在截止日期之前完成它(因为上交日期是一周前),但我仍然很好奇,所以现在我要上网 ;)。
基本任务是在顺序和 CUDA 中实现红黑连续过度松弛的 SCSS 版本,确保在两者中获得相同的结果,然后比较加速。就像我说的,使用共享内存来做是一个可选的 +10% 附加组件。
我将发布我的工作版本和我尝试做的伪代码,因为我目前没有代码,但如果有人需要,我可以稍后用实际代码更新它。
在任何人说出来之前:是的,我知道使用 CUtil 是蹩脚的,但它使比较和计时器更容易。
工作全局内存版本:
#include <stdlib.h>
#include <stdio.h>
#include <cutil_inline.h>
#define N 1024
__global__ void kernel(int *d_A, int *d_B) {
unsigned int index_x = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int index_y = blockIdx.y * blockDim.y + threadIdx.y;
// map the two 2D indices to a single linear, 1D index
unsigned int grid_width = gridDim.x * blockDim.x;
unsigned int index = index_y * grid_width + index_x;
// check for boundaries and write out the result
if((index_x > 0) && (index_y > 0) && (index_x < N-1) && (index_y < N-1))
d_B[index] = (d_A[index-1]+d_A[index+1]+d_A[index+N]+d_A[index-N])/4;
}
main (int argc, char **argv) {
int A[N][N], B[N][N];
int *d_A, *d_B; // These are the copies of A and B on the GPU
int *h_B; // This is a host copy of the output of B from the GPU
int i, j;
int num_bytes = N * N * sizeof(int);
// Input is randomly generated
for(i=0;i<N;i++) {
for(j=0;j<N;j++) {
A[i][j] = rand()/1795831;
//printf("%d\n",A[i][j]);
}
}
cudaEvent_t start_event0, stop_event0;
float elapsed_time0;
CUDA_SAFE_CALL( cudaEventCreate(&start_event0) );
CUDA_SAFE_CALL( cudaEventCreate(&stop_event0) );
cudaEventRecord(start_event0, 0);
// sequential implementation of main computation
for(i=1;i<N-1;i++) {
for(j=1;j<N-1;j++) {
B[i][j] = (A[i-1][j]+A[i+1][j]+A[i][j-1]+A[i][j+1])/4;
}
}
cudaEventRecord(stop_event0, 0);
cudaEventSynchronize(stop_event0);
CUDA_SAFE_CALL( cudaEventElapsedTime(&elapsed_time0,start_event0, stop_event0) );
h_B = (int *)malloc(num_bytes);
memset(h_B, 0, num_bytes);
//ALLOCATE MEMORY FOR GPU COPIES OF A AND B
cudaMalloc((void**)&d_A, num_bytes);
cudaMalloc((void**)&d_B, num_bytes);
cudaMemset(d_A, 0, num_bytes);
cudaMemset(d_B, 0, num_bytes);
//COPY A TO GPU
cudaMemcpy(d_A, A, num_bytes, cudaMemcpyHostToDevice);
// create CUDA event handles for timing purposes
cudaEvent_t start_event, stop_event;
float elapsed_time;
CUDA_SAFE_CALL( cudaEventCreate(&start_event) );
CUDA_SAFE_CALL( cudaEventCreate(&stop_event) );
cudaEventRecord(start_event, 0);
// TODO: CREATE BLOCKS AND THREADS AND INVOKE GPU KERNEL
dim3 block_size(256,1,1); //values experimentally determined to be fastest
dim3 grid_size;
grid_size.x = N / block_size.x;
grid_size.y = N / block_size.y;
kernel<<<grid_size,block_size>>>(d_A,d_B);
cudaEventRecord(stop_event, 0);
cudaEventSynchronize(stop_event);
CUDA_SAFE_CALL( cudaEventElapsedTime(&elapsed_time,start_event, stop_event) );
//COPY B BACK FROM GPU
cudaMemcpy(h_B, d_B, num_bytes, cudaMemcpyDeviceToHost);
// Verify result is correct
CUTBoolean res = cutComparei( (int *)B, (int *)h_B, N*N);
printf("Test %s\n",(1 == res)?"Passed":"Failed");
printf("Elapsed Time for Sequential: \t%.2f ms\n", elapsed_time0);
printf("Elapsed Time for CUDA:\t%.2f ms\n", elapsed_time);
printf("CUDA Speedup:\t%.2fx\n",(elapsed_time0/elapsed_time));
cudaFree(d_A);
cudaFree(d_B);
free(h_B);
cutilDeviceReset();
}
#define N 1024
__global__ void kernel(int *d_A, int *d_B, int width) {
//assuming width is 64 because that's the biggest number I can make it
//each MP has 48KB of shared mem, which is 12K ints, 32 threads/warp, so max 375 ints/thread?
__shared__ int A_sh[3][66];
//get x and y index and turn it into linear index
for(i=0; i < width+2; i++) //have to load 2 extra values due to the -1 and +1 in algo
A_sh[index_y%3][i] = d_A[index+i-1]; //so A_sh[index_y%3][0] is actually d_A[index-1]
__syncthreads(); //and hope that previous and next row have been loaded by other threads in the block?
//ignore boundary conditions because it's pseudocode
for(i=0; i < width; i++)
d_B[index+i] = A_sh[index_y%3][i] + A_sh[index_y%3][i+2] + A_sh[index_y%3-1][i+1] + A_sh[index_y%3+1][i+1];
}
main(){
//same init as above until threads/grid init
dim3 threadsperblk(32,16);
dim3 numblks(32,64);
kernel<<<numblks,threadsperblk>>>(d_A,d_B,64);
//rest is the same
}
最佳答案
在 CUDA 中充分利用这些模板运算符的关键是数据重用。我发现最好的方法通常是让每个块“走过”网格的一个维度。在块将初始数据块加载到共享内存后,只需要从全局内存中读取一个维度(因此行主序二维问题中的行),以便在共享内存中为第二个和后续提供必要的数据行计算。其余的数据可以重复使用。为了可视化共享内存缓冲区如何查看此类算法的前四个步骤:
template<int width>
__device__ void rowfetch(int *in, int *out, int col)
{
*out = *in;
if (col == 1) *(out-1) = *(in-1);
if (col == width) *(out+1) = *(in+1);
}
template<int width>
__global__ operator(int *in, int *out, int nrows, unsigned int lda)
{
// shared buffer holds three rows x (width+2) cols(threads)
__shared__ volatile int buffer [3][2+width];
int colid = threadIdx.x + blockIdx.x * blockDim.x;
int tid = threadIdx.x + 1;
int * rowpos = &in[colid], * outpos = &out[colid];
// load the first three rows (compiler will unroll loop)
for(int i=0; i<3; i++, rowpos+=lda) {
rowfetch<width>(rowpos, &buffer[i][tid], tid);
}
__syncthreads(); // shared memory loaded and all threads ready
int brow = 0; // brow is the next buffer row to load data onto
for(int i=0; i<nrows; i++, rowpos+=lda, outpos+=lda) {
// Do stencil calculations - use the value of brow to determine which
// stencil to use
result = ();
// write result to outpos
*outpos = result;
// Fetch another row
__syncthreads(); // Wait until all threads are done calculating
rowfetch<width>(rowpos, &buffer[brow][tid], tid);
brow = (brow < 2) ? (brow+1) : 0; // Increment or roll brow over
__syncthreads(); // Wait until all threads have updated the buffer
}
}
关于CUDA:在大型二维阵列上共享内存,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/5794617/
只是想知道是否有可能找出谁从 Windows 共享中读取了文件(最好使用 .NET,但 win32 native 可以)? 我想做的是创建类似 awstats 的东西对于 Windows 共享,这样我
是否可以列出 Intent.ACTION_SEND ?我的意思是我需要知道是否有人通过 action_send 在 Facebook 上分享或在 Twitter 上发推文。 最佳答案 也许你想要一个更
我正在使用 Google Apps 应用程序。实际上,我想在不使用密码的情况下访问另一个 ID。我使用了 OAuth,它运行良好。但我无法分享特定人的日历。我尝试了以下代码。 GoogleOAuthP
我怎样才能只创建模拟器...可能吗?我知道,设备需要分发证书。 最佳答案 您只需将应用程序目录从 iPhone 模拟器复制到另一个实例/操作系统版本,它就应该可以工作。 因此,如果您想分发 3.1.3
我想使用多阶段构建来避免每次构建应用程序时都下载我的 Java 项目所需的所有 Maven 依赖项。 我正在考虑在第一阶段解决 Maven 依赖项,然后在第二阶段构建应用程序,这将需要访问在前一阶段下
我正在寻找保护用户下载内容的初步想法。用户下载充满有趣资源的 zip 文件,这些资源被提取到本地文件系统中以供应用程序使用。我的目标是防止用户通过互联网将下载的资源共享给其他用户(假设他们获得了对文件
我想知道在具有移动和桌面版本的网站上共享身份验证、 session 管理等的最佳方法是什么。我们正在运行 Tomcat,并且更愿意将移动站点和桌面站点的应用程序保持在不同的节点上。 我看过类似的帖子,
我发现了这个单例的实现。我怎样才能创建指向它的指针或共享指针?` 为什么这不起作用?自动测试 = Singleton::Instance(); class Singleton { public: st
我有一个 heroku 项目,我想与其他人分享。作为the instructions describe ,我使用 virtualenv 来管理环境和依赖项。有没有办法在新机器上从 requiremen
Maven 将所有 jar 存储在本地存储库 ~/.m2/repository/ 下。用户多时占用空间大。 那么,是否可以由多个用户共享这个本地存储库,或许在不同的目录结构下? 最佳答案 简单的回答
为什么共享 worker 在重新加载页面时死了?应该是复活了我该如何解决这个问题? 重装前 重新加载后(在example.com上按F5) parent worker var port = new S
我正在开发多个小型应用程序,这些应用程序将共享通用和共享模块和 Assets 。 关于如何创建项目结构的部分在这里回答:https://stackoverflow.com/a/61254557/135
我在 RHEL 上安装了 jenkins (localhost:8080),我能够成功地构建代码 现在,我想设置主/从代理。 我的笔记本电脑将充当“Master Jenkins”,而我同事的笔记本电脑
我有这种方法可以根据我使用的 EXTRA_STREAM 共享文本文件或图片。我有这两个我可以选择 i.putExtra(Intent.EXTRA_STREAM, uri); i.putExtra(In
我正在使用 R 中的一个数据分析项目,我正在使用 R 中的敏感私有(private)数据进行一些逻辑和多级建模。我爱上了 。预订 包,我已经创建了一本关于我们的工作流程和分析管道的相当广泛的书。问题是
我正在构建的应用程序需要在 UITabBarController 框架内为多个 View (及其 subview )显示共享的自定义 UIToolbar。自定义工具栏的内容在所有 View 中都是相同
我有多个应用程序,我想共享相同的 eslint 配置: - project_root/ - app1/ - node_modules/ - eslint.rc
我有多个 Electron 应用程序。一个是主应用程序,其他几个功能应用程序。主应用程序上的按钮很少,这将导致功能应用程序打开。这里的问题是每个应用程序都有一个主进程,该进程导致要利用更多的CPU。是
我正在开发一个 Node.js 后端,它通过 websocket 与一些桌面客户端进行通信,而服务器端的通信是从 Web 前端发起的。一切正常,因为我将 SockJS Connection 实例存储在
我对托管多个网站的服务器上的多个用户帐户使用私有(private) SSH key 和无密码条目。 我为每个用户帐户使用相同的私钥。 (因为我很懒?或者那是“正确”的方式)。 我现在想授权该国不同地区
我是一名优秀的程序员,十分优秀!