- html - 出于某种原因,IE8 对我的 Sass 文件中继承的 html5 CSS 不友好?
- JMeter 在响应断言中使用 span 标签的问题
- html - 在 :hover and :active? 上具有不同效果的 CSS 动画
- html - 相对于居中的 html 内容固定的 CSS 重复背景?
大约两年前,我编写了一个内核,用于同时处理多个数值网格。出现了一些非常奇怪的行为,导致错误的结果。当利用内核内部的printf()语句查找bug时,该bug消失了。
由于截止日期的限制,我一直保持这种方式,尽管最近我发现这不是合适的编码风格。因此,我重新审视了内核,并将其归结为您在下面看到的内容。
__launch_bounds__(672, 2)
__global__ void heisenkernel(float *d_u, float *d_r, float *d_du, int radius,
int numNodesPerGrid, int numBlocksPerSM, int numGridsPerSM, int numGrids)
{
__syncthreads();
int id_sm = blockIdx.x / numBlocksPerSM; // (arbitrary) ID of Streaming Multiprocessor (SM) this thread works upon - (constant over lifetime of thread)
int id_blockOnSM = blockIdx.x % numBlocksPerSM; // Block number on this specific SM - (constant over lifetime of thread)
int id_r = id_blockOnSM * (blockDim.x - 2*radius) + threadIdx.x - radius; // Grid point number this thread is to work upon - (constant over lifetime of thread)
int id_grid = id_sm * numGridsPerSM; // Grid ID this thread is to work upon - (not constant over lifetime of thread)
while(id_grid < numGridsPerSM * (id_sm + 1)) // this loops over numGridsPerSM grids
{
__syncthreads();
int id_numInArray = id_grid * numNodesPerGrid + id_r; // Entry in array this thread is responsible for (read and possibly write) - (not constant over lifetime of thread)
float uchange = 0.0f;
//uchange = 1.0f; // if this line is uncommented, results will be computed correctly ("Solution 1")
float du = 0.0f;
if((threadIdx.x > radius-1) && (threadIdx.x < blockDim.x - radius) && (id_r < numNodesPerGrid) && (id_grid < numGrids))
{
if (id_r == 0) // FO-forward difference
du = (d_u[id_numInArray+1] - d_u[id_numInArray])/(d_r[id_numInArray+1] - d_r[id_numInArray]);
else if (id_r == numNodesPerGrid - 1) // FO-rearward difference
du = (d_u[id_numInArray] - d_u[id_numInArray-1])/(d_r[id_numInArray] - d_r[id_numInArray-1]);
else if (id_r == 1 || id_r == numNodesPerGrid - 2) //SO-central difference
du = (d_u[id_numInArray+1] - d_u[id_numInArray-1])/(d_r[id_numInArray+1] - d_r[id_numInArray-1]);
else if(id_r > 1 && id_r < numNodesPerGrid - 2)
du = d_fourpoint_constant * ((d_u[id_numInArray+1] - d_u[id_numInArray-1])/(d_r[id_numInArray+1] - d_r[id_numInArray-1])) + (1-d_fourpoint_constant) * ((d_u[id_numInArray+2] - d_u[id_numInArray-2])/(d_r[id_numInArray+2] - d_r[id_numInArray-2]));
else
du = 0;
}
__syncthreads();
if((threadIdx.x > radius-1 && threadIdx.x < blockDim.x - radius) && (id_r < numNodesPerGrid) && (id_grid < numGrids))
{
d_u[ id_numInArray] = d_u[id_numInArray] * uchange; // if this line is commented out, results will be computed correctly ("Solution 2")
d_du[ id_numInArray] = du;
}
__syncthreads();
++id_grid;
}
u_arr
,
du_arr
和
r_arr
(及其对应的设备数组
d_u
,
d_du
和
d_r
)中。
printf()
进行调试将改变
cuda-memcheck --tool racecheck
仅适用于共享内存争用条件,而当前的问题具有
d_u
的争用条件,即全局内存。
printf()
语句(通过一些黑魔法也必须利用设备和主机内存)和memcheck算法(cuda-memcheck和valgrind)会影响这一事实
d_u
时输入/输出值
uchange = 1.0
不会改变。
d_u
,另一个计算派生的
d_du
。尽管我不知道如何使用
-arch=sm_20
完成此操作,但最好仅执行一次内核调用而不是两次。尽管第二次内核调用的开销可以忽略不计,但使用
-arch=sm_35
可能可以使用动态并行性来实现。
#include <cuda.h>
#include <cuda_runtime.h>
#include <stdio.h>
const float r_sol = 6.955E8f;
__constant__ float d_fourpoint_constant = 0.2f;
__launch_bounds__(672, 2)
__global__ void heisenkernel(float *d_u, float *d_r, float *d_du, int radius,
int numNodesPerGrid, int numBlocksPerSM, int numGridsPerSM, int numGrids)
{
__syncthreads();
int id_sm = blockIdx.x / numBlocksPerSM; // (arbitrary) ID of Streaming Multiprocessor (SM) this thread works upon - (constant over lifetime of thread)
int id_blockOnSM = blockIdx.x % numBlocksPerSM; // Block number on this specific SM - (constant over lifetime of thread)
int id_r = id_blockOnSM * (blockDim.x - 2*radius) + threadIdx.x - radius; // Grid point number this thread is to work upon - (constant over lifetime of thread)
int id_grid = id_sm * numGridsPerSM; // Grid ID this thread is to work upon - (not constant over lifetime of thread)
while(id_grid < numGridsPerSM * (id_sm + 1)) // this loops over numGridsPerSM grids
{
__syncthreads();
int id_numInArray = id_grid * numNodesPerGrid + id_r; // Entry in array this thread is responsible for (read and possibly write) - (not constant over lifetime of thread)
float uchange = 0.0f;
//uchange = 1.0f; // if this line is uncommented, results will be computed correctly ("Solution 1")
float du = 0.0f;
if((threadIdx.x > radius-1) && (threadIdx.x < blockDim.x - radius) && (id_r < numNodesPerGrid) && (id_grid < numGrids))
{
if (id_r == 0) // FO-forward difference
du = (d_u[id_numInArray+1] - d_u[id_numInArray])/(d_r[id_numInArray+1] - d_r[id_numInArray]);
else if (id_r == numNodesPerGrid - 1) // FO-rearward difference
du = (d_u[id_numInArray] - d_u[id_numInArray-1])/(d_r[id_numInArray] - d_r[id_numInArray-1]);
else if (id_r == 1 || id_r == numNodesPerGrid - 2) //SO-central difference
du = (d_u[id_numInArray+1] - d_u[id_numInArray-1])/(d_r[id_numInArray+1] - d_r[id_numInArray-1]);
else if(id_r > 1 && id_r < numNodesPerGrid - 2)
du = d_fourpoint_constant * ((d_u[id_numInArray+1] - d_u[id_numInArray-1])/(d_r[id_numInArray+1] - d_r[id_numInArray-1])) + (1-d_fourpoint_constant) * ((d_u[id_numInArray+2] - d_u[id_numInArray-2])/(d_r[id_numInArray+2] - d_r[id_numInArray-2]));
else
du = 0;
}
__syncthreads();
if((threadIdx.x > radius-1 && threadIdx.x < blockDim.x - radius) && (id_r < numNodesPerGrid) && (id_grid < numGrids))
{
d_u[ id_numInArray] = d_u[id_numInArray] * uchange; // if this line is commented out, results will be computed correctly ("Solution 2")
d_du[ id_numInArray] = du;
}
__syncthreads();
++id_grid;
}
}
bool gridValuesEqual(float *matarray, uint id0, uint id1, const char *label, int numNodesPerGrid){
bool retval = true;
for(uint i=0; i<numNodesPerGrid; ++i)
if(matarray[id0 * numNodesPerGrid + i] != matarray[id1 * numNodesPerGrid + i])
{
printf("value %s at position %u of grid %u not equal that of grid %u: %E != %E, diff: %E\n",
label, i, id0, id1, matarray[id0 * numNodesPerGrid + i], matarray[id1 * numNodesPerGrid + i],
matarray[id0 * numNodesPerGrid + i] - matarray[id1 * numNodesPerGrid + i]);
retval = false;
}
return retval;
}
int main(int argc, const char* argv[])
{
float *d_u;
float *d_du;
float *d_r;
float *u_arr;
float *du_arr;
float *r_arr;
int numNodesPerGrid = 1300;
int numBlocksPerSM = 2;
int numGridsPerSM = 37;
int numSM = 7;
int TPB = 672;
int radius = 2;
int numGrids = 259;
int memsize_grid = sizeof(float) * numNodesPerGrid;
int numBlocksPerGrid = numNodesPerGrid / (TPB - 2 * radius) + (numNodesPerGrid%(TPB - 2 * radius) == 0 ? 0 : 1);
printf("---------------------------------------------------------------------------\n");
printf("--- Heisenbug Extermination Tracker ---------------------------------------\n");
printf("---------------------------------------------------------------------------\n\n");
cudaSetDevice(0);
cudaDeviceReset();
cudaMalloc((void **) &d_u, memsize_grid * numGrids);
cudaMalloc((void **) &d_du, memsize_grid * numGrids);
cudaMalloc((void **) &d_r, memsize_grid * numGrids);
u_arr = new float[numGrids * numNodesPerGrid];
du_arr = new float[numGrids * numNodesPerGrid];
r_arr = new float[numGrids * numNodesPerGrid];
for(uint k=0; k<numGrids; ++k)
for(uint i=0; i<numNodesPerGrid; ++i)
{
uint index = k * numNodesPerGrid + i;
if (i < 585)
r_arr[index] = i * (6000.0f);
else
{
if (i == 585)
r_arr[index] = r_arr[index - 1] + 8.576E-6f * r_sol;
else
r_arr[index] = r_arr[index - 1] + 1.02102f * ( r_arr[index - 1] - r_arr[index - 2] );
}
u_arr[index] = 1E-10f * (i+1);
du_arr[index] = 0.0f;
}
/*
printf("\n\nbefore kernel start\n\n");
for(uint k=0; k<numGrids; ++k)
printf("matrix->du_arr[k*paramH.numNodes + 668]:\t%E\n", du_arr[k*numNodesPerGrid + 668]);//*/
bool equal = true;
for(int k=1; k<numGrids; ++k)
{
equal &= gridValuesEqual(u_arr, 0, k, "u", numNodesPerGrid);
equal &= gridValuesEqual(du_arr, 0, k, "du", numNodesPerGrid);
equal &= gridValuesEqual(r_arr, 0, k, "r", numNodesPerGrid);
}
if(!equal)
printf("Input values are not identical for different grids!\n\n");
else
printf("All grids contain the same values at same grid points.!\n\n");
cudaMemcpy(d_u, u_arr, memsize_grid * numGrids, cudaMemcpyHostToDevice);
cudaMemcpy(d_du, du_arr, memsize_grid * numGrids, cudaMemcpyHostToDevice);
cudaMemcpy(d_r, r_arr, memsize_grid * numGrids, cudaMemcpyHostToDevice);
printf("Configuration:\n\n");
printf("numNodesPerGrid:\t%i\nnumBlocksPerSM:\t\t%i\nnumGridsPerSM:\t\t%i\n", numNodesPerGrid, numBlocksPerSM, numGridsPerSM);
printf("numSM:\t\t\t\t%i\nTPB:\t\t\t\t%i\nradius:\t\t\t\t%i\nnumGrids:\t\t\t%i\nmemsize_grid:\t\t%i\n", numSM, TPB, radius, numGrids, memsize_grid);
printf("numBlocksPerGrid:\t%i\n\n", numBlocksPerGrid);
printf("Kernel launch parameters:\n\n");
printf("moduleA2_3<<<%i, %i, %i>>>(...)\n\n", numBlocksPerSM * numSM, TPB, 0);
printf("Launching Kernel...\n\n");
heisenkernel<<<numBlocksPerSM * numSM, TPB, 0>>>(d_u, d_r, d_du, radius, numNodesPerGrid, numBlocksPerSM, numGridsPerSM, numGrids);
cudaDeviceSynchronize();
cudaMemcpy(u_arr, d_u, memsize_grid * numGrids, cudaMemcpyDeviceToHost);
cudaMemcpy(du_arr, d_du, memsize_grid * numGrids, cudaMemcpyDeviceToHost);
cudaMemcpy(r_arr, d_r, memsize_grid * numGrids, cudaMemcpyDeviceToHost);
/*
printf("\n\nafter kernel finished\n\n");
for(uint k=0; k<numGrids; ++k)
printf("matrix->du_arr[k*paramH.numNodes + 668]:\t%E\n", du_arr[k*numNodesPerGrid + 668]);//*/
equal = true;
for(int k=1; k<numGrids; ++k)
{
equal &= gridValuesEqual(u_arr, 0, k, "u", numNodesPerGrid);
equal &= gridValuesEqual(du_arr, 0, k, "du", numNodesPerGrid);
equal &= gridValuesEqual(r_arr, 0, k, "r", numNodesPerGrid);
}
if(!equal)
printf("Results are wrong!!\n");
else
printf("All went well!\n");
cudaFree(d_u);
cudaFree(d_du);
cudaFree(d_r);
delete [] u_arr;
delete [] du_arr;
delete [] r_arr;
return 0;
}
CUDA = 1
DEFINES =
ifeq ($(CUDA), 1)
DEFINES += -DCUDA
CUDAPATH = /usr/local/cuda-6.5
CUDAINCPATH = -I$(CUDAPATH)/include
CUDAARCH = -arch=sm_20
endif
CXX = g++
CXXFLAGS = -pipe -g -std=c++0x -fPIE -O0 $(DEFINES)
VALGRIND = valgrind
VALGRIND_FLAGS = -v --leak-check=yes --log-file=out.memcheck
CUDAMEMCHECK = cuda-memcheck
CUDAMC_FLAGS = --tool memcheck
RACECHECK = $(CUDAMEMCHECK)
RACECHECK_FLAGS = --tool racecheck
INCPATH = -I. $(CUDAINCPATH)
LINK = g++
LFLAGS = -O0
LIBS =
ifeq ($(CUDA), 1)
NVCC = $(CUDAPATH)/bin/nvcc
LIBS += -L$(CUDAPATH)/lib64/
LIBS += -lcuda -lcudart -lcudadevrt
NVCCFLAGS = -g -G -O0 --ptxas-options=-v
NVCCFLAGS += -lcuda -lcudart -lcudadevrt -lineinfo --machine 64 -x cu $(CUDAARCH) $(DEFINES)
endif
all:
$(NVCC) $(NVCCFLAGS) $(INCPATH) -c -o $(DST_DIR)heisenbug.o $(SRC_DIR)heisenbug.cu
$(LINK) $(LFLAGS) -o heisenbug heisenbug.o $(LIBS)
clean:
rm heisenbug.o
rm heisenbug
memrace: all
./heisenbug > out
$(VALGRIND) $(VALGRIND_FLAGS) ./heisenbug > out.memcheck.log
$(CUDAMEMCHECK) $(CUDAMC_FLAGS) ./heisenbug > out.cudamemcheck
$(RACECHECK) $(RACECHECK_FLAGS) ./heisenbug > out.racecheck
最佳答案
请注意,在您撰写本文的全部内容中,我没有看到明确提出的问题,因此我在答复:
我期待着学习这里发生的事情。
您在d_u
上存在比赛条件。
根据您自己的陈述:
•为了使块彼此独立,在网格上引入了一个小的重叠部分(尽管只有一个线程正在写入,但每个网格的网格点666、667、668、669是由来自不同块的两个线程读取的)它们,就是发生问题的地方就是这种重叠)
此外,如果您注释掉对d_u
的写入,则根据代码中的语句,问题将消失。
CUDA线程块可以按任何顺序执行。您至少有2个不同的块正在从网格点666、667、668、669读取。根据实际发生的情况,结果将有所不同:
两个块都在发生任何写操作之前读取值。
一个块读取该值,然后进行写操作,然后另一个块读取该值。
如果一个块正在读取可被另一块写入的值,则这些块不是彼此独立的(与您的陈述相反)。在这种情况下,块执行的顺序将确定结果,而CUDA未指定块执行的顺序。
请注意,带有cuda-memcheck
选项only captures race conditions related to -tool racecheck
memory usage的__shared__
。您发布的内核不使用__shared__
内存,因此我不希望cuda-memcheck
报告任何内容。
为了收集数据,cuda-memcheck
确实会影响块执行的顺序,因此它会影响行为也就不足为奇了。
内核中的printf
表示代价高昂的函数调用,它写入全局内存缓冲区。因此,它也会影响执行行为/模式。而且,如果您打印出大量数据,超出了输出的缓冲区行,则在缓冲区溢出的情况下,这样做的代价是非常昂贵的(就执行时间而言)。
顺便说一句,据我所知,Linux Mint是not a supported distro for CUDA。但是,我认为这与您的问题无关;我可以在受支持的配置上重现该行为。
关于cuda - CUDA内核中的Heisenbug,全局内存访问,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/27219649/
我的应用程序中有一个 settings.php 页面,它使用 $GLOBALS 来存储网络应用程序中使用的配置。 例如,他是我使用的一个示例设置变量: $GLOBALS["new_login_page
我正在尝试编译我们在 OS 类上获得的简单操作系统代码。它在 Ubuntu 下运行良好,但我想在 OS X 上编译它。我得到的错误是: [compiling] arch/i386/arch/start
我知道distcp无法使用通配符。 但是,我将需要在更改的目录上安排distcp。 (即,仅在星期一等“星期五”目录中复制数据),还从指定目录下的所有项目中复制数据。 是否有某种设计模式可用于编写此类
是否可以在config.groovy中全局定义资源格式(json,xml)的优先级,而不是在每个Resource上指定?例如,不要在@Resource Annotation的参数中指定它,例如: @R
是否有一些简单的方法来获取大对象图的所有关联,而不必“左连接获取”所有关联?我不能只告诉 Hibernate 默认获取 eager 关联吗? 最佳答案 即使有可能有一个全局 lazy=false(谷歌
我正在尝试实现一个全局加载对话框...我想调用一些静态函数来显示对话框和一些静态函数来关闭它。与此同时,我正在主线程或子线程中做一些工作...... 我尝试了以下操作,但对话框没有更新...最后一次,
当我偶然发现 this question 时,我正在阅读更改占位符文本。 无论如何,我回去学习了占位符。一个 SO 的回答大致如下: Be careful when designing your pl
例如,如果我有这样的文字: "hello800 more text 1234 and 567" 它应该匹配 1234 和 567,而不是 800(因为它遵循 hello 的 o,这不是一个数字)。 这
我一直在尝试寻找一种无需使用 SMS 验证系统即可验证电话号码(Android 和 iPhone)的方法。原因纯粹是围绕成本。我想要一个免费的解决方案。 我可以安全地假设 Android 操作系统会向
解决此类问题的规范 C++ 设计模式是什么? 我有一些共享多个类的多线程服务器。我需要为大多数类提供各种运行时参数(例如服务器名称、日志记录级别)。 在下面的伪 C++ 代码中,我使用了一个日志记录类
这个问题在这里已经有了答案: Using global variables in a function (25 个答案) 关闭 9 年前。 我是 python 的新手,所以可能有一个简单的答案,但我
这个问题在这里已经有了答案: 关闭 10 年前。 Possible Duplicate: Does C++ call destructors for global and class static
我正在尝试使用 Objective-C 中的 ArrayList 的等价物。我知道我必须使用 NSMutableArray。我想要一个字符串列表 (NSString)。关键是我的列表应该可以从我类(c
今天刚开始学习 Android 开发,我找不到任何关于如何定义 Helper 类或将全局加载的函数集合的信息,我会能够在我创建的任何 Activity 中使用它们。 我的计划是创建(至少目前)2 个几
为什么这段代码有效: var = 0 def func(num): print num var = 1 if num != 0: func(num-1) fun
$GLOBALS["items"] = array('one', 'two', 'three', 'four', 'five' ,'six', 'seven'); $alter = &$GLOBALS
我想知道如何实现一个可以在任何地方使用您自己的设置的全局记录器: 我目前有一个自定义记录器类: class customLogger(logging.Logger): ... 该类位于一个单独的
我需要使用 React 测试库和 Jest 在我的测试中模拟不同的窗口大小。 目前我必须在每个测试文件中包含这个beforeAll: import matchMediaPolyfill from 'm
每次我遇到单例模式或任何静态类(即(几乎)只有静态成员的类)的实现时,我想知道这是否实际上不是一种黑客行为,因此只是为了设计而严重滥用类和实例的原则单个对象,而不是设计类和创建单个实例。对我来说,看起
这个问题在这里已经有了答案: Help understanding global flag in perl (2 个回答) 7年前关闭。 my $test = "There was once an\n
我是一名优秀的程序员,十分优秀!