pycuda parece no determinista

Question

Tengo un extraño problema con CUDA,

En el siguiente fragmento,

#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"); 
     } 
    } 
} 

Cuando compilar y ejecutar el código (a través de pycuda), funciona como se esperaba. Cuando elimino el printf, los resultados son raros: la mayor parte de la matriz está poblada correctamente, pero parte de ella parece completamente aleatoria.

Aquí está el código Python completo:

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() 

Básicamente, con el printf en el núcleo, el err es 0 - sin que se imprime un cierto error aleatorio (en mi máquina alrededor de 2452 (para el máximo), y -2,583 (para el min))

no tengo ni idea de por qué.

He corriendo CUDA 4.2 en pycuda 2012.2 (Windows 7 de 64 bits) con una GeForce 570.

Gracias.

Answer 1

Esto es muy probablemente debido a la optimización del compilador. Está configurando un bloque de memoria OUTPUT_SIZE en longitud al valor de constante de bucle de id. En mi experiencia, el compilador lo optimizará a una memcpy o lo que sea que a menos que haya algo más en el ciclo, es decir, su estado de impresión. Además, si no utiliza ese bloque de memoria, el compilador puede optimizar todo el bucle. Pruebe con sus niveles de optimización y vea si tiene resultados diferentes.