- html - 出于某种原因,IE8 对我的 Sass 文件中继承的 html5 CSS 不友好?
- JMeter 在响应断言中使用 span 标签的问题
- html - 在 :hover and :active? 上具有不同效果的 CSS 动画
- html - 相对于居中的 html 内容固定的 CSS 重复背景?
编辑:这个问题是原版的重做版本,所以前几个回答可能不再相关。
我很好奇强制无内联的设备函数调用对设备函数内的同步有什么影响。我有一个简单的测试内核来说明相关行为。
内核获取一个缓冲区并将其传递给设备函数,以及一个共享缓冲区和一个指示变量,该变量将单个线程标识为“老板”线程。 device 函数有不同的代码:boss 线程首先花时间在共享缓冲区上做一些琐碎的操作,然后写入全局缓冲区。同步调用后,所有线程都写入全局缓冲区。内核调用后,主机打印全局缓冲区的内容。这是代码:
CUDA 代码:
test_main.cu
#include<cutil_inline.h>
#include "test_kernel.cu"
int main()
{
int scratchBufferLength = 100;
int *scratchBuffer;
int *d_scratchBuffer;
int b = 1;
int t = 64;
// copy scratch buffer to device
scratchBuffer = (int *)calloc(scratchBufferLength,sizeof(int));
cutilSafeCall( cudaMalloc(&d_scratchBuffer,
sizeof(int) * scratchBufferLength) );
cutilSafeCall( cudaMemcpy(d_scratchBuffer, scratchBuffer,
sizeof(int)*scratchBufferLength, cudaMemcpyHostToDevice) );
// kernel call
testKernel<<<b, t>>>(d_scratchBuffer);
cudaThreadSynchronize();
// copy data back to host
cutilSafeCall( cudaMemcpy(scratchBuffer, d_scratchBuffer,
sizeof(int) * scratchBufferLength, cudaMemcpyDeviceToHost) );
// print results
printf("Scratch buffer contents: \t");
for(int i=0; i < scratchBufferLength; ++i)
{
if(i % 25 == 0)
printf("\n");
printf("%d ", scratchBuffer[i]);
}
printf("\n");
//cleanup
cudaFree(d_scratchBuffer);
free(scratchBuffer);
return 0;
}
#ifndef __TEST_KERNEL_CU
#define __TEST_KERNEL_CU
#define IS_BOSS() (threadIdx.x == blockDim.x - 1)
__device__
__noinline__
void testFunc(int *sA, int *scratchBuffer, bool isBoss) {
if(isBoss) { // produces unexpected output-- "broken" code
//if(IS_BOSS()) { // produces expected output-- "working" code
for (int c = 0; c < 10000; c++) {
sA[0] = 1;
}
}
if(isBoss) {
scratchBuffer[0] = 1;
}
__syncthreads();
scratchBuffer[threadIdx.x ] = threadIdx.x;
return;
}
__global__
void testKernel(int *scratchBuffer)
{
__shared__ int sA[4];
bool isBoss = IS_BOSS();
testFunc(sA, scratchBuffer, isBoss);
return;
}
#endif
0 1 2 3 ... blockDim.x-1
1 1 2 3 ... blockDim.x-1
if(isBoss) {
if(IS_BOSS()) {
code for sm_20
Function : _Z10testKernelPi
/*0000*/ /*0x00005de428004404*/ MOV R1, c [0x1] [0x100];
/*0008*/ /*0x20009de428004000*/ MOV R2, c [0x0] [0x8];
/*0010*/ /*0x84001c042c000000*/ S2R R0, SR_Tid_X;
/*0018*/ /*0xfc015de428000000*/ MOV R5, RZ;
/*0020*/ /*0x00011de428004000*/ MOV R4, c [0x0] [0x0];
/*0028*/ /*0xfc209c034800ffff*/ IADD R2, R2, 0xfffff;
/*0030*/ /*0x9001dde428004000*/ MOV R7, c [0x0] [0x24];
/*0038*/ /*0x80019de428004000*/ MOV R6, c [0x0] [0x20];
/*0040*/ /*0x08001c03110e0000*/ ISET.EQ.U32.AND R0, R0, R2, pt;
/*0048*/ /*0x01221f841c000000*/ I2I.S32.S32 R8, -R0;
/*0050*/ /*0x2001000750000000*/ CAL 0x60;
/*0058*/ /*0x00001de780000000*/ EXIT;
/*0060*/ /*0x20201e841c000000*/ I2I.S32.S8 R0, R8;
/*0068*/ /*0xfc01dc231a8e0000*/ ISETP.NE.AND P0, pt, R0, RZ, pt;
/*0070*/ /*0xc00021e740000000*/ @!P0 BRA 0xa8;
/*0078*/ /*0xfc001de428000000*/ MOV R0, RZ;
/*0080*/ /*0x04001c034800c000*/ IADD R0, R0, 0x1;
/*0088*/ /*0x04009de218000000*/ MOV32I R2, 0x1;
/*0090*/ /*0x4003dc231a8ec09c*/ ISETP.NE.AND P1, pt, R0, 0x2710, pt;
/*0098*/ /*0x00409c8594000000*/ ST.E [R4], R2;
/*00a0*/ /*0x600005e74003ffff*/ @P1 BRA 0x80;
/*00a8*/ /*0x040001e218000000*/ @P0 MOV32I R0, 0x1;
/*00b0*/ /*0x0060008594000000*/ @P0 ST.E [R6], R0;
/*00b8*/ /*0xffffdc0450ee0000*/ BAR.RED.POPC RZ, RZ;
/*00c0*/ /*0x84001c042c000000*/ S2R R0, SR_Tid_X;
/*00c8*/ /*0x10011c03200dc000*/ IMAD.U32.U32 R4.CC, R0, 0x4, R6;
/*00d0*/ /*0x10009c435000c000*/ IMUL.U32.U32.HI R2, R0, 0x4;
/*00d8*/ /*0x08715c4348000000*/ IADD.X R5, R7, R2;
/*00e0*/ /*0x00401c8594000000*/ ST.E [R4], R0;
/*00e8*/ /*0x00001de790000000*/ RET;
.................................
code for sm_20
Function : _Z10testKernelPi
/*0000*/ /*0x00005de428004404*/ MOV R1, c [0x1] [0x100];
/*0008*/ /*0x20009de428004000*/ MOV R2, c [0x0] [0x8];
/*0010*/ /*0x84001c042c000000*/ S2R R0, SR_Tid_X;
/*0018*/ /*0xfc015de428000000*/ MOV R5, RZ;
/*0020*/ /*0x00011de428004000*/ MOV R4, c [0x0] [0x0];
/*0028*/ /*0xfc209c034800ffff*/ IADD R2, R2, 0xfffff;
/*0030*/ /*0x9001dde428004000*/ MOV R7, c [0x0] [0x24];
/*0038*/ /*0x80019de428004000*/ MOV R6, c [0x0] [0x20];
/*0040*/ /*0x08001c03110e0000*/ ISET.EQ.U32.AND R0, R0, R2, pt;
/*0048*/ /*0x01221f841c000000*/ I2I.S32.S32 R8, -R0;
/*0050*/ /*0x2001000750000000*/ CAL 0x60;
/*0058*/ /*0x00001de780000000*/ EXIT;
/*0060*/ /*0x20009de428004000*/ MOV R2, c [0x0] [0x8];
/*0068*/ /*0x8400dc042c000000*/ S2R R3, SR_Tid_X;
/*0070*/ /*0x20201e841c000000*/ I2I.S32.S8 R0, R8;
/*0078*/ /*0x4000000760000001*/ SSY 0xd0;
/*0080*/ /*0xfc209c034800ffff*/ IADD R2, R2, 0xfffff;
/*0088*/ /*0x0831dc031a8e0000*/ ISETP.NE.U32.AND P0, pt, R3, R2, pt;
/*0090*/ /*0xc00001e740000000*/ @P0 BRA 0xc8;
/*0098*/ /*0xfc009de428000000*/ MOV R2, RZ;
/*00a0*/ /*0x04209c034800c000*/ IADD R2, R2, 0x1;
/*00a8*/ /*0x04021de218000000*/ MOV32I R8, 0x1;
/*00b0*/ /*0x4021dc231a8ec09c*/ ISETP.NE.AND P0, pt, R2, 0x2710, pt;
/*00b8*/ /*0x00421c8594000000*/ ST.E [R4], R8;
/*00c0*/ /*0x600001e74003ffff*/ @P0 BRA 0xa0;
/*00c8*/ /*0xfc01dc33190e0000*/ ISETP.EQ.AND.S P0, pt, R0, RZ, pt;
/*00d0*/ /*0x040021e218000000*/ @!P0 MOV32I R0, 0x1;
/*00d8*/ /*0x0060208594000000*/ @!P0 ST.E [R6], R0;
/*00e0*/ /*0xffffdc0450ee0000*/ BAR.RED.POPC RZ, RZ;
/*00e8*/ /*0x10311c03200dc000*/ IMAD.U32.U32 R4.CC, R3, 0x4, R6;
/*00f0*/ /*0x10309c435000c000*/ IMUL.U32.U32.HI R2, R3, 0x4;
/*00f8*/ /*0x84001c042c000000*/ S2R R0, SR_Tid_X;
/*0100*/ /*0x08715c4348000000*/ IADD.X R5, R7, R2;
/*0108*/ /*0x00401c8594000000*/ ST.E [R4], R0;
/*0110*/ /*0x00001de790000000*/ RET;
.................................
最佳答案
这似乎只是 CUDA 4.1/4.2 中修复的编译器错误。不会在 CUDA 4.2 上为提问者重现。
关于CUDA 设备堆栈和同步; SSY指令,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/12377398/
出于好奇 - 我知道有 LAMP - Linux、Apache、MySQL 和 PHP。但是还有哪些其他 Web 堆栈替代方案的缩写呢?像 LAMR - Linux、Apache、MySQL Ruby
我有以下代码。 var stackMapIn = []; var stackMapOut = []; var stackBack = []; stackMapOut.push("m1"); $scop
我遇到了导致我的堆栈无法恢复的情况,我别无选择,只能将其删除。使用完全相同的模板,我继续创建了另一个同名的堆栈。 The following resource(s) failed to create:
这是我第一次查看 Node 堆栈,自从我学习使用 Ruby on Rails 进行 Web 开发以来,我对一些基本的东西有点困惑。我了解 Rails 目录是什么样的。 demo/ ..../app .
本文实例讲述了C语言使用深度优先搜索算法解决迷宫问题。分享给大家供大家参考,具体如下: 深度优先搜索 伪代码 (Pseudocode)如下: ?
我正在按照指南 here ,它告诉我: The stack setup will download the compiler if necessary in an isolatedlocation (
同时 trying to debug a different question ,我安装了一个似乎与我安装的其他一些软件包冲突的软件包。 我跑了 $ stack install regex-pcre-
我花了几个小时创建了一个方法,该方法将从堆栈 s1 中获取 null 元素,并将它们放入 s2 中。然后该类应该打印堆栈。方法如下 import net.datastructures.ArraySta
我有一个类Floor,它有一个Stack block ,但我不知道如何初始化它。我曾尝试过这样的: public class Floor { private Stack stack;
我知道这个问题已经问过很多次了,但搜索一个小时后我仍然遇到问题。 我想使用一个 lifo 堆栈,它可以存储最大数量的元素。达到最大数量后,首先删除该元素并将其替换为新元素,这样在第一次弹出时我可以获取
我需要编写一个方法,压缩以执行以下操作; 目标compress方法是从栈s1中移除所有null元素。剩余(非空)元素应按其初始顺序保留在 s1 上。辅助堆栈 s2 应用作s1 中元素的临时存储。在该方
我正在尝试验证以下代码发生的顺序。 function square(n) { return n * n; } setTimeout(function(){ console.log("H
我需要一个字符数组,其中包含基于特定文件夹中文件数量的动态数量的字符数组。我能够通过初始化 char (*FullPathNames)[MAX_FILENAME_AND_PATHNAME_LENGTH
我正在编写一些日志逻辑并想要进行一些缩进。了解是否存在任何函数调用或某个函数是否已完成的最简单方法是查看堆栈/帧的当前地址。让我们假设堆栈颠倒增长。然后,如果 log() 调用中的堆栈地址小于前一次调
所以内存分段在x86-64中被放弃了,但是当我们使用汇编时,我们可以在代码中指定.code和.data段/段,并且还有堆栈指针寄存器。 还有堆栈段、数据段和代码段寄存器。 代码/数据/堆栈的划分是如何
void main() { int x = 5; // stack-allocated Console.WriteLine(x); } 我知道 x 是堆栈分配的。但是关于 x 的堆栈中
这是我关于 SO 的第一个问题。这可能是一个愚蠢的问题,但到目前为止我还没弄明白。 考虑下面的程序 Reader.java: public class Reader { public
java中有没有一种快速的方法来获取嵌套/递归级别? 我正在编写一个函数来创建组及其成员的列表。成员也可以是团体。我们最终可能会得到一组循环的组/成员。 我想在某个任意级别停止。 我知道我可以将变量保
考虑以下代码: struct A{...}; A a[100]; A* pa = new A[100]; delete[] pa; a/pa 元素的销毁顺序是由标准定义的还是实现定义的(对于第二种情况
我在下面有一些代码。此代码是一个基本的压入/弹出堆栈类,我将其创建为模板以允许某人压入/弹出堆栈。我有一个家庭作业,我现在要做的是创建一个具有多个值的堆栈。 所以我希望能够创建一个基本上可以发送三个整
我是一名优秀的程序员,十分优秀!