Pycuda似乎是不确定性的。

Question

我遇到了一个关于CUDA的奇怪问题，

在下面的代码片段中，

#include <stdio.h>

#define OUTPUT_SIZE         26

typedef $PRECISION REAL;

extern "C"    
{
    __global__ void test_coeff ( REAL* results )
    {
        int id      = blockDim.x * blockIdx.x + threadIdx.x;

        int out_index  = OUTPUT_SIZE * id;
        for (int i=0; i<OUTPUT_SIZE; i++)
        {               
            results[out_index+i]=id;
            printf("q");
        }
    }
}

当我通过pycuda编译并运行代码时，结果符合预期。但是如果我移除printf语句后，结果会很奇怪 - 大多数数组元素被正确填充，但部分元素则完全随机。

以下是完整的Python代码：

import numpy as np
import string

#pycuda stuff
import pycuda.driver as drv
import pycuda.autoinit

from pycuda.compiler import SourceModule

class MC:

    cudacodetemplate = """
    #include <stdio.h>

    #define OUTPUT_SIZE         26

    typedef $PRECISION REAL;

    extern "C"    
    {
        __global__ void test_coeff ( REAL* results )
        {
            int id      = blockDim.x * blockIdx.x + threadIdx.x;

            int out_index  = OUTPUT_SIZE * id;
            for (int i=0; i<OUTPUT_SIZE; i++)
            {               
                results[out_index+i]=id;
                //printf("q");
            }
        }
    }
    """

    def __init__(self, size, prec = np.float32):
        #800 meg should be enough . . .
        drv.limit.MALLOC_HEAP_SIZE = 1024*1024*800

        self.size       = size
        self.prec       = prec
        template        = string.Template(MC.cudacodetemplate)
        self.cudacode   = template.substitute( PRECISION = 'float' if prec==np.float32 else 'double')

        #self.module     = pycuda.compiler.SourceModule(self.cudacode, no_extern_c=True, options=['--ptxas-options=-v'])
        self.module     = SourceModule(self.cudacode, no_extern_c=True)

    def test(self, out_size):
        #try to precalc the co-efficients for just the elements of the vector that changes
        test  = np.zeros( ( 128, out_size*(2**self.size) ), dtype=self.prec )
        test2 = np.zeros( ( 128, out_size*(2**self.size) ), dtype=self.prec )

        test_coeff =  self.module.get_function ('test_coeff')
        test_coeff( drv.Out(test), block=(2**self.size,1,1), grid=( 128, 1 ) )
        test_coeff( drv.Out(test2), block=(2**self.size,1,1), grid=( 128, 1 ) )
        error = (test-test2)
        return error

if __name__ == '__main__':
    p1  = MC ( 5, np.float64 )
    err = p1.test(26)
    print err.max()
    print err.min()

基本上，在内核中使用printf，err为0 - 如果没有它，则会打印一些随机错误（在我的机器上最大约为2452，最小为-2583）。

我不知道为什么。

我正在运行cuda 4.2 on pycuda 2012.2（windows 7 64位），使用geforce 570。

谢谢。

- user1726633

---

抱歉，但我无法在64位Linux主机和GTX 670上使用CUDA 4.2重现此问题。每次我运行单精度和双精度版本时，都会通过您发布的内核。 - talonmies

我认为我的硬件有问题 - 尽管我不确定为什么4.2 GPU SDK中的所有其他CUDA程序都能正常工作。我将尝试在Linux上使用相同的硬件运行此程序 - 然后我会尝试在Windows上使用不同的硬件来测试... - user1726633

我不了解pycuda，但在C/C++中，您不能在__global__或__device__代码内使用printf函数。在pycuda中是否可能？ - szamil

@szamil，是可以的，在Fermi及其之后的GPU上，使用CUDA C/C++或pycuda。 - harrism

谢谢！在我的Quadro Fx 1600M上，使用计算能力1.1是不可能的。 - szamil

在搭载CUDA 5.0的GT 650M上没有问题。这可能是一个已经被修复的错误，无论是在CUDA还是驱动程序中。您可以在NVIDIA论坛上发布此类问题，并可能在其错误报告平台上进行报告。 - BenC

Answer 1

这很可能是编译器优化导致的。您将长度为OUTPUT_SIZE的内存块设置为循环常量值id。根据我的经验，编译器会将其优化为memcpy或其他操作，除非在循环中发生了其他事情，比如您的打印语句。此外，如果您不使用该内存块，则编译器可能会将整个循环优化掉。尝试调整优化级别，看看是否有不同的结果。

- Ethereal