小编har*_*ism的帖子

为什么没有atomicAdd()双打作为CUDA 4.0或更高版本的一部分明确实现？

从CUDA编程指南4.1的附录F第97页开始,已经实现了以下版本的atomicAdd.

int atomicAdd(int* address, int val);
unsigned int atomicAdd(unsigned int* address,
                       unsigned int val);
unsigned long long int atomicAdd(unsigned long long int* address,
                                 unsigned long long int val);
float atomicAdd(float* address, float val)

同样的页面继续为我的项目中刚刚开始使用的下面的双打提供了一个小型的atomicAdd实现.

__device__ double atomicAdd(double* address, double val)
{
    unsigned long long int* address_as_ull =
                             (unsigned long long int*)address;
    unsigned long long int old = *address_as_ull, assumed;
    do {
        assumed = old;
old = atomicCAS(address_as_ull, assumed,
                        __double_as_longlong(val +
                               __longlong_as_double(assumed)));
    } while (assumed …

我正在编写自己的图形库(是的,它的作业:)并使用cuda快速完成所有渲染和计算.

我有绘制填充三角形的问题.我这样编写了一个进程绘制一个三角形.当场景中有很多小三角形时它工作得很好,但是当三角形很大时它会完全破坏性能.

我的想法是做两次通过.首先计算仅包含扫描线信息的选项卡(从此处绘制到此处).这将是每个过程计算的三角形,就像当前算法一样.在第二遍中,确实绘制了每个三角形有多个进程的扫描线.

但它会足够快吗？也许有更好的解决方案？

在带有FindCUDA的CMAKE中,给定一个输入文件filename.cu,该cuda_compile_ptx命令生成表单cuda_compile_ptx_generated_filename.cu.ptx的输出文件名,但我需要输出文件名为表单filename.ptx.

是否有一种简单的方法可以使这项工作？

我正在学习OpenACC(使用PGI的编译器)并尝试优化矩阵乘法示例.到目前为止我提出的最快的实现如下:

void matrix_mul(float *restrict r, float *a, float *b, int N, int accelerate){

#pragma acc data copyin (a[0: N * N ], b[0: N * N]) copyout (r [0: N * N ]) if(accelerate)
{
# pragma acc region if(accelerate)
{
# pragma acc loop independent vector(32) 
for (int j = 0; j < N; j ++)
{    
   # pragma acc loop independent vector(32) 
   for (int i = 0; i < N ; i ++ )
   {
      float sum = …

我想在CUDA代码中实例化一个类,它在同一个块中与其他线程共享一些成员.

但是,在尝试编译以下代码时,我收到错误:»属性"shared"在这里不适用«(nvcc版本4.2).

class SharedSomething {

public:
    __shared__ int i; // this is not allowed
};

__global__ void run() {

    SharedSomething something;
}

这背后的理由是什么？是否有解决方案来实现所需的行为(跨越一个块的类的共享成员)？

更新:while()下面的条件由编译器优化,因此两个线程都跳过条件并进入CS甚至带有-O0标志.有谁知道为什么编译器这样做？顺便说一句,声明全局变量volatile会导致程序因某些奇怪的原因而挂起...

我阅读了CUDA编程指南,但我仍然不清楚CUDA如何处理与全局内存相关的内存一致性.(这与内存层次结构不同)基本上,我正在运行试图破坏顺序一致性的测试.我使用的算法是Peterson的内核函数内两个线程之间互斥的算法:

flag[threadIdx.x] = 1; // both these are global
turn = 1-threadIdx.x;

while(flag[1-threadIdx.x] == 1 && turn == (1- threadIdx.x));
shared_gloabl_variable_x ++;

flag[threadIdx.x] = 0;

这非常简单.每个线程通过将其标志设置为1并通过将转向设置为另一个线程来获得关键部分来请求关键部分.在评估时while(),如果其他线程没有设置其标志,则请求线程可以安全地进入临界区.现在,这种方法的一个微妙问题是,如果编译器重新排序写入,那么写入将turn在写入之前执行flag.如果发生这种情况,两个线程将同时在CS中结束.这很容易用普通的Pthreads来证明,因为大多数处理器都没有实现顺序一致性.但GPU呢？

这两个线程都将处于相同的warp中.并且他们将以锁步模式执行他们的语句.但是当它们到达turn变量时,它们正在写入相同的变量,因此内部执行变为序列化(无论顺序是什么).现在,此时,获胜的线程是否继续进入while条件,还是等待另一个线程完成其写入,以便两者可以同时进行评估while()？路径将再次分叉while(),因为只有其中一个会赢,而另一个等待.

运行代码后,我得到它一直打破SC.我读取的值总是1,这意味着两个线程每次都以某种方式进入CS.这怎么可能(GPU按顺序执行指令)？(注意:我已经编译了它-O0,因此没有编译器优化,因此没有使用volatile).

我知道FFT实现是如何工作的(Cooley-Tuckey算法),我知道有一个CUFFT CUDA库可以快速计算1D或2D FFT,但我想知道在这个过程中如何利用CUDA并行性.

它与蝴蝶计算有关吗？(类似每个线程将部分数据加载到共享内存中,然后每个线程计算一个偶数项或一个奇项？)

我有:

• 使用cudaHostAlloc(..., cudaHostAllocMapped)或成功固定和映射的主机内存cudaHostRegister(..., cudaHostRegisterMapped);
• 已经使用获得了设备指针cudaHostGetDevicePointer(...).

我cudaMemcpy(..., cudaMemcpyDeviceToDevice)在src和dest设备指针上启动,指向通过上述技术获得的两个不同的固定+映射内存区域.一切正常.

问题:我应该继续这样做还是只使用传统的CPU风格,memcpy()因为一切都在系统内存中？...或者它们是否相同(即当src和dest固定时,是否cudaMemcpy映射到直线memcpy)？

(我仍在使用该cudaMemcpy方法,因为之前所有内容都在设备全局内存中,但由于gmem大小限制,因此已切换到固定内存)

我正在用CUDA中的原子做一些实验.我的一个大问题是当同一个块中运行的两个线程原子地访问同一个地址时,它们是如何表现的.我尝试使用atomicAdd进行一些测试并且它原子地工作但是当我使用atomicCAS尝试下面的代码时,结果不是我所期望的.有人有解释吗？

#include <cuda_runtime.h>
#include <iostream>
#include <cuComplex.h>
using namespace std;
__global__ void kernel(int * pointer)
{
    *pointer=0;
    *(pointer+threadIdx.x+1)=0;
    __syncthreads();
    *(pointer+threadIdx.x+1)=atomicCAS(pointer,0,100);
}
int main(int argc,char ** argv)
{
    int numThreads=40;
    dim3 threadsPerBlock;
    dim3 blocks;
    int o[numThreads+1];
    int * pointer;
    cudaMalloc(&pointer,sizeof(int)*(numThreads+1));
    cudaMemset(pointer,0,sizeof(int)*(numThreads+1));
    threadsPerBlock.x=numThreads;
    threadsPerBlock.y=1;
    threadsPerBlock.z=1;
    blocks.x=1;
    blocks.y=1;
    blocks.z=1;
    kernel <<<threadsPerBlock,blocks>>> (pointer);
    cudaMemcpy(o,pointer,sizeof(int)*(numThreads+1),cudaMemcpyDeviceToHost);



    for (int i=0;i<numThreads+1;i++)
            cout << o[i] << " ";

    cout << endl;

}

在上面的代码中,在同一个块内运行的atomicCAS访问相同的地址进行比较和交换...我的期望是只有一个atomicCAS会找到要比较的值0而其他所有人都会找到100,但奇怪的是输出我的计划是:

 100 100 0 0 0 0 0 0 0 0 0 0 0 0 0 0 …

定位不同浮点精度(float/double)的常用方法是使用typedef

typedef float Real;
//typedef double Real;

或者使用模板

template<typename Real>
...

这很方便,但是任何人都有想法如何使用CUDA类型float2/float3/...和make_float2/make_float3/...？当然,我可以为所有人制作#defines或typedef,但这看起来不是很优雅.