RFR(M) GRAAL-234 - PTX code loader

[Morris Meyer]

Folks,

I have extended the Graal PTX back-end to incorporate a light and tight 
linkage to Nvidia GPUs.  I've managed to have the Graal-generated PTX 
code from Java methods properly compile using the Cuda 5.0 drivers on my 
GTX 660 graphics card.

If you are wondering about this comment in src/share/vm/runtime/thread.cpp:

     // Initialize the os module before using TLS
     os::init();

+   // probe for warp capability
+   gpu::init();
+

Per http://docs.nvidia.com/cuda/pdf/ptx_isa_3.1.pdf - the Parallel 
Thread Execution ISA Version 3.1 from Nvidia:

"A warp is a maximal subset of threads from a single cooperative thread 
array (CTA), such that the threads execute the same instructions at the 
same time. ...Each grid of CTAs has a 1D, 2D or 3D shape"

I am currently assigned JDK-8013168, which is to extend the set of code 
pointers from Method to support multiple architectures.  This will be 
necessary to adapt the loaded GPU kernel to internal HotSpot method 
invocation, as well as architectures that have a heterogeneous CPU+APU 
configuration.

WEBREV - http://cr.openjdk.java.net/~morris/GRAAL-234.01
JIRA - https://lafo.ssw.uni-linz.ac.at/jira/browse/GRAAL-234

         --morris

"to boldly go"

__________________________________________________________________

Output:

vendor: 0x000010DE
device: 0x000011C0
model: NVIDIA GeForce GTX 660
gpu_bsd::probe_gpu(APPLE): 1
gpu::Ptx::probe_linkage
gpu_ptx::probe_linkage(APPLE): 1
gpu::initialize_gpu
gpu_ptx::_cuda_cu_init: 0
gpu_ptx::_cuda_cu_device_get_count(1): 0
gpu_ptx::_cuda_cu_device_get(0): 0
gpu_ptx::_cuda_cu_device_compute_capability(major 3, minor 0): 0
gpu_ptx::_cuda_cu_device_get_name(GeForce GTX 660): 0
gpu_ptx::_cuda_cu_ctx_create(3b013800): 0
gpu_ptx::initialize_gpu(): 1
gpu::generate_kernel
gpu::Ptx::generate_kernel
gpu_ptx::_cuda_cu_module_load_data_ex(39b575a0): 0
gpu_ptx::jit_log_buffer

gpu_ptx::_cuda_cu_module_get_function(testAddConst1I):3b119600 0
testAddConst1I:
     .version 1.4
     .target sm_10
.entry testAddConst1I (
     .param .u32 param0
) {
       .reg .pred %p,%q;
       .reg .u32 %r<16>;
L121:
     add.s32 %r2, %r2, 1;
     mov.s32 %r0, %r2;
     exit;
}