spirv-1.2 branch?

[hughperkins]

so many versions and repos, and out of date readmes :-) . I'm currently following #44 , but shouldnt we be using branch spirv-1.2 instead of spirv-1.1 now? Per https://www.khronos.org/spir , seems like 1.2 is the current latest 'release' version of SPIR-V?

[hughperkins]

(by the way, changing the name from spir to spirv, and then restarting versioning at 1.0 seems confusing as anything. Were you ever confused by CA-OpenIngres versioning, or Java versioning? Same idea :-( )

[hughperkins]

Oh.... am I right in thinking that the SPIR-V 1.2 branch is actually called spir_12?

[pierremoreau]

SPIR and SPIR-V are different intermediate representations: the former was based on LLVM IR, while the latter was built from scratch with no links to LLVM IR as far as I understand.

There is no SPIR-V 1.2 branch yet. I have no idea if someone is working on it.

[hughperkins]

ok. spir-v 1.2 is marked as stable. how are people using spir-v 1.2, without source-code for it?

[pierremoreau]

I guess they aren’t using it. SPIR-V 1.2 is only needed by OpenCL 2.2, which is probably not that used yet.

[hughperkins]

Hmmm. Well. I was thinking of pointing coriander at it, and using Coriander to compile/run SPIRV-1.2 for OpenCL 1.2 devices. So... I need to get SPIRV-1.2 from upstream khronos repos? Or .... ?

[pierremoreau]

SPIR-V versions do not match 1:1 to OpenCL versions. :-)
OpenCL 1.0 to OpenCL 2.0 (included) do not support SPIR-V at all. OpenCL 2.1 was the first version to introduce SPIR-V (1.0) support.

[hughperkins]

Coriander is a compatilibyt layer. eg it can run large chunks of Tensorflow and Eigen on OpenCL 1.2 devices. Do you see any obvious technical reasons why Coriander would not be able to run OpenCL 2.2 code on OpenCL 1.2 devices?

[keryell]

The hard part will be to implement OpenCL 2.2 hardware capability features running on lower OpenCL 1.2 hardware capability features... :-(
Think about pipes, subgroups, fine-grain system SVM...

Same as, say, running CUDA 8.0 PTX on some CUDA 2.0 hardware.

[hughperkins]

pipes is for FPGA guys ;-)

subgroups... that means like __shfl, right?

"fine-grain system svm". Unsure on this point. I already started iplementing device-side virtualmem. eg/ie, the float *..[] in this struct works ok now :-) :

https://github.com/hughperkins/coriander/blob/master/test/endtoend/test_floatstarstar.cu#L91-L103

struct BoundedArrayUnion {
    float *bounded_array[8];
    float **unbounded_array;
};

__global__ void run_bounded_array_two(struct BoundedArrayUnion mystruct, int useUnbounded, int numBuffers, int N) {
    float **starstar = useUnbounded ? mystruct.unbounded_array : mystruct.bounded_array;
    for(int i = 0; i < numBuffers; i++) {
        for(int j = 0; j < N; j++) {
            starstar[i][j] = 123.0f + i + 1 + j;
        }
    }
}

[keryell]

There is also an easy solution of course: just run everything on the CPU. :-)

The issue is to have all OpenCL 2.2 running on top of OpenCL 1.2, not the OpenCL 1.2 subset of OpenCL 2.2 capabilities...

If you have an OpenCL 2.2 GPU with independent forward progress between subgroups and you have cooperative concurrent programming between the CPU and the OpenCL 2.2 subgroups using atomic operations through normal memory (via fine-grain system SVM), it seems difficult...

[johnkslang]

Sounds like most these questions are about tooling, but regarding location of the specifications, they are all kept in the SPIR-V registry, hosted by Khronos, at https://www.khronos.org/registry/spir-v.

[hughperkins]

No. its about the branches in the SPIR repo. Here is a list:

https://github.com/KhronosGroup/SPIR/branches

You can see that there is a branch for spirv_1.1, but we are missing the spirv_1.2 branch.