www.digitalmars.com         C & C++   DMDScript  

digitalmars.D - dcompute is nearing minimal functionality

reply Nicholas Wilson <iamthewilsonator hotmail.com> writes:
DCompute is my effort to target CUDA and SPIR to enable hassle 
free native execution on the gpu.

It is a D library +modification of LDC and is available at 
https://github.com/thewilsonator/ldc/tree/dcompute

The compiler is nearing minimal functionality and is able to 
generate code for cuda's pxt format.

Required to build is Khronos's llvm 
https://github.com/KhronosGroup/SPIRV-LLVM
and the usual build requirements of LDC,i.e. a dmd compatible d 
compiler, cmake.

Post cmake configuration and generation you will possibly need to 
add -L-lLLVMSPIRVLib to the (l|g)dmd invocation.

The code at the moment is very hacky and very hardcoded, with a 
lot of magic, and currently requires a manual invocation of llc, 
but hopefully this will improve.
$car pointer.d
 compute module dcompute.types.pointer;
import ldc.attributes;


enum Private = 0;
enum Global = 1;
enum Shared = 2;
enum Constant = 3;
enum Generic = 4;

 kernel //magic attribute
void xdfgcmain(Pointer!(1,float) a,Pointer!(1,float) b , float p)
{
     *a = *b + p;
}

struct compute {} //really hacky and somewhat magic
pure:  trusted: nothrow:  nogc:

struct Pointer(uint p, T) if(p <= Generic) //magic type
{
     T* ptr;
     alias ptr this;
}

$../ldcbuild/bin/ldc2 pointer.d -output-ll -O1
$llc -mcpu=sm_20 gpusuff_ptx.ll -o kernel.ptx
$cat kernel.ptx
//
// Generated by LLVM NVPTX Back-End
//

.version 3.2
.target sm_20
.address_size 64

	// 
.globl	_D8dcompute5types7pointer9xdfgcmainFS8dcompute5types7pointer18__T7PointerVki1TfZ7PointerS8dcompute5types7pointer18__T7PointerVki1TfZ7PointerfZv
                                         // 
 _D8dcompute5types7pointer9xdfgcmainFS8dcompute5types7pointer18__T7PointerVki1TfZ7PointerS8dcompute5types7pointer18__T7PointerVki1TfZ7PointerfZv
.visible .entry 
_D8dcompute5types7pointer9xdfgcmainFS8dcompute5types7pointer18__T7PointerVki1TfZ7PointerS8dcompute5types7pointer18__T7PointerVki1TfZ7PointerfZv(
	.param .align 8 .b8 
_D8dcompute5types7pointer9xdfgcmainFS8dcompute5types7pointer18__T7PointerVki1TfZ7PointerS8dcompute5types7pointer18__T7PointerVki1TfZ7PointerfZv_param_0[8],
	.param .align 8 .b8 
_D8dcompute5types7pointer9xdfgcmainFS8dcompute5types7pointer18__T7PointerVki1TfZ7PointerS8dcompute5types7pointer18__T7PointerVki1TfZ7PointerfZv_param_1[8],
	.param .f32 
_D8dcompute5types7pointer9xdfgcmainFS8dcompute5types7pointer18__T7PointerVki1TfZ7PointerS8dcompute5types7pointer18__T7PointerVki1TfZ7PointerfZv_param_2
)
{
	.reg .f32 	%f<4>;
	.reg .s64 	%rd<3>;


	ld.param.u64 	%rd1, 
[_D8dcompute5types7pointer9xdfgcmainFS8dcompute5types7pointer18__T7PointerVki1TfZ7PointerS8dcompute5types7pointer18__T7PointerVki1TfZ7PointerfZv_param_0];
	ld.param.u64 	%rd2, 
[_D8dcompute5types7pointer9xdfgcmainFS8dcompute5types7pointer18__T7PointerVki1TfZ7PointerS8dcompute5types7pointer18__T7PointerVki1TfZ7PointerfZv_param_1];
	ld.param.f32 	%f1, 
[_D8dcompute5types7pointer9xdfgcmainFS8dcompute5types7pointer18__T7PointerVki1TfZ7PointerS8dcompute5types7pointer18__T7PointerVki1TfZ7PointerfZv_param_2];
	ld.global.f32 	%f2, [%rd2];
	add.rn.f32 	%f3, %f2, %f1;
	st.global.f32 	[%rd1], %f3;
	ret;
}
Jul 13 2016
next sibling parent reply Nicholas Wilson <iamthewilsonator hotmail.com> writes:
On Thursday, 14 July 2016 at 03:30:44 UTC, Nicholas Wilson wrote:

 DCompute is my effort to target CUDA and SPIR to enable hassle 
 free native execution on the gpu.

 [...]


