- html - 出于某种原因,IE8 对我的 Sass 文件中继承的 html5 CSS 不友好?
- JMeter 在响应断言中使用 span 标签的问题
- html - 在 :hover and :active? 上具有不同效果的 CSS 动画
- html - 相对于居中的 html 内容固定的 CSS 重复背景?
我正在学习有关并行编程的 Udacity 类(class)(作业 3),但无法弄清楚为什么我无法使用并行缩减获得数组中的最大值(Udacity 论坛尚未提供解决方案)。我很确定我已经正确设置了数组并且算法是正确的。我怀疑我的内存管理有问题(越界访问、不正确的数组大小、来回复制)。请帮忙!我在 Udacity 环境中运行它,而不是在本地。下面是我目前正在使用的代码。出于某种原因,当我将 fmaxf
更改为 fminf
时,它确实找到了最小值。
#include "reference_calc.cpp"
#include "utils.h"
#include "math.h"
#include <stdio.h>
#include <cmath>
__global__ void reduce_max_kernel(float *d_out, const float *d_logLum, int size) {
// Reduce log Lum with Max Operator
int myId = threadIdx.x + blockDim.x * blockIdx.x;
int tid = threadIdx.x;
extern __shared__ float temp[];
if (myId < size) {
temp[tid] = d_logLum[myId];
}
else {
temp[tid] = d_logLum[tid];
}
for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) {
if (tid < s) {
if (myId < size) {
temp[tid] = fmaxf(d_logLum[myId + s], d_logLum[myId]);
} else {
temp[tid] = d_logLum[tid];
}
}
__syncthreads();
}
if (tid == 0) {
d_out[blockIdx.x] = temp[0];
}
}
__global__ void reduce_max_kernel2(float *d_out, float *d_in) {
// Reduce log Lum with Max Operator
int myId = threadIdx.x + blockDim.x * blockIdx.x;
int tid = threadIdx.x;
for (unsigned int s = blockDim.x >> 1; s > 0; s >>= 1) {
if (tid < s) {
d_in[myId] = fmaxf(d_in[myId + s], d_in[myId]);
}
__syncthreads();
}
if (tid == 0) {
d_out[0] = d_in[0];
}
}
void your_histogram_and_prefixsum(const float* const d_logLuminance,
unsigned int* const d_cdf,
float &min_logLum,
float &max_logLum,
const size_t numRows,
const size_t numCols,
const size_t numBins)
{
//TODO
/*Here are the steps you need to implement
1) find the minimum and maximum value in the input logLuminance channel
store in min_logLum and max_logLum
2) subtract them to find the range
3) generate a histogram of all the values in the logLuminance channel using
the formula: bin = (lum[i] - lumMin) / lumRange * numBins
4) Perform an exclusive scan (prefix sum) on the histogram to get
the cumulative distribution of luminance values (this should go in the
incoming d_cdf pointer which already has been allocated for you) */
//int size = 1 << 18;
int points = numRows * numCols;
int logPoints = ceil(log(points)/log(2));
int sizePow = logPoints;
int size = pow(2, sizePow);
int numThreads = 1024;
int numBlocks = size / numThreads;
float *d_out;
float *d_max_out;
checkCudaErrors(cudaMalloc((void **) &d_out, numBlocks * sizeof(float)));
checkCudaErrors(cudaMalloc((void **) &d_max_out, sizeof(float)));
cudaDeviceSynchronize();
reduce_max_kernel<<<numBlocks, numThreads, sizeof(float)*numThreads>>>(d_out, d_logLuminance, points);
cudaDeviceSynchronize();
reduce_max_kernel2<<<1, numBlocks>>>(d_max_out, d_out);
float h_out_max;
checkCudaErrors(cudaMemcpy(&h_out_max, d_max_out, sizeof(float), cudaMemcpyDeviceToHost));
printf("%f\n", h_out_max);
checkCudaErrors(cudaFree(d_max_out));
checkCudaErrors(cudaFree(d_out));
}
最佳答案
您正在尝试重现 CUDA SDK 缩减示例的 reduce2
缩减内核。 Robert Crovella 已经发现了您在代码中犯的两个错误。除了它们,我认为你还错误地初始化了共享内存。
请在下面找到围绕您的尝试构建的完整工作示例。我留下了错误的方法说明。
#include <thrust\device_vector.h>
#define BLOCKSIZE 256
/********************/
/* CUDA ERROR CHECK */
/********************/
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
if (code != cudaSuccess)
{
fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
if (abort) { getchar(); exit(code); }
}
}
/*******************************************************/
/* CALCULATING THE NEXT POWER OF 2 OF A CERTAIN NUMBER */
/*******************************************************/
unsigned int nextPow2(unsigned int x)
{
--x;
x |= x >> 1;
x |= x >> 2;
x |= x >> 4;
x |= x >> 8;
x |= x >> 16;
return ++x;
}
__global__ void reduce_max_kernel(float *d_out, const float *d_logLum, int size) {
int tid = threadIdx.x; // Local thread index
int myId = blockIdx.x * blockDim.x + threadIdx.x; // Global thread index
extern __shared__ float temp[];
// --- Loading data to shared memory. All the threads contribute to loading the data to shared memory.
temp[tid] = (myId < size) ? d_logLum[myId] : -FLT_MAX;
// --- Your solution
// if (myId < size) { temp[tid] = d_logLum[myId]; } else { temp[tid] = d_logLum[tid]; }
// --- Before going further, we have to make sure that all the shared memory loads have been completed
__syncthreads();
// --- Reduction in shared memory. Only half of the threads contribute to reduction.
for (unsigned int s=blockDim.x/2; s>0; s>>=1)
{
if (tid < s) { temp[tid] = fmaxf(temp[tid], temp[tid + s]); }
// --- At the end of each iteration loop, we have to make sure that all memory operations have been completed
__syncthreads();
}
// --- Your solution
//for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) {
// if (tid < s) { if (myId < size) { temp[tid] = fmaxf(d_logLum[myId + s], d_logLum[myId]); } else { temp[tid] = d_logLum[tid]; } }
// __syncthreads();
//}
if (tid == 0) {
d_out[blockIdx.x] = temp[0];
}
}
/********/
/* MAIN */
/********/
int main()
{
const int N = 10;
thrust::device_vector<float> d_vec(N,3.f); d_vec[4] = 4.f;
int NumThreads = (N < BLOCKSIZE) ? nextPow2(N) : BLOCKSIZE;
int NumBlocks = (N + NumThreads - 1) / NumThreads;
// when there is only one warp per block, we need to allocate two warps
// worth of shared memory so that we don't index shared memory out of bounds
int smemSize = (NumThreads <= 32) ? 2 * NumThreads * sizeof(int) : NumThreads * sizeof(int);
// --- reduce2
thrust::device_vector<float> d_vec_block(NumBlocks);
reduce_max_kernel<<<NumBlocks, NumThreads, smemSize>>>(thrust::raw_pointer_cast(d_vec_block.data()), thrust::raw_pointer_cast(d_vec.data()), N);
// --- The last part of the reduction, which would be expensive to perform on the device, is executed on the host
thrust::host_vector<float> h_vec_block(d_vec_block);
float result_reduce0 = -FLT_MAX;
for (int i=0; i<NumBlocks; i++) result_reduce0 = fmax(h_vec_block[i], result_reduce0);
printf("Result = %f\n",result_reduce0);
}
关于CUDA减少以找到数组的最大值,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/24475872/
我是 Bison 解析的新手,我无法理解它是如何工作的。我有以下语法,其中我保留了最低限度的语法来突出问题。 %left '~' %left '+' %token T_VARIABLE %% star
我链接了 2 个映射器和 1 个缩减器。是否可以将中间输出(链中每个映射器的 o/p)写入 HDFS?我尝试为每个设置 OutputPath,但它似乎不起作用。现在,我不确定是否可以完成。有什么建议吗
我正在编写一些代码来管理自定义磁盘文件结构并将其同步到未连接的系统。我的要求之一是能够在实际生成同步内容之前估计同步的大小。作为一个简单的解决方案,我整理了一个包含完整路径文件名的 map ,作为高效
我来自一个 SQL 世界,其中查找由多个对象属性(published = TRUE 或 user_id = X)完成,并且有 任何地方都没有加入 (因为 1:1 缓存层)。文档数据库似乎很适合我的数据
在 R 中,我有一个整数向量。从这个向量中,我想随机减少每个整数元素的值,以获得向量的总和,即初始总和的百分比。 在这个例子中,我想将向量“x”减少到向量“y”,其中每个元素都被随机减少以获得等于初始
我发现自己遇到过几次我有一个 reducer /组合 fn 的情况,如下所示: def combiner(a: String, b: String): Either[String, String]
Ubuntu 12.04 nginx 1.2.4 avconv版本 avconv version 0.8.10-4:0.8.10-0ubuntu0.12.04.1, Copyright (c) 200
我是 R 编程语言的新手。我有一个包含 2 列(ID 和 Num)的数据集,如下所示: ID Num 3 8 3 12 4 15 4 18 4
我正在使用高阶函数将函数应用于向量中的每个元素并将结果作为标量值返回。 假设我有: v = c(0, 1, 2, 3, 4, 5, 6, 7, 8) 我想计算以左边 5 个整数为中心的所有这些整数的总
关闭。这个问题需要debugging details .它目前不接受答案。 编辑问题以包含 desired behavior, a specific problem or error, and th
这个问题在这里已经有了答案: How to write the dataframes in a list to a single csv file (2 个回答) 5年前关闭。 我正在尝试使用 Red
刚开始学习CUDA编程,对归约有些迷茫。 我知道与共享内存相比,全局内存有很多访问延迟,但我可以使用全局内存来(至少)模拟类似于共享内存的行为吗? 例如,我想对长度恰好为 BLOCK_SIZE * T
我经常使用OptiPNG或pngcrush减小PNG图像的文件大小。 我希望能够从.NET应用程序中以编程方式执行此类操作。我正在动态生成要发送到移动设备的PNG,因此我想减小文件大小。 图像质量很重
减少和减少让您在序列上累积状态。 序列中的每个元素都会修改累积的状态,直到 到达序列的末尾。 在无限列表上调用reduce 或reductions 有什么含义? (def c (cycle [0]))
这与R: use the newly generated data in the previous row有关 我意识到我面临的实际问题比我在上面的线程中给出的示例要复杂一些 - 似乎我必须将 3 个
有什么办法可以减少.ttf字体的大小?即如果我们要删除一些我们不使用的glyps。 最佳答案 使用Google Web Fonts,您可以限制字符集,例如: //fonts.googleapis.co
我需要在iOS中制作一个应用程序,在她的工作过程中发出类似“哔”的声音。 我已经使用MPMusicPlayerController实现了与背景ipod的交互。 问题: 由于来自ipod的音乐音量很大,
我有一个嵌套 map m,如下所示: m = Map("电子邮件"-> "a@b.com", "背景"-> Map("语言"-> "英语")) 我有一个数组arr = Array("backgroun
有什么原因为什么不应该转发map / reduce函数中收到的可写内容? 我的意思是-每个map / reduce函数都有一个可写的键/值,并可能发出一个键/值对。如果我想执行一些过滤,我应该只发出接
假设我有一个数据列表 val data = listOf("F 1", "D 2", "U 1", "D 3", "F 10") 我想执行每个元素的给定逻辑。 我必须在外部添加 var acc2 =
我是一名优秀的程序员,十分优秀!