Alignement de structures avec PyOpenCL

mise à jour: l’ int4 dans mon kernel était faux.

J’utilise pyopencl mais je ne parviens pas à faire en sorte que l’alignement des structures fonctionne correctement. Dans le code ci-dessous, qui appelle le kernel deux fois, la valeur b est renvoyée correctement (sous la forme 1), mais la valeur c a une valeur “aléatoire”.

En d’autres termes: j’essaie de lire deux membres d’une structure. Je peux lire le premier mais pas le second. Pourquoi?

Le même problème se produit si j’utilise des tableaux structurés numpy ou un pack avec struct. Et les parameters _-atsortingbute__ dans les commentaires ne sont d’aucun secours.

Je soupçonne que je fais quelque chose de stupide ailleurs dans le code, mais je ne le vois pas. Toute aide appréciée.

 import struct as s import pyopencl as cl import numpy as n ctx = cl.create_some_context() queue = cl.CommandQueue(ctx) for use_struct in (True, False): if use_struct: a = s.pack('=ii',1,2) print(a, len(a)) a_dev = cl.Buffer(ctx, cl.mem_flags.WRITE_ONLY, len(a)) else: # a = n.array([(1,2)], dtype=n.dtype('2i4', align=True)) a = n.array([(1,2)], dtype=n.dtype('2i4')) print(a, a.itemsize, a.nbytes) a_dev = cl.Buffer(ctx, cl.mem_flags.WRITE_ONLY, a.nbytes) b = n.array([0], dtype='i4') print(b, b.itemsize, b.nbytes) b_dev = cl.Buffer(ctx, cl.mem_flags.READ_ONLY, b.nbytes) c = n.array([0], dtype='i4') print(c, c.itemsize, c.nbytes) c_dev = cl.Buffer(ctx, cl.mem_flags.READ_ONLY, c.nbytes) prg = cl.Program(ctx, """ typedef struct s { int4 f0; int4 f1 __atsortingbute__ ((packed)); // int4 f1 __atsortingbute__ ((aligned (4))); // int4 f1; } s; __kernel void test(__global const s *a, __global int4 *b, __global int4 *c) { *b = a->f0; *c = a->f1; } """).build() cl.enqueue_copy(queue, a_dev, a) event = prg.test(queue, (1,), None, a_dev, b_dev, c_dev) event.wait() cl.enqueue_copy(queue, b, b_dev) print(b) cl.enqueue_copy(queue, c, c_dev) print(c) 

La sortie (j’ai dû reformater en coupant + collant, donc j’ai peut-être légèrement modifié les sauts de ligne; j’ai également ajouté des commentaires indiquant les différentes valeurs d’impression):

 # first using struct /home/andrew/projects/personal/kultrung/env/bin/python3.2 /home/andrew/projects/personal/kultrung/src/kultrung/test6.py b'\x01\x00\x00\x00\x02\x00\x00\x00' 8 # the struct packed values [0] 4 4 # output buffer 1 [0] 4 4 # output buffer 2 /home/andrew/projects/personal/kultrung/env/lib/python3.2/site-packages/pyopencl/cache.py:343: UserWarning: Build succeeded, but resulted in non-empty logs: Build on  succeeded, but said: Build started Kernel  was successfully vectorized Done. warn("Build succeeded, but resulted in non-empty logs:\n"+message) [1] # the first value (correct) [240] # the second value (wrong) # next using numpy [[1 2]] 4 8 # the numpy struct [0] 4 4 # output buffer [0] 4 4 # output buffer /home/andrew/projects/personal/kultrung/env/lib/python3.2/site-packages/pyopencl/__init__.py:174: UserWarning: Build succeeded, but resulted in non-empty logs: Build on  succeeded, but said: Build started Kernel  was successfully vectorized Done. warn("Build succeeded, but resulted in non-empty logs:\n"+message) [1] # first value (ok) [67447488] # second value (wrong) Process finished with exit code 0 

Dans le programme OpenCL, essayez l’atsortingbut compacté sur la structure elle-même, à la place de l’un des membres:

 typedef struct s { int4 f0; int4 f1; } __atsortingbute__((packed)) s; 

Il se peut que, dans la mesure où l’atsortingbut packed ne figure que sur un seul membre de la structure, il est possible que la structure ne soit pas entièrement packed .

ok, je ne sais pas d’où je int4 – je pense que ce doit être une extension intel. le passage à AMD avec int comme type de kernel fonctionne comme prévu. Je vais poster plus à http://acooke.org/cute/Somesimple0.html une fois que j’ai nettoyé les choses.