- iOS/Objective-C 元类和类别
- objective-c - -1001 错误,当 NSURLSession 通过 httpproxy 和/etc/hosts
- java - 使用网络类获取 url 地址
- ios - 推送通知中不播放声音
更新:标题具有误导性。最初我可以通过展开下面代码中的 block
循环来使错误消失。现在,即使是简单的代码更改也可以让它消失。 (参见下面的代码示例)。
背景:
12x12 矩阵的 Cholesky 分解的 CUDA C 内核实现导致相当大的 CUDA 内核(280 行代码,大量循环)。
我用简化的设置重现了错误(下面的代码)。 NVCC (CUDA 4.2) 调用
nvcc -arch sm_20 -o main main.cu
在 Linux 上的 Fermi 架构上执行:
kernel call: unspecified launch failure
内核主体包含一个条件预处理器 block #if 1
、#else
、#endif
。我插入它是为了在工作版本和非工作版本之间轻松切换。编译第一个替代结果会导致 未指定的启动失败
。而第二种选择运行良好。
棘手的部分是,实际执行的代码在任何一种情况下都应该相同。 (hasOrderedRep 是真的!!)
即使 #if 1
语句保持不变,仍然可以使错误消失。因此必须展开 block
循环。这就是标题的来源。
#include "cuda.h"
#include "stdio.h"
#include <iostream>
#include <string>
using namespace std;
/////////////////////////////////////
//
// First some basic types I need
// Implementation of a templated
// scalar and complex number
template<class T> class RScalar
{
public:
__device__ RScalar() {}
__device__ ~RScalar() {}
template<class T1> __device__
RScalar(const RScalar<T1>& rhs) : F(rhs.elem()) {}
template<class T1> __device__
RScalar(const T1& rhs) : F(rhs) {}
template<class T1> __device__ inline
RScalar& operator=(const RScalar<T1>& rhs) {
elem() = rhs.elem();
return *this;
}
public:
__device__ T& elem() {return F;}
__device__ const T& elem() const {return F;}
private:
T F;
};
template<class T> class RComplex
{
public:
__device__ RComplex() {}
__device__ ~RComplex() {}
template<class T1, class T2> __device__
RComplex(const RScalar<T1>& _re, const RScalar<T2>& _im):
re(_re.elem()), im(_im.elem()) {}
template<class T1, class T2> __device__
RComplex(const T1& _re, const T2& _im): re(_re), im(_im) {}
template<class T1> __device__
RComplex(const T1& _re): re(_re), im() {}
template<class T1>
__device__ inline
RComplex& operator*=(const RScalar<T1>& rhs)
{
real() *= rhs.elem();
imag() *= rhs.elem();
return *this;
}
template<class T1>
__device__ inline
RComplex& operator-=(const RComplex<T1>& rhs)
{
real() -= rhs.real();
imag() -= rhs.imag();
return *this;
}
template<class T1> __device__ inline
RComplex& operator/=(const RComplex<T1>& rhs)
{
RComplex<T> d;
d = *this / rhs;
real() = d.real();
imag() = d.imag();
return *this;
}
public:
__device__ T& real() {return re;}
__device__ const T& real() const {return re;}
__device__ T& imag() {return im;}
__device__ const T& imag() const {return im;}
private:
T re;
T im;
};
template<class T> __device__ RComplex<T>
operator*(const RComplex<T>& __restrict__ l,
const RComplex<T>& __restrict__ r)
{
return RComplex<T>(l.real()*r.real() - l.imag()*r.imag(),
l.real()*r.imag() + l.imag()*r.real());
}
template<class T> __device__ RComplex<T>
operator/(const RComplex<T>& l, const RComplex<T>& r)
{
T tmp = T(1.0) / (r.real()*r.real() + r.imag()*r.imag());
return RComplex<T>((l.real()*r.real() + l.imag()*r.imag()) * tmp,
(l.imag()*r.real() - l.real()*r.imag()) * tmp);
}
template<class T> __device__ RComplex<T>
operator*(const RComplex<T>& l, const RScalar<T>& r)
{
return RComplex<T>(l.real()*r.elem(),
l.imag()*r.elem());
}
//
//
//////////////////////////////////////////////
#define REALT float
#define Nc 3
struct PrimitiveClovTriang
{
RScalar<REALT> diag[2][2*Nc];
RComplex<REALT> offd[2][2*Nc*Nc-Nc];
};
__global__ void kernel(bool hasOrderedRep, int * siteTable,
PrimitiveClovTriang* tri)
{
RScalar<REALT> zip=0;
int N = 2*Nc;
int site;
//
// First if-block results in an error,
// second, runs fine! Since hasOrderedRep
// is true, the code blocks should be
// identical.
//
#if 1
if (hasOrderedRep) {
site = blockDim.x * blockIdx.x +
blockDim.x * gridDim.x * blockIdx.y +
threadIdx.x;
} else {
int idx0 = blockDim.x * blockIdx.x +
blockDim.x * gridDim.x * blockIdx.y +
threadIdx.x;
site = ((int*)(siteTable))[idx0];
}
#else
site = blockDim.x * blockIdx.x + blockDim.x * gridDim.x * blockIdx.y + threadIdx.x;
#endif
int site_neg_logdet=0;
for(int block=0; block < 2; block++) {
RScalar<REALT> inv_d[6];
RComplex<REALT> inv_offd[15];
RComplex<REALT> v[6];
RScalar<REALT> diag_g[6];
for(int i=0; i < N; i++) {
inv_d[i] = tri[site].diag[block][i];
}
for(int i=0; i < 15; i++) {
inv_offd[i] =tri[site].offd[block][i];
}
for(int j=0; j < N; ++j) {
for(int i=0; i < j; i++) {
int elem_ji = j*(j-1)/2 + i;
RComplex<REALT> A_ii = RComplex<REALT>( inv_d[i], zip );
v[i] = A_ii*RComplex<REALT>(inv_offd[elem_ji].real(),-inv_offd[elem_ji].imag());
}
v[j] = RComplex<REALT>(inv_d[j],zip);
for(int k=0; k < j; k++) {
int elem_jk = j*(j-1)/2 + k;
v[j] -= inv_offd[elem_jk]*v[k];
}
inv_d[j].elem() = v[j].real();
for(int k=j+1; k < N; k++) {
int elem_kj = k*(k-1)/2 + j;
for(int l=0; l < j; l++) {
int elem_kl = k*(k-1)/2 + l;
inv_offd[elem_kj] -= inv_offd[elem_kl] * v[l];
}
inv_offd[elem_kj] /= v[j];
}
}
RScalar<REALT> one;
one.elem() = (REALT)1;
for(int i=0; i < N; i++) {
diag_g[i].elem() = one.elem()/inv_d[i].elem();
// ((PScalar<PScalar<RScalar<float> > > *)(args->dev_ptr[ 1 ] ))[site] .elem().elem().elem() += log(fabs(inv_d[i].elem()));
if( inv_d[i].elem() < 0 ) {
site_neg_logdet++;
}
}
RComplex<REALT> sum;
for(int k = 0; k < N; ++k) {
for(int i = 0; i < k; ++i) {
v[i].real()=v[i].imag()=0;
}
v[k] = RComplex<REALT>(diag_g[k],zip);
for(int i = k+1; i < N; ++i) {
v[i].real()=v[i].imag()=0;
for(int j = k; j < i; ++j) {
int elem_ij = i*(i-1)/2+j;
v[i] -= inv_offd[elem_ij] *inv_d[j]*v[j];
}
v[i] *= diag_g[i];
}
for(int i = N-2; (int)i >= (int)k; --i) {
for(int j = i+1; j < N; ++j) {
int elem_ji = j*(j-1)/2 + i;
v[i] -= RComplex<REALT>(inv_offd[elem_ji].real(),-inv_offd[elem_ji].imag()) * v[j];
}
}
inv_d[k].elem() = v[k].real();
for(int i = k+1; i < N; ++i) {
int elem_ik = i*(i-1)/2+k;
inv_offd[elem_ik] = v[i];
}
}
for(int i=0; i < N; i++) {
tri[site].diag[block][i] = inv_d[i];
}
for(int i=0; i < 15; i++) {
tri[site].offd[block][i] = inv_offd[i];
}
}
if( site_neg_logdet != 0 ) {
}
}
int main()
{
int sites=1;
dim3 blocksPerGrid( 1 , 1 , 1 );
dim3 threadsPerBlock( sites , 1, 1);
PrimitiveClovTriang* tri_dev;
int * siteTable;
cudaMalloc( (void**)&tri_dev , sizeof(PrimitiveClovTriang) * sites );
cudaMalloc( (void**)&siteTable , sizeof(int) * sites );
bool ord=true;
kernel<<< blocksPerGrid , threadsPerBlock , 0 >>>( ord , siteTable , tri_dev );
cudaDeviceSynchronize();
cudaError_t kernel_call = cudaGetLastError();
cout << "kernel call: " << string(cudaGetErrorString(kernel_call)) << endl;
cudaFree(tri_dev);
cudaFree(siteTable);
return(0);
}
请注意,siteTable 包含未初始化的数据。这很好,因为它没有被使用。我需要它来显示错误。
更新:
刚在另一台安装了 CUDA 4.0 的机器上试过。那里没有出现错误(相同的费米卡模型)。可能真的是 CUDA 4.2 的 NVCC 错误。由于他们从 CUDA 4.1 切换到 LLVM,这很可能是一个错误。
最佳答案
此问题似乎是由基于 LLVM 的 CUDA 工具链的早期版本中的细微编译器错误引起的。该错误不会在更现代的工具链版本(CUDA 6.x 和 7.x)上重现。
[此答案已添加为来自评论和编辑的社区 wiki 条目,以便将问题从 CUDA 标签的已回答列表中删除]
关于c++ - 只有在展开循环时才会出现未指定的启动错误?,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/10548151/
我找到了 this excellent question and answer它以 x/y(加上 center x/y 和 degrees/radians)开始并计算旋转- 到 x'/y'。这个计算很
全部: 我已经创建了一个 Windows 窗体和一个按钮。在另一个线程中,我试图更改按钮的文本,但它崩溃了;但是如果我尝试更改按钮的颜色,它肯定会成功。我认为如果您更改任何 Windows 窗体控件属
本网站的另一个问题已证实,C 中没有缩写的字面后缀,并且可以执行以下操作: short Number = (short)1; 但是转换它和不这样做有什么区别: short Number = 1; 您使
我有下表: ID (int) EMAIL (varchar(50)) CAMPAIGNID (int) isSubscribe (bit) isActionByUser (bit) 此表存储了用户对事
也就是说,无需触发Javascript事件即可改变的属性,如何保留我手动选中或取消选中的复选框的状态,然后复制到另一个地方? 运行下面的代码片段并选中或取消选中其中的一些,然后点击“复制”: $('#
我在网上找到的所有关于递增指针导致段错误的示例都涉及指针的取消引用 - 如果我只想递增它(例如在 for 循环的末尾)并且我不在乎它是否最终进入无效内存,因为我不会再使用它。例如,在这个程序中,每次迭
我有一个 Spring MVC REST 服务,它使用 XStream 将消息与 XML 相互转换。 有什么方法可以将请求和响应中的 xml(即正文)打印到普通的 log4j 记录器? 在 Contr
做我的任务有一个很大的挑战,那就是做相互依赖的任务我在这张照片中说的。假设我们有两个任务 A 和 B,执行子任务 A1、A2 和 B1、B2,假设任务 B 依赖于 A。 要理想地执行任务 B,您应该执
通过阅读该网站上的几个答案,我了解到 CoInitialize(Ex) should be called by the creator of a thread 。然后,在该线程中运行的任何代码都可以使
这个问题已经困扰我一段时间了。我以前从未真正使用过 ListViews,也没有使用过 FirebaseListAdapters。我想做的就是通过显示 id 和用户位置来启动列表的基础,但由于某种原因,
我很难解释这两个(看似简单)句子的含义: “受检异常由编译器在编译时检查” 这是什么意思?编译器检查是否捕获了所有已检查的异常(在代码中抛出)? “未经检查的异常在运行时检查,而不是编译时” 这句话中
我有一个包含排除子字符串的文本文件,我想迭代该文件以检查并返回不带排除子字符串的输入项。 这里我使用 python 2.4,因此下面的代码可以实现此目的,因为 with open 和 any 不起作用
Spring 的缓存框架能否了解请求上下文的身份验证状态,或者更容易推出自己的缓存解决方案? 最佳答案 尽管我发现这个用例 super 奇怪,但您可以为几乎任何与 SpEL 配合使用的内容设置缓存条件
我有以下函数模板: template HeldAs* duplicate(MostDerived *original, HeldAs *held) { // error checking omi
如果我的应用程序具有设备管理员/设备所有者权限(未获得 root 权限),我如何才能从我的应用程序中终止(或阻止启动)另一个应用程序? 最佳答案 设备所有者可以阻止应用程序: DevicePolicy
非常简单的问题,但我似乎无法让它正常工作。 我有一个组件,其中有一些 XSLT(用于导航)。它通过 XSLT TBB 使用 XSLT Mediator 发布。 发布后
我正在将一个对象拖动到一个可拖放的对象内,该对象也是可拖动的。放置对象后,它会嵌套在可放置对象内。同样,如果我将对象拖到可放置的外部,它就不再嵌套。 但是,如果我经常拖入和拖出可放置对象,则可拖动对象
我正在尝试为按钮和弹出窗口等多个指令实现“取消选择”功能。也就是说,我希望当用户单击不属于指令模板一部分的元素时触发我的函数。目前,我正在使用以下 JQuery 代码: $('body').click
我从 this question 得到了下面的代码,该脚本用于在 Google tasks 上更改 iframe[src="about:blank"] 内的 CSS使用 Chrome 扩展 Tempe
我有一些 @Mock 对象,但没有指定在该对象上调用方法的返回值。该方法返回 int (不是 Integer)。我很惊讶地发现 Mockito 没有抛出 NPE 并返回 0。这是预期的行为吗? 例如:
我是一名优秀的程序员,十分优秀!