Execution of a simple Lambda Method on GPU

[Doug Simon]

Tom,

You can now file bug reports for Graal via JBS: https://bugs.openjdk.java.net/browse/GRAAL

For smaller issues such as the one below, email is maybe still a better option.

-Doug

On Jan 20, 2014, at 5:45 PM, Deneau, Tom <tom.deneau at amd.com> wrote:

> Bharadwaj --
> 
> There is a bug in Method::is_lambda() in that it does not return false
> if methodPrefix == 0.  Actually I don't understand why the gcc
> compiler would not complain about this.
> 
> Also, I noticed the CompilationPolicy::must_be_compiled, is_lambda path
> is not hit in the debug build.  Is this by design?
> 
> -- Tom
> 
>> -----Original Message-----
>> From: sumatra-dev-bounces at openjdk.java.net [mailto:sumatra-dev-
>> bounces at openjdk.java.net] On Behalf Of S. Bharadwaj Yadavalli
>> Sent: Friday, January 17, 2014 4:26 PM
>> To: sumatra-dev at openjdk.java.net; graal-dev at openjdk.java.net
>> Subject: Execution of a simple Lambda Method on GPU
>> 
>> 
>> A quick update on reaching the next milestone in executing Lambda
>> methods on GPUs as a step towards the goals of Project Sumatra.
>> 
>> The JVM in Graal repo, now has the the ability to recognize a simple
>> Lambda method defined in main class method, schedule compilation of such
>> a method to target PTX backend, offload execution of the generated PTX
>> code on a supported nVidia GPU hardware and get back the result to the
>> VM.
>> 
>> This was made possible owing to recent enhancements and code
>> reorganization made in Graal compiler to enhance support for GPU
>> backends.
>> 
>> The current implementation is in experimental stage. For example,
>> executing the Java class [1] results in [2]. Additional refinements to
>> choose candidate Java methods for offloading to GPU are planned. We look
>> forward to any open source community involvement that will help as move
>> faster towards the next milestones of Project Sumatra.
>> 
>> Regards,
>> 
>> Bharadwaj
>> 
>> [1]
>> interface BinaryOperation {
>>     int apply(int a, int b);
>> }
>> 
>> class FindSumL {
>>    public static void main(String args[]) {
>>        BinaryOperation add = (x, y) -> x + y;
>>        int result = add.apply(8, 12);
>>        System.out.println("Sum is " + result);
>>    }
>> }
>> 
>> [2]
>> $ ./mx.sh vmg -XX:+TraceGPUInteraction -XX:-BootstrapGraal -G:Threads=1
>> FindSumL
>> Found supported nVidia GPU device vendor : 0x10de device 0x06dd
>> gpu_linux::probe_gpu(): 1
>> [CUDA] Success: library linkage
>> CUDA driver initialization: Success
>> [CUDA] Number of compute-capable devices found: 2
>> [CUDA] Got the handle of first compute-device
>> [CUDA] Unified addressing support on device 0: 1
>> [CUDA] Using GeForce GTX 780
>> Compiling Lambda method FindSumL::lambda$main$0 to PTX
>> [CUDA] Success: Created context for device: 0
>> [CUDA] Success: Set current context for device: 0
>> [CUDA] PTX Kernel
>>     .version 3.0
>>     .target sm_30
>> .entry lambda$main$0 (
>>     .param .s32 param1,
>>     .param .s32 param2,
>>     .param .u64 param0
>> ) {
>>     .reg .s32 %r3;
>>     .reg .s32 %r4;
>>     .reg .s32 %r5;
>>     .reg .u64 %r6;
>> L0:
>>     ld.param.s32 %r3, [param1 + 0];
>>     ld.param.s32 %r4, [param2 + 0];
>>     add.s32 %r5, %r4, %r3;
>>     ld.param.u64 %r6, [param0 + 0];
>>     st.global.s32 [%r6 + 0], %r5;
>>     ret;
>> }
>> 
>> [CUDA] Function name : lambda$main$0
>> [CUDA] Loaded data for PTX Kernel
>> [CUDA] Got function handle for lambda$main$0 kernel address
>> 0x7fba686bfd80
>> [CUDA] Generated kernel
>> External method:FindSumL.lambda$main$0(II)I
>> installCode0: ExternalCompilationResult
>> [CUDA] launching kernel
>> [CUDA] Success: Kernel Launch: X: 1 Y: 1 Z: 1
>> [CUDA] Success: Synchronized launch kernel
>> [CUDA] Success: Freed device memory of return value
>> [CUDA] Success: Destroy context
>> Sum is 20
>> 
>> 
>> 
> 
>