pycuda wydaje niedeterministyczny

Question

Mam dziwny problem z CUDA

w poniższym fragmencie,

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

Kiedy skompilować i uruchomić kod (poprzez pycuda), to działa zgodnie z oczekiwaniami. Kiedy usuwam printf, wyniki są dziwne - większość tablic jest wypełniona poprawnie, ale niektóre z nich wydają się zupełnie przypadkowe.

oto pełny kod 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() 

Zasadniczo z printf w jądrze, ERR jest 0 - bez niego wypisuje jakiś przypadkowy błąd (na moim komputerze około 2452 (na max) i -2583 (dla min))

Nie mam pojęcia dlaczego.

mam działa CUDA 4.2 na pycuda 2012.2 (Windows 7 64bit) z geforce 570.

dzięki.

Answer 1

Jest to najprawdopodobniej spowodowane optymalizacją kompilatora. Ustawiasz długość bloku OUTPUT_SIZE na wartość stałej pętli id. Z mojego doświadczenia kompilator zoptymalizuje to do memcpy lub whathaveyou chyba że jest coś innego dzieje się w pętli - tj. Wyciągu drukowania. Ponadto, jeśli nie korzystasz z tego bloku pamięci, kompilator może zoptymalizować całą pętlę. Spróbuj manipulować poziomami optymalizacji i sprawdź, czy masz różne wyniki.