and a HUGE thanks to Johan Engelen for helping me get to grips 
with ldc.

Any and all suggestions and feedback (and PRs) welcome! I'd like 
to try to get this done by the start of august when I start my 
engineering honours thesis.
Jul 13 2016
parent reply Andrei Alexandrescu <SeeWebsiteForEmail erdani.org> writes:
On 07/13/2016 11:53 PM, Nicholas Wilson wrote:

 I'd like to try to get this done by the start of august when I start my
 engineering honours thesis.


Is your thesis related to this project in any way? -- Andrei
Jul 14 2016
parent reply Nicholas Wilson <iamthewilsonator hotmail.com> writes:
On Thursday, 14 July 2016 at 13:06:05 UTC, Andrei Alexandrescu 
wrote:

 On 07/13/2016 11:53 PM, Nicholas Wilson wrote:

 I'd like to try to get this done by the start of august when I 
 start my
 engineering honours thesis.


 Is your thesis related to this project in any way? -- Andrei


No, it will be integrating light, infrared camera signals with 
wind direction and speed direction and try to predict the 
movement of fire fronts at a very local scale using a drone. I 
intend to use d to program the brains of the onboard or attached 
computer though. It's an electronics thesis as much as it is 
software.
Jul 14 2016
parent Nicholas Wilson <iamthewilsonator hotmail.com> writes:
On Thursday, 14 July 2016 at 13:23:48 UTC, Nicholas Wilson wrote:

 On Thursday, 14 July 2016 at 13:06:05 UTC, Andrei Alexandrescu 
 wrote:

 On 07/13/2016 11:53 PM, Nicholas Wilson wrote:

 I'd like to try to get this done by the start of august when 
 I start my
 engineering honours thesis.


 Is your thesis related to this project in any way? -- Andrei


 No, it will be integrating light, infrared camera signals with 
 wind direction and speed direction and try to predict the 
 movement of fire fronts at a very local scale using a drone. I 
 intend to use d to program the brains of the onboard or 
 attached computer though. It's an electronics thesis as much as 
 it is software.


Whoops forgot GPS. I'd like to get it finished before because I'm 
going to have much less time after i start.
Jul 14 2016
prev sibling next sibling parent reply Andrei Alexandrescu <SeeWebsiteForEmail erdani.org> writes:
On 07/13/2016 11:30 PM, Nicholas Wilson wrote:

 DCompute is my effort to target CUDA and SPIR to enable hassle free
 native execution on the gpu.

 It is a D library +modification of LDC and is available at
 https://github.com/thewilsonator/ldc/tree/dcompute


Whoa, now that's a sight for sore eyes. This is awesome, 
congratulations! -- Andrei
Jul 14 2016
parent Nicholas Wilson <iamthewilsonator hotmail.com> writes:
On Thursday, 14 July 2016 at 13:05:12 UTC, Andrei Alexandrescu 
wrote:

 On 07/13/2016 11:30 PM, Nicholas Wilson wrote:

 DCompute is my effort to target CUDA and SPIR to enable hassle 
 free
 native execution on the gpu.

 It is a D library +modification of LDC and is available at
 https://github.com/thewilsonator/ldc/tree/dcompute


 Whoa, now that's a sight for sore eyes. This is awesome, 
 congratulations! -- Andrei


Thanks! Still a way to go though, the backend of llvm is not 
cooperating. :(

But as soon as thats done its just a case of de-hackifying it and 
adding things like command line parameters, instead of hardcoding 
things. Then I get to start on wrapping those magic types and 
making a library out of it (writing in D instead if C++, so that 
should be more fun and less time spent waiting for the linker).
Jul 14 2016
prev sibling parent Nicholas Wilson <iamthewilsonator hotmail.com> writes:
On Thursday, 14 July 2016 at 03:30:44 UTC, Nicholas Wilson wrote:

 DCompute is my effort to target CUDA and SPIR to enable hassle 
 free native execution on the gpu.

 It is a D library +modification of LDC and is available at 
 https://github.com/thewilsonator/ldc/tree/dcompute

 The compiler is nearing minimal functionality and is able to 
 generate code for cuda's pxt format.


now generates .ptx's without the aid of llc.
Jul 14 2016