CUDA surface with Unity OpenGL interop (source code)

Procedural images (1024 x 1024) are generated in real-time from CUDA kernel mainImage.
Everything works fine in Editor (OpenGL 4.5), but with standalone build it fails, error:

OPENGL NATIVE PLUG-IN ERROR: GL_INVALID_OPERATION: Operation illegal in current state
(Filename: C:\buildslave\unity\build\Runtime\GfxDevice\opengles\GfxDeviceGLES.cpp Line: 340)

Code (CSharp):

1. using UnityEngine;
2. using System;
3. using System.Collections;
4. using System.Runtime.InteropServices;
5. public class Surface : MonoBehaviour
6. {
7.     [DllImport("Surface")]
8.     static extern IntPtr Execute();
9.  
10.     [DllImport("Surface")]
11.     static extern void SetTime(float t);
12.    
13.     public RenderTexture RT;
14.     private Material InternalMaterial;
15.     void Awake()
16.     {
17.         InternalMaterial = new Material(Shader.Find("Sprites/Default"));
18.     }
19.     void RenderToBuffer(RenderTexture destination, Material material)
20.     {
21.         RenderTexture.active = destination;
22.         GL.PushMatrix();
23.         GL.LoadOrtho();
24.         material.SetPass(0);
25.         GL.Begin(GL.QUADS);
26.         GL.MultiTexCoord2(0, 0.0f, 0.0f);
27.         GL.Vertex3(0.0f, 0.0f, 0.0f);
28.         GL.MultiTexCoord2(0, 1.0f, 0.0f);
29.         GL.Vertex3(1.0f, 0.0f, 0.0f);
30.         GL.MultiTexCoord2(0, 1.0f, 1.0f);
31.         GL.Vertex3(1.0f, 1.0f, 0.0f);
32.         GL.MultiTexCoord2(0, 0.0f, 1.0f);
33.         GL.Vertex3(0.0f, 1.0f, 0.0f);
34.         GL.End();
35.         GL.Clear(false, true, Color.black);
36.         GL.IssuePluginEvent(Execute(), 1);
37.         GL.PopMatrix();
38.     }
39.     void Update()
40.     {
41.         SetTime(Time.time);
42.         RenderToBuffer(RT,InternalMaterial);
43.     }
44. }

Code (CSharp):

1. // nvcc -o Surface.dll --shared Surface.cu -arch=sm_30 opengl32.lib user32.lib gdi32.lib
2. #include <windows.h>
3. #include <cuda_gl_interop.h>
4. #include <GL/gl.h>
5.  
6. #define width 1024
7. #define height 1024
8.  
9. typedef GLuint(APIENTRY *PFNGLCREATEPROGRAMPROC) ();
10. typedef GLuint(APIENTRY *PFNGLCREATESHADERPROC) (GLenum t);
11. typedef void(APIENTRY *PFNGLSHADERSOURCEPROC) (GLuint s, GLsizei c, const char*const*string, const GLint* i);
12. typedef void(APIENTRY *PFNGLCOMPILESHADERPROC) (GLuint s);
13. typedef void(APIENTRY *PFNGLATTACHSHADERPROC) (GLuint p, GLuint s);
14. typedef void(APIENTRY *PFNGLLINKPROGRAMPROC) (GLuint p);
15. typedef void(APIENTRY *PFNGLUSEPROGRAMPROC) (GLuint p);
16. typedef void(APIENTRY *PFNGLGENBUFFERSPROC) (GLsizei n, GLuint *b);
17. typedef void(APIENTRY *PFNGLBINDBUFFERPROC) (GLenum t, GLuint b);
18. typedef void(APIENTRY *PFNGLBUFFERDATAPROC) (GLenum t, ptrdiff_t s, const GLvoid *d, GLenum u);
19. typedef void(APIENTRY *PFNGLBINDVERTEXARRAYPROC) (GLuint a);
20. typedef void(APIENTRY *PFNGLENABLEVERTEXATTRIBARRAYPROC) (GLuint i);
21. typedef void(APIENTRY *PFNGLVERTEXATTRIBPOINTERPROC) (GLuint i, GLint s, GLenum t, GLboolean n, GLsizei k, const void *p);
22. typedef void(APIENTRY *PFNGLGENVERTEXARRAYSPROC) (GLsizei n, GLuint *a);
23. typedef void(APIENTRY *PFNGLDELETEVERTEXARRAYSPROC) (GLsizei n, const GLuint *a);
24. typedef GLint(APIENTRY *PFNGLGETUNIFORMLOCATIONPROC) (GLuint p, const char *n);
25. typedef void (APIENTRY *PFNGLUNIFORM1FVPROC) (GLint k, GLsizei c, const GLfloat *v);
26. typedef void (APIENTRY *PFNGLUNIFORM1IPROC) (GLint l, GLint v);
27. typedef void (APIENTRY *PFNGLACTIVETEXTUREPROC) (GLenum t);
28.  
29. unsigned int PS, VertexArrayID, VertexBuffer, store;
30. static const GLfloat vertices[] = {-1.0f,-1.0f,0.0f,1.0f,-1.0f,0.0f,-1.0f,1.0f,0.0f,1.0f,-1.0f,0.0f,1.0f,1.0f,0.0f,-1.0f,1.0f,0.0f};
31. surface<void, cudaSurfaceType2D> RenderSurface;
32. static float _Time;
33.  
34. extern "C" void __declspec(dllexport) __stdcall SetTime (float t) { _Time = t; }
35.  
36. static const char* VertexShader = \
37.     "#version 450 core \n"
38.     "layout (location=0) in vec3 position;"
39.     "void main()"
40.     "{"
41.         "gl_Position=vec4(position,1.0);"
42.     "}";
43.  
44. static const char* FragmentShader = \
45.     "#version 450 core \n"
46.     "layout (location=0) out vec4 color;"
47.     "uniform sampler2D pattern;"
48.     "void main()"
49.     "{"
50.         "vec2 uv = gl_FragCoord.xy / vec2(1024,1024);"
51.         "color = texture(pattern,uv);"
52.     "}";
53.  
54. __device__ float smoothstep(float a, float b, float x)
55. {
56.     float t = fmaxf(0.0f, fminf((x - a)/(b - a), 1.0f));
57.     return t*t*(3.0f-(2.0f*t));
58. }
59.    
60. __global__ void mainImage (float iTime)
61. {
62.     int x = blockIdx.x*blockDim.x + threadIdx.x;
63.     int y = blockIdx.y*blockDim.y + threadIdx.y;
64.     float2 iResolution = make_float2((float)width, (float)height);
65.     float2 fragCoord = make_float2((float)x, (float)y);
66.     float2 uv = make_float2((2.0f * fragCoord.x / iResolution.x - 1.0f),(2.0f * fragCoord.y / iResolution.y - 1.0f));
67.     float L = sqrt(uv.x*uv.x+uv.y*uv.y)*4.0f;
68.     float K = atan2(uv.y, uv.x)+iTime;
69.     float X = fmod(sin(K*3.0f), cos(K*3.0f));
70.     float Y = fmod(cos(K*3.0f), sin(K*3.0f));
71.     float3 A = make_float3(X, X, 1.0f-X);
72.     float3 B = make_float3(Y+1.0f, Y+2.0f, Y+5.0f);
73.     float3 T = make_float3(1.0f-sin(iTime)*L, 1.0f-cos(iTime)*L, 1.0f-cos(iTime)*L);
74.     float3 color = make_float3(0.9f-smoothstep(A.x,B.x,T.x), 0.9f-smoothstep(A.y,B.y,T.y), 0.9f-smoothstep(A.z,B.z,T.z));
75.     uchar4 fragColor = make_uchar4(0.9f*color.x*255, 0.9f*color.y*255, 0.5f*color.z*255, 255);  
76.     surf2Dwrite(fragColor, RenderSurface, x*sizeof(uchar4), y, cudaBoundaryModeClamp);
77. }  
78.  
79. void Init(int p)
80. {
81.     cudaGraphicsResource *resource;
82.     cudaArray *image;
83.     ((PFNGLUSEPROGRAMPROC)wglGetProcAddress("glUseProgram"))(PS);
84.     ((PFNGLACTIVETEXTUREPROC)wglGetProcAddress("glActiveTexture"))(0x84C0);
85.     glGenTextures(1, &store);
86.     glBindTexture(GL_TEXTURE_2D, store);
87.     glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
88.     glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
89.     glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, width, height, 0, GL_RGBA, GL_UNSIGNED_BYTE, 0);
90.     int num = ((PFNGLGETUNIFORMLOCATIONPROC)wglGetProcAddress("glGetUniformLocation"))(p, "pattern");
91.     ((PFNGLUNIFORM1IPROC)wglGetProcAddress("glUniform1i"))(num, 0);  
92.     cudaGraphicsGLRegisterImage(&resource, store, GL_TEXTURE_2D, cudaGraphicsRegisterFlagsSurfaceLoadStore);
93.     cudaGraphicsMapResources(1, &resource, 0);
94.     cudaGraphicsSubResourceGetMappedArray(&image, resource, 0, 0);
95.     cudaBindSurfaceToArray(RenderSurface, image);
96. }  
97.  
98. int MakeShaders(const char* VS, const char* FS)
99. {
100.     int p = ((PFNGLCREATEPROGRAMPROC)wglGetProcAddress("glCreateProgram"))();
101.     int s1 = ((PFNGLCREATESHADERPROC)wglGetProcAddress("glCreateShader"))(0x8B31);  
102.     int s2 = ((PFNGLCREATESHADERPROC)wglGetProcAddress("glCreateShader"))(0x8B30);  
103.     ((PFNGLSHADERSOURCEPROC)wglGetProcAddress("glShaderSource"))(s1,1,&VS,0);
104.     ((PFNGLSHADERSOURCEPROC)wglGetProcAddress("glShaderSource"))(s2,1,&FS,0);  
105.     ((PFNGLCOMPILESHADERPROC)wglGetProcAddress("glCompileShader"))(s1);
106.     ((PFNGLCOMPILESHADERPROC)wglGetProcAddress("glCompileShader"))(s2);  
107.     ((PFNGLATTACHSHADERPROC)wglGetProcAddress("glAttachShader"))(p,s1);
108.     ((PFNGLATTACHSHADERPROC)wglGetProcAddress("glAttachShader"))(p,s2);  
109.     ((PFNGLLINKPROGRAMPROC)wglGetProcAddress("glLinkProgram"))(p);
110.     return p;
111. }
112.  
113. void Rendering()
114. {
115.     dim3 block(8, 8, 1);
116.     dim3 grid(width/block.x, height/block.y, 1);
117.     mainImage <<< grid, block >>> (_Time);
118.     glBindTexture(GL_TEXTURE_2D, store);
119.     ((PFNGLUSEPROGRAMPROC)wglGetProcAddress("glUseProgram"))(PS);
120.     ((PFNGLGENVERTEXARRAYSPROC)wglGetProcAddress("glGenVertexArrays")) (1, &VertexArrayID);      
121.     ((PFNGLBINDVERTEXARRAYPROC)wglGetProcAddress("glBindVertexArray")) (VertexArrayID);  
122.     ((PFNGLGENBUFFERSPROC)wglGetProcAddress("glGenBuffers"))(1, &VertexBuffer);
123.     ((PFNGLBINDBUFFERPROC)wglGetProcAddress("glBindBuffer"))(0x8892, VertexBuffer);
124.     ((PFNGLBUFFERDATAPROC)wglGetProcAddress("glBufferData"))(0x8892, sizeof(vertices), vertices, 0x88E4);
125.     ((PFNGLENABLEVERTEXATTRIBARRAYPROC)wglGetProcAddress("glEnableVertexAttribArray"))(0);
126.     ((PFNGLBINDBUFFERPROC)wglGetProcAddress("glBindBuffer"))(0x8892, VertexBuffer);
127.     ((PFNGLVERTEXATTRIBPOINTERPROC)wglGetProcAddress("glVertexAttribPointer"))(0,3, GL_FLOAT, GL_FALSE, 0,(void*)0 );
128.     glDrawArrays(GL_TRIANGLES, 0, 6);
129.     ((PFNGLDELETEVERTEXARRAYSPROC)wglGetProcAddress("glDeleteVertexArrays"))(1, &VertexArrayID);
130. }
131.  
132. typedef enum UnityGfxRenderer
133. {
134.     kUnityGfxRendererNull = 4,
135.     kUnityGfxRendererOpenGLCore = 17,
136. } UnityGfxRenderer;
137.  
138. typedef enum UnityGfxDeviceEventType
139. {
140.     kUnityGfxDeviceEventInitialize = 0,
141.     kUnityGfxDeviceEventShutdown = 1,
142.     kUnityGfxDeviceEventBeforeReset = 2,
143.     kUnityGfxDeviceEventAfterReset = 3,
144. } UnityGfxDeviceEventType;
145.    
146. struct UnityInterfaceGUID
147. {
148.     UnityInterfaceGUID(unsigned long long high, unsigned long long low) : m_GUIDHigh(high) , m_GUIDLow(low) { }
149.     unsigned long long m_GUIDHigh;
150.     unsigned long long m_GUIDLow;
151. };
152.  
153. struct IUnityInterface {};
154. typedef void (__stdcall * IUnityGraphicsDeviceEventCallback)(UnityGfxDeviceEventType eventType);
155.  
156. struct IUnityInterfaces
157. {
158.     IUnityInterface* (__stdcall* GetInterface)(UnityInterfaceGUID guid);
159.     void(__stdcall* RegisterInterface)(UnityInterfaceGUID guid, IUnityInterface * ptr);
160.     template<typename INTERFACE>
161.     INTERFACE* Get()
162.     {
163.         return static_cast<INTERFACE*>(GetInterface(UnityInterfaceGUID(0x7CBA0A9CA4DDB544ULL, 0x8C5AD4926EB17B11ULL)));
164.     }
165.     void Register(IUnityInterface* ptr)
166.     {
167.         RegisterInterface(UnityInterfaceGUID(0x7CBA0A9CA4DDB544ULL, 0x8C5AD4926EB17B11ULL), ptr);
168.     }
169. };
170.  
171. struct IUnityGraphics : IUnityInterface
172. {
173.     void(__stdcall* RegisterDeviceEventCallback)(IUnityGraphicsDeviceEventCallback callback);
174. };
175.  
176. typedef void (__stdcall* UnityRenderingEvent)(int eventId);
177. typedef void(__stdcall* UnregisterDeviceEventCallback)(IUnityGraphicsDeviceEventCallback callback);
178. static UnityGfxRenderer DeviceType = kUnityGfxRendererNull;
179.  
180. static void __stdcall OnGraphicsDeviceEvent(UnityGfxDeviceEventType eventType)
181. {
182.     if (eventType == kUnityGfxDeviceEventInitialize)
183.     {
184.         DeviceType = kUnityGfxRendererOpenGLCore;
185.         PS = MakeShaders(VertexShader, FragmentShader);
186.         Init(PS);
187.     }
188.     if (eventType == kUnityGfxDeviceEventShutdown)
189.     {
190.         DeviceType = kUnityGfxRendererNull;
191.     }
192. }
193.  
194. static void __stdcall OnRenderEvent(int eventID)
195. {
196.     Rendering();
197. }
198.  
199. extern "C" void    __declspec(dllexport) __stdcall UnityPluginLoad(IUnityInterfaces* unityInterfaces)
200. {
201.     IUnityInterfaces* s_UnityInterfaces = unityInterfaces;
202.     IUnityGraphics* s_Graphics = s_UnityInterfaces->Get<IUnityGraphics>();
203.     s_Graphics->RegisterDeviceEventCallback(OnGraphicsDeviceEvent);
204.     OnGraphicsDeviceEvent(kUnityGfxDeviceEventInitialize);
205. }
206.  
207. extern "C" void __declspec(dllexport) __stdcall UnityPluginUnload()
208. {
209.     UnregisterDeviceEventCallback(OnGraphicsDeviceEvent);  
210. }
211.  
212. extern "C" UnityRenderingEvent __declspec(dllexport) __stdcall Execute()
213. {
214.     return OnRenderEvent;
215. }

Screenshot: