Au stade du développement du module d'interopérabilité OpenCL de Abstrasy 2.0 qui correspond au moment où j'écris cet article, il est possible de créer un programme OpenCL prêt l'emploi (donc chargé et compilé dans la mémoire de la carte - ceci ne comprend pas l'utilisation proprement dites).
Voyez le code suivant :
(import 'OpenCL)
(define 'plateform (OpenCL:Platforms:ALL 0))
(define 'devices (plateform:get-all-devices))
(define 'context (OpenCL:Context devices))
(define 'program
(context:program
`
__kernel void hello_kernel( __global const float *a, __global const float *b, __global float *result )
{
int gid = get_global_id(0);
result[gid] = a[gid] + b[gid];
}
`
)
)
# ---------------------------------------------------------------------------------------------------------
(display "context:NUM_DEVICES : " (context:NUM_DEVICES))
(for 'd (iterator context:DEVICES) {display " - " d})
(display)
# ---------------------------------------------------------------------------------------------------------
(display "program:NUM_DEVICES : " (program:NUM_DEVICES))
(for 'd (iterator program:DEVICES) {display " - " d})
(display)
(display "program:SOURCE : " (program:SOURCE))
(display)
(display "program:BUILD_OPTIONS :")
(for 'd (iterator program:BUILD_OPTIONS) {display " - " d})
(display)
(display "program:BUILD_STATUS :")
(for 'd (iterator program:BUILD_STATUS) {display " - " d})
(display)
(display "program:BUILD_LOG :")
(for 'd (iterator program:BUILD_LOG) {
(display " - " d)
(display)
})
(display)
(display "program:BUILD_BINARIES :")
(for 'd (iterator program:BUILD_BINARIES) {
(display " - " d)
(try{ display "\n --- convertible to human readable code : ---\n\n" (string d '!utf8) } catch{ })
(display)
})
(display)
Bien sûr, il y a plus de 5 lignes dans ce script, mais en réalité, tout est fait dans la première partie et si vous comptez bien les lignes effectives en Abstrasy, vous constaterez qu'il n'y a que 5 lignes (1 import et 4 define).
La variable 'context contient le contexte et 'program le programme prêt à l'emploi pour toutes les cartes répertoriées dans le contexte. Rappelons qu'un programme peut être exécuté sur plusieurs cartes (accélérateurs, GPUs ou CPUs). Chaque carte n'utilise donc pas nécessairement le même code binaire (le source compilé). OpenCL permet l'utilisation de systèmes hétérogènes.
Sur mon poste de tests, je n'ai qu'une seule carte NVIDIA QUADRO. Voyons le résultat :
context:NUM_DEVICES : 1
- Quadro K5000 : NVIDIA Corporation : OpenCL 1.1 CUDA
program:NUM_DEVICES : 1
- Quadro K5000 : NVIDIA Corporation : OpenCL 1.1 CUDA
program:SOURCE :
__kernel void hello_kernel( __global const float *a, __global const float *b, __global float *result )
{
int gid = get_global_id(0);
result[gid] = a[gid] + b[gid];
}
program:BUILD_OPTIONS :
-
program:BUILD_STATUS :
- '@!SUCCESS
program:BUILD_LOG :
-
program:BUILD_BINARIES :
- (bytes 0x2Fy 0x2Fy 0x0Ay 0x2Fy 0x2Fy 0x20y 0x47y 0x65y 0x6Ey 0x65y 0x72y 0x61y 0x74y 0x65y 0x64y 0x20y 0x62y 0x79y
0x20y 0x4Ey 0x56y 0x49y 0x44y 0x49y 0x41y 0x20y 0x4Ey 0x56y 0x56y 0x4Dy 0x20y 0x43y 0x6Fy 0x6Dy 0x70y 0x69y 0x6Cy 0x65y
0x72y 0x0Ay 0x2Fy 0x2Fy 0x20y 0x43y 0x6Fy 0x6Dy 0x70y 0x69y 0x6Cy 0x65y 0x72y 0x20y 0x62y 0x75y 0x69y 0x6Cy 0x74y 0x20y
0x6Fy 0x6Ey 0x20y 0x54y 0x68y 0x75y 0x20y 0x4Ay 0x61y 0x6Ey 0x20y 0x20y 0x39y 0x20y 0x30y 0x34y 0x3Ay 0x35y 0x35y 0x3Ay
0x32y 0x34y 0x20y 0x32y 0x30y 0x31y 0x34y 0x20y 0x28y 0x31y 0x33y 0x38y 0x39y 0x32y 0x33y 0x39y 0x37y 0x32y 0x34y 0x29y
0x0Ay 0x2Fy 0x2Fy 0x20y 0x44y 0x72y 0x69y 0x76y 0x65y 0x72y 0x20y 0x33y 0x33y 0x31y 0x2Ey 0x33y 0x38y 0x0Ay 0x2Fy 0x2Fy
0x0Ay 0x0Ay 0x2Ey 0x76y 0x65y 0x72y 0x73y 0x69y 0x6Fy 0x6Ey 0x20y 0x33y 0x2Ey 0x30y 0x0Ay 0x2Ey 0x74y 0x61y 0x72y 0x67y
0x65y 0x74y 0x20y 0x73y 0x6Dy 0x5Fy 0x33y 0x30y 0x2Cy 0x20y 0x74y 0x65y 0x78y 0x6Dy 0x6Fy 0x64y 0x65y 0x5Fy 0x69y 0x6Ey
0x64y 0x65y 0x70y 0x65y 0x6Ey 0x64y 0x65y 0x6Ey 0x74y 0x0Ay 0x2Ey 0x61y 0x64y 0x64y 0x72y 0x65y 0x73y 0x73y 0x5Fy 0x73y
0x69y 0x7Ay 0x65y 0x20y 0x33y 0x32y 0x0Ay 0x0Ay 0x0Ay 0x2Ey 0x65y 0x6Ey 0x74y 0x72y 0x79y 0x20y 0x68y 0x65y 0x6Cy 0x6Cy
0x6Fy 0x5Fy 0x6By 0x65y 0x72y 0x6Ey 0x65y 0x6Cy 0x28y 0x0Ay 0x09y 0x2Ey 0x70y 0x61y 0x72y 0x61y 0x6Dy 0x20y 0x2Ey 0x75y
0x33y 0x32y 0x20y 0x2Ey 0x70y 0x74y 0x72y 0x20y 0x2Ey 0x67y 0x6Cy 0x6Fy 0x62y 0x61y 0x6Cy 0x20y 0x2Ey 0x61y 0x6Cy 0x69y
0x67y 0x6Ey 0x20y 0x34y 0x20y 0x68y 0x65y 0x6Cy 0x6Cy 0x6Fy 0x5Fy 0x6By 0x65y 0x72y 0x6Ey 0x65y 0x6Cy 0x5Fy 0x70y 0x61y
0x72y 0x61y 0x6Dy 0x5Fy 0x30y 0x2Cy 0x0Ay 0x09y 0x2Ey 0x70y 0x61y 0x72y 0x61y 0x6Dy 0x20y 0x2Ey 0x75y 0x33y 0x32y 0x20y
0x2Ey 0x70y 0x74y 0x72y 0x20y 0x2Ey 0x67y 0x6Cy 0x6Fy 0x62y 0x61y 0x6Cy 0x20y 0x2Ey 0x61y 0x6Cy 0x69y 0x67y 0x6Ey 0x20y
0x34y 0x20y 0x68y 0x65y 0x6Cy 0x6Cy 0x6Fy 0x5Fy 0x6By 0x65y 0x72y 0x6Ey 0x65y 0x6Cy 0x5Fy 0x70y 0x61y 0x72y 0x61y 0x6Dy
0x5Fy 0x31y 0x2Cy 0x0Ay 0x09y 0x2Ey 0x70y 0x61y 0x72y 0x61y 0x6Dy 0x20y 0x2Ey 0x75y 0x33y 0x32y 0x20y 0x2Ey 0x70y 0x74y
0x72y 0x20y 0x2Ey 0x67y 0x6Cy 0x6Fy 0x62y 0x61y 0x6Cy 0x20y 0x2Ey 0x61y 0x6Cy 0x69y 0x67y 0x6Ey 0x20y 0x34y 0x20y 0x68y
0x65y 0x6Cy 0x6Cy 0x6Fy 0x5Fy 0x6By 0x65y 0x72y 0x6Ey 0x65y 0x6Cy 0x5Fy 0x70y 0x61y 0x72y 0x61y 0x6Dy 0x5Fy 0x32y 0x0Ay
0x29y 0x0Ay 0x7By 0x0Ay 0x09y 0x2Ey 0x72y 0x65y 0x67y 0x20y 0x2Ey 0x66y 0x33y 0x32y 0x20y 0x09y 0x25y 0x66y 0x3Cy 0x34y
0x3Ey 0x3By 0x0Ay 0x09y 0x2Ey 0x72y 0x65y 0x67y 0x20y 0x2Ey 0x73y 0x33y 0x32y 0x20y 0x09y 0x25y 0x72y 0x3Cy 0x31y 0x34y
0x3Ey 0x3By 0x0Ay 0x0Ay 0x0Ay 0x09y 0x6Cy 0x64y 0x2Ey 0x70y 0x61y 0x72y 0x61y 0x6Dy 0x2Ey 0x75y 0x33y 0x32y 0x20y 0x09y
0x25y 0x72y 0x35y 0x2Cy 0x20y 0x5By 0x68y 0x65y 0x6Cy 0x6Cy 0x6Fy 0x5Fy 0x6By 0x65y 0x72y 0x6Ey 0x65y 0x6Cy 0x5Fy 0x70y
0x61y 0x72y 0x61y 0x6Dy 0x5Fy 0x30y 0x5Dy 0x3By 0x0Ay 0x09y 0x6Cy 0x64y 0x2Ey 0x70y 0x61y 0x72y 0x61y 0x6Dy 0x2Ey 0x75y
0x33y 0x32y 0x20y 0x09y 0x25y 0x72y 0x36y 0x2Cy 0x20y 0x5By 0x68y 0x65y 0x6Cy 0x6Cy 0x6Fy 0x5Fy 0x6By 0x65y 0x72y 0x6Ey
0x65y 0x6Cy 0x5Fy 0x70y 0x61y 0x72y 0x61y 0x6Dy 0x5Fy 0x31y 0x5Dy 0x3By 0x0Ay 0x09y 0x6Cy 0x64y 0x2Ey 0x70y 0x61y 0x72y
0x61y 0x6Dy 0x2Ey 0x75y 0x33y 0x32y 0x20y 0x09y 0x25y 0x72y 0x37y 0x2Cy 0x20y 0x5By 0x68y 0x65y 0x6Cy 0x6Cy 0x6Fy 0x5Fy
0x6By 0x65y 0x72y 0x6Ey 0x65y 0x6Cy 0x5Fy 0x70y 0x61y 0x72y 0x61y 0x6Dy 0x5Fy 0x32y 0x5Dy 0x3By 0x0Ay 0x09y 0x2Fy 0x2Fy
0x20y 0x69y 0x6Ey 0x6Cy 0x69y 0x6Ey 0x65y 0x20y 0x61y 0x73y 0x6Dy 0x0Ay 0x09y 0x6Dy 0x6Fy 0x76y 0x2Ey 0x75y 0x33y 0x32y
0x20y 0x09y 0x25y 0x72y 0x31y 0x2Cy 0x20y 0x25y 0x65y 0x6Ey 0x76y 0x72y 0x65y 0x67y 0x33y 0x3By 0x0Ay 0x09y 0x2Fy 0x2Fy
0x20y 0x69y 0x6Ey 0x6Cy 0x69y 0x6Ey 0x65y 0x20y 0x61y 0x73y 0x6Dy 0x0Ay 0x09y 0x2Fy 0x2Fy 0x20y 0x69y 0x6Ey 0x6Cy 0x69y
0x6Ey 0x65y 0x20y 0x61y 0x73y 0x6Dy 0x0Ay 0x09y 0x6Dy 0x6Fy 0x76y 0x2Ey 0x75y 0x33y 0x32y 0x20y 0x09y 0x25y 0x72y 0x32y
0x2Cy 0x20y 0x25y 0x6Ey 0x74y 0x69y 0x64y 0x2Ey 0x78y 0x3By 0x0Ay 0x09y 0x2Fy 0x2Fy 0x20y 0x69y 0x6Ey 0x6Cy 0x69y 0x6Ey
0x65y 0x20y 0x61y 0x73y 0x6Dy 0x0Ay 0x09y 0x2Fy 0x2Fy 0x20y 0x69y 0x6Ey 0x6Cy 0x69y 0x6Ey 0x65y 0x20y 0x61y 0x73y 0x6Dy
0x0Ay 0x09y 0x6Dy 0x6Fy 0x76y 0x2Ey 0x75y 0x33y 0x32y 0x20y 0x09y 0x25y 0x72y 0x33y 0x2Cy 0x20y 0x25y 0x63y 0x74y 0x61y
0x69y 0x64y 0x2Ey 0x78y 0x3By 0x0Ay 0x09y 0x2Fy 0x2Fy 0x20y 0x69y 0x6Ey 0x6Cy 0x69y 0x6Ey 0x65y 0x20y 0x61y 0x73y 0x6Dy
0x0Ay 0x09y 0x2Fy 0x2Fy 0x20y 0x69y 0x6Ey 0x6Cy 0x69y 0x6Ey 0x65y 0x20y 0x61y 0x73y 0x6Dy 0x0Ay 0x09y 0x6Dy 0x6Fy 0x76y
0x2Ey 0x75y 0x33y 0x32y 0x20y 0x09y 0x25y 0x72y 0x34y 0x2Cy 0x20y 0x25y 0x74y 0x69y 0x64y 0x2Ey 0x78y 0x3By 0x0Ay 0x09y
0x2Fy 0x2Fy 0x20y 0x69y 0x6Ey 0x6Cy 0x69y 0x6Ey 0x65y 0x20y 0x61y 0x73y 0x6Dy 0x0Ay 0x09y 0x61y 0x64y 0x64y 0x2Ey 0x73y
0x33y 0x32y 0x20y 0x09y 0x25y 0x72y 0x38y 0x2Cy 0x20y 0x25y 0x72y 0x34y 0x2Cy 0x20y 0x25y 0x72y 0x31y 0x3By 0x0Ay 0x09y
0x6Dy 0x61y 0x64y 0x2Ey 0x6Cy 0x6Fy 0x2Ey 0x73y 0x33y 0x32y 0x20y 0x09y 0x25y 0x72y 0x39y 0x2Cy 0x20y 0x25y 0x72y 0x33y
0x2Cy 0x20y 0x25y 0x72y 0x32y 0x2Cy 0x20y 0x25y 0x72y 0x38y 0x3By 0x0Ay 0x09y 0x73y 0x68y 0x6Cy 0x2Ey 0x62y 0x33y 0x32y
0x20y 0x09y 0x25y 0x72y 0x31y 0x30y 0x2Cy 0x20y 0x25y 0x72y 0x39y 0x2Cy 0x20y 0x32y 0x3By 0x0Ay 0x09y 0x61y 0x64y 0x64y
0x2Ey 0x73y 0x33y 0x32y 0x20y 0x09y 0x25y 0x72y 0x31y 0x31y 0x2Cy 0x20y 0x25y 0x72y 0x35y 0x2Cy 0x20y 0x25y 0x72y 0x31y
0x30y 0x3By 0x0Ay 0x09y 0x61y 0x64y 0x64y 0x2Ey 0x73y 0x33y 0x32y 0x20y 0x09y 0x25y 0x72y 0x31y 0x32y 0x2Cy 0x20y 0x25y
0x72y 0x36y 0x2Cy 0x20y 0x25y 0x72y 0x31y 0x30y 0x3By 0x0Ay 0x09y 0x6Cy 0x64y 0x2Ey 0x67y 0x6Cy 0x6Fy 0x62y 0x61y 0x6Cy
0x2Ey 0x66y 0x33y 0x32y 0x20y 0x09y 0x25y 0x66y 0x31y 0x2Cy 0x20y 0x5By 0x25y 0x72y 0x31y 0x32y 0x5Dy 0x3By 0x0Ay 0x09y
0x6Cy 0x64y 0x2Ey 0x67y 0x6Cy 0x6Fy 0x62y 0x61y 0x6Cy 0x2Ey 0x66y 0x33y 0x32y 0x20y 0x09y 0x25y 0x66y 0x32y 0x2Cy 0x20y
0x5By 0x25y 0x72y 0x31y 0x31y 0x5Dy 0x3By 0x0Ay 0x09y 0x61y 0x64y 0x64y 0x2Ey 0x66y 0x33y 0x32y 0x20y 0x09y 0x25y 0x66y
0x33y 0x2Cy 0x20y 0x25y 0x66y 0x32y 0x2Cy 0x20y 0x25y 0x66y 0x31y 0x3By 0x0Ay 0x09y 0x61y 0x64y 0x64y 0x2Ey 0x73y 0x33y
0x32y 0x20y 0x09y 0x25y 0x72y 0x31y 0x33y 0x2Cy 0x20y 0x25y 0x72y 0x37y 0x2Cy 0x20y 0x25y 0x72y 0x31y 0x30y 0x3By 0x0Ay
0x09y 0x73y 0x74y 0x2Ey 0x67y 0x6Cy 0x6Fy 0x62y 0x61y 0x6Cy 0x2Ey 0x66y 0x33y 0x32y 0x20y 0x09y 0x5By 0x25y 0x72y 0x31y
0x33y 0x5Dy 0x2Cy 0x20y 0x25y 0x66y 0x33y 0x3By 0x0Ay 0x09y 0x72y 0x65y 0x74y 0x3By 0x0Ay 0x7Dy 0x0Ay 0x0Ay 0x0Ay)
--- convertible to human readable code : ---
//
// Generated by NVIDIA NVVM Compiler
// Compiler built on Thu Jan 9 04:55:24 2014 (1389239724)
// Driver 331.38
//
.version 3.0
.target sm_30, texmode_independent
.address_size 32
.entry hello_kernel(
.param .u32 .ptr .global .align 4 hello_kernel_param_0,
.param .u32 .ptr .global .align 4 hello_kernel_param_1,
.param .u32 .ptr .global .align 4 hello_kernel_param_2
)
{
.reg .f32 %f<4>;
.reg .s32 %r<14>;
ld.param.u32 %r5, [hello_kernel_param_0];
ld.param.u32 %r6, [hello_kernel_param_1];
ld.param.u32 %r7, [hello_kernel_param_2];
// inline asm
mov.u32 %r1, %envreg3;
// inline asm
// inline asm
mov.u32 %r2, %ntid.x;
// inline asm
// inline asm
mov.u32 %r3, %ctaid.x;
// inline asm
// inline asm
mov.u32 %r4, %tid.x;
// inline asm
add.s32 %r8, %r4, %r1;
mad.lo.s32 %r9, %r3, %r2, %r8;
shl.b32 %r10, %r9, 2;
add.s32 %r11, %r5, %r10;
add.s32 %r12, %r6, %r10;
ld.global.f32 %f1, [%r12];
ld.global.f32 %f2, [%r11];
add.f32 %f3, %f2, %f1;
add.s32 %r13, %r7, %r10;
st.global.f32 [%r13], %f3;
ret;
}
Ready...
Cas particulier, le code binaire produit sur le matériel NVIDIA est «human readable». Il suffit donc de convertir le vecteur de bytes résultant de la compilation pour obtenir le code désassemblé du résultat.
Les lignes qui suivent permettent d'énumérer les états du programme. Remarquez que le code OpenCL est délimité dans une chaîne de caractères multi-lignes (pour cela on utilise l'apostrophe oblique «`»).
Pour vérifier la prise ne compte d'erreurs éventuelles, je modifie le code OpenCL en commettant volontairement une erreur:
(define 'program
(context:program
`
__kernel void hello_kernel( __goba const float *a, __global const float *b, __global float *result )
{
int gid = get_global_id(0);
result[gid] = a[gid] + b[gid];
}
`
)
)
Notez l'erreur dans la liste des arguments du kernel, __goba au lieu de __global.
Voyons le résultat:
context:NUM_DEVICES : 1
- Quadro K5000 : NVIDIA Corporation : OpenCL 1.1 CUDA
program:NUM_DEVICES : 1
- Quadro K5000 : NVIDIA Corporation : OpenCL 1.1 CUDA
program:SOURCE :
__kernel void hello_kernel( __goba const float *a, __global const float *b, __global float *result )
{
int gid = get_global_id(0);
result[gid] = a[gid] + b[gid];
}
program:BUILD_OPTIONS :
-
program:BUILD_STATUS :
- '@!ERROR
program:BUILD_LOG :
- :2:35: error: unknown type name '__goba'
__kernel void hello_kernel( __goba const float *a, __global const float *b, __global float *result )
^
:2:42: error: expected ')'
__kernel void hello_kernel( __goba const float *a, __global const float *b, __global float *result )
^
:2:33: note: to match this '('
__kernel void hello_kernel( __goba const float *a, __global const float *b, __global float *result )
^
:2:42: error: parameter name omitted
__kernel void hello_kernel( __goba const float *a, __global const float *b, __global float *result )
^
:5:9: error: use of undeclared identifier 'result'
result[gid] = a[gid] + b[gid];
^
program:BUILD_BINARIES :
- (bytes)
--- convertible to human readable code : ---
Ready...
Comme prévu, program:BUILD_STATUS indique bien une erreur.
Cependant, program:BUILD_LOG fournit une description suffisante du problème rencontré.
Bien entendu, le code binaire est un vecteur de bytes vide… Logique!…
En conclusion, nous disposons des premiers éléments de base du module d'interopérabilité OpenCL : Platforms, Platform, Device, Context et Program. C'est déjà très prometteur.
 |
 |
 |
 |
|
|
|
||||||||
