fVertexGenD.dfm 30 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265266267268269270271272273274275276277278279280281282283284285286287288289290291292293294295296297298299300301302303304305306307308309310311312313314315316317318319320321322323324325326327328329330331332333334335336337338339340341342343344345346347348349350351352353354355356357358359360361362363364365366367368369370371372373374375376377378379380381382383384385386387388389390391392393394395396397398399400401402403404405406407408409410411412413414415416417418419420421422423424425426427428429430431432433434435436437438439440441442443444445446447448449450451452453454455456457458459460461462463464465466467468469470471472473474475476477478479480481482483484485486487488489490491492493494495496497498499500501502503504505506507508509510511512513514515516517518519520521522523524525526527528529530531532533534535536537538539540541542543544545546547548549550551552553554555556557558559560561562563564565566567568569570571572573574575576577578579580581582583584585586587588589590591592593594595596597598599600601602603604605606607608609610611612613614615616617618619620621622623624625626627628629630631632633634635636637638639640641642643644645646647648649650651652653654655656657658659660661662663664665666667668669670671672673674675676677678679680681682683684685686687688689690691692693694695696697698699700701702703704705706707708709710711712713714715716717718719720721722723724725726727728729730731732733734735736737738739740741742743744745746747748749750751752753754755756757758759760761762763764765766767768769770771772773774775776777778779780781782783784785
  1. object Form1: TForm1
  2. Left = 423
  3. Top = 62
  4. Caption = 'CUDA fit GLScene'
  5. ClientHeight = 512
  6. ClientWidth = 512
  7. Color = clBtnFace
  8. Font.Charset = DEFAULT_CHARSET
  9. Font.Color = clWindowText
  10. Font.Height = -11
  11. Font.Name = 'Tahoma'
  12. Font.Style = []
  13. OldCreateOrder = False
  14. OnCreate = FormCreate
  15. PixelsPerInch = 96
  16. TextHeight = 13
  17. object GLSceneViewer1: TGLSceneViewer
  18. Left = 0
  19. Top = 0
  20. Width = 512
  21. Height = 512
  22. Camera = GLCamera1
  23. Buffer.BackgroundColor = clBlack
  24. FieldOfView = 157.897079467773400000
  25. Align = alClient
  26. TabOrder = 0
  27. end
  28. object GLScene1: TGLScene
  29. Left = 24
  30. Top = 16
  31. object GLCamera1: TGLCamera
  32. DepthOfView = 100.000000000000000000
  33. FocalLength = 50.000000000000000000
  34. TargetObject = GLDummyCube1
  35. Position.Coordinates = {0000C03F0000C03F000000400000803F}
  36. end
  37. object GLDummyCube1: TGLDummyCube
  38. CubeSize = 2.000000000000000000
  39. VisibleAtRunTime = True
  40. object CUDAFeedBackMesh1: TCUDAFeedBackMesh
  41. Attributes = <
  42. item
  43. Name = 'Position'
  44. GLSLType = GLSLType4F
  45. end>
  46. Shader = GLSLShader1
  47. CommonKernelFunction = MakeDotField
  48. Visible = False
  49. end
  50. end
  51. end
  52. object GLCadencer1: TGLCadencer
  53. Scene = GLScene1
  54. OnProgress = GLCadencer1Progress
  55. Left = 96
  56. Top = 16
  57. end
  58. object GLSimpleNavigation1: TGLSimpleNavigation
  59. Form = Owner
  60. GLSceneViewer = GLSceneViewer1
  61. FormCaption = 'Form1 - %FPS'
  62. KeyCombinations = <
  63. item
  64. ShiftState = [ssLeft, ssRight]
  65. Action = snaZoom
  66. end
  67. item
  68. ShiftState = [ssLeft]
  69. Action = snaMoveAroundTarget
  70. end
  71. item
  72. ShiftState = [ssRight]
  73. Action = snaMoveAroundTarget
  74. end>
  75. Left = 24
  76. Top = 72
  77. end
  78. object GLCUDADevice1: TGLCUDADevice
  79. SelectDevice = 'GeForce GTX 260 (1)'
  80. Left = 448
  81. Top = 16
  82. end
  83. object GLCUDA1: TGLCUDA
  84. ComputingDevice = GLCUDADevice1
  85. OnOpenGLInteropInit = GLCUDA1OpenGLInteropInit
  86. Left = 448
  87. Top = 72
  88. object MainModule: TCUDAModule
  89. Code.Strings = (
  90. #9'.version 1.4'
  91. #9'.target sm_13'
  92. #9'// compiled with C:\Program Files\NVIDIA GPU Computing Toolkit\' +
  93. 'CUDA\v3.2\bin/../open64/lib//be.exe'
  94. #9'// nvopencc 3.2 built on 2010-11-06'
  95. ''
  96. #9'//-----------------------------------------------------------'
  97. #9'// Compiling C:/Users/YARUNA~1/AppData/Local/Temp/tmpxft_00000a' +
  98. 'b0_00000000-11_temp.cpp3.i (C:/Users/YARUNA~1/AppData/Local/Temp' +
  99. '/ccBI#.a02728)'
  100. #9'//-----------------------------------------------------------'
  101. ''
  102. #9'//-----------------------------------------------------------'
  103. #9'// Options:'
  104. #9'//-----------------------------------------------------------'
  105. #9'// Target:ptx, ISA:sm_13, Endian:little, Pointer Size:32'
  106. #9'// -O3'#9'(Optimization level)'
  107. #9'// -g0'#9'(Debug level)'
  108. #9'// -m2'#9'(Report advisories)'
  109. #9'//-----------------------------------------------------------'
  110. ''
  111. #9'.file'#9'1'#9'"C:/Users/YARUNA~1/AppData/Local/Temp/tmpxft_00000ab0_0' +
  112. '0000000-10_temp.cudafe2.gpu"'
  113. #9'.file'#9'2'#9'"C:\Program Files\Microsoft Visual Studio 9.0\VC\INCLUD' +
  114. 'E\crtdefs.h"'
  115. #9'.file'#9'3'#9'"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v3.' +
  116. '2\include\crt/device_runtime.h"'
  117. #9'.file'#9'4'#9'"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v3.' +
  118. '2\include\host_defines.h"'
  119. #9'.file'#9'5'#9'"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v3.' +
  120. '2\include\builtin_types.h"'
  121. #9'.file'#9'6'#9'"c:\program files\nvidia gpu computing toolkit\cuda\v3.' +
  122. '2\include\device_types.h"'
  123. #9'.file'#9'7'#9'"c:\program files\nvidia gpu computing toolkit\cuda\v3.' +
  124. '2\include\driver_types.h"'
  125. #9'.file'#9'8'#9'"c:\program files\nvidia gpu computing toolkit\cuda\v3.' +
  126. '2\include\surface_types.h"'
  127. #9'.file'#9'9'#9'"c:\program files\nvidia gpu computing toolkit\cuda\v3.' +
  128. '2\include\texture_types.h"'
  129. #9'.file'#9'10'#9'"c:\program files\nvidia gpu computing toolkit\cuda\v3' +
  130. '.2\include\vector_types.h"'
  131. #9'.file'#9'11'#9'"c:\program files\nvidia gpu computing toolkit\cuda\v3' +
  132. '.2\include\builtin_types.h"'
  133. #9'.file'#9'12'#9'"c:\program files\nvidia gpu computing toolkit\cuda\v3' +
  134. '.2\include\host_defines.h"'
  135. #9'.file'#9'13'#9'"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v3' +
  136. '.2\include\device_launch_parameters.h"'
  137. #9'.file'#9'14'#9'"c:\program files\nvidia gpu computing toolkit\cuda\v3' +
  138. '.2\include\crt\storage_class.h"'
  139. #9'.file'#9'15'#9'"C:\Program Files\Microsoft Visual Studio 9.0\VC\INCLU' +
  140. 'DE\time.h"'
  141. #9'.file'#9'16'#9'"c:\program files\nvidia gpu computing toolkit\cuda\v3' +
  142. '.2\include\texture_fetch_functions.h"'
  143. #9'.file'#9'17'#9'"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v3' +
  144. '.2\include\common_functions.h"'
  145. #9'.file'#9'18'#9'"c:\program files\nvidia gpu computing toolkit\cuda\v3' +
  146. '.2\include\math_functions.h"'
  147. #9'.file'#9'19'#9'"c:\program files\nvidia gpu computing toolkit\cuda\v3' +
  148. '.2\include\math_constants.h"'
  149. #9'.file'#9'20'#9'"c:\program files\nvidia gpu computing toolkit\cuda\v3' +
  150. '.2\include\device_functions.h"'
  151. #9'.file'#9'21'#9'"c:\program files\nvidia gpu computing toolkit\cuda\v3' +
  152. '.2\include\sm_11_atomic_functions.h"'
  153. #9'.file'#9'22'#9'"c:\program files\nvidia gpu computing toolkit\cuda\v3' +
  154. '.2\include\sm_12_atomic_functions.h"'
  155. #9'.file'#9'23'#9'"c:\program files\nvidia gpu computing toolkit\cuda\v3' +
  156. '.2\include\sm_13_double_functions.h"'
  157. #9'.file'#9'24'#9'"c:\program files\nvidia gpu computing toolkit\cuda\v3' +
  158. '.2\include\sm_20_atomic_functions.h"'
  159. #9'.file'#9'25'#9'"c:\program files\nvidia gpu computing toolkit\cuda\v3' +
  160. '.2\include\sm_20_intrinsics.h"'
  161. #9'.file'#9'26'#9'"c:\program files\nvidia gpu computing toolkit\cuda\v3' +
  162. '.2\include\surface_functions.h"'
  163. #9'.file'#9'27'#9'"c:\program files\nvidia gpu computing toolkit\cuda\v3' +
  164. '.2\include\math_functions_dbl_ptx3.h"'
  165. #9'.file'#9'28'#9'"C:/Users/YARUNA~1/AppData/Local/Temp/temp.cu"'
  166. ''
  167. #9'.const .align 4 .b8 __cudart_i2opi_f[24] = {65,144,67,60,153,14' +
  168. '9,98,219,192,221,52,245,209,87,39,252,41,21,68,78,110,131,249,16' +
  169. '2};'
  170. ''
  171. #9'.entry _Z6kernelP6float4jjf ('
  172. #9#9'.param .u32 __cudaparm__Z6kernelP6float4jjf_pos,'
  173. #9#9'.param .u32 __cudaparm__Z6kernelP6float4jjf_width,'
  174. #9#9'.param .u32 __cudaparm__Z6kernelP6float4jjf_height,'
  175. #9#9'.param .f32 __cudaparm__Z6kernelP6float4jjf_time)'
  176. #9'{'
  177. #9'.reg .u16 %rh<6>;'
  178. #9'.reg .u32 %r<180>;'
  179. #9'.reg .f32 %f<92>;'
  180. #9'.reg .pred %p<26>;'
  181. #9'.local .align 4 .b8 __cuda_result_16[28];'
  182. #9'.local .align 4 .b8 __cuda_result_44[28];'
  183. #9'.loc'#9'28'#9'9'#9'0'
  184. '$LDWbegin__Z6kernelP6float4jjf:'
  185. #9'.loc'#9'18'#9'1638'#9'0'
  186. #9'mov.u16 '#9'%rh1, %ctaid.x;'
  187. #9'mov.u16 '#9'%rh2, %ntid.x;'
  188. #9'mul.wide.u16 '#9'%r1, %rh1, %rh2;'
  189. #9'ld.param.u32 '#9'%r2, [__cudaparm__Z6kernelP6float4jjf_width];'
  190. #9'cvt.rn.f32.u32 '#9'%f1, %r2;'
  191. #9'cvt.u32.u16 '#9'%r3, %tid.x;'
  192. #9'add.u32 '#9'%r4, %r3, %r1;'
  193. #9'cvt.rn.f32.u32 '#9'%f2, %r4;'
  194. #9'div.full.f32 '#9'%f3, %f2, %f1;'
  195. #9'add.f32 '#9'%f4, %f3, %f3;'
  196. #9'mov.f32 '#9'%f5, 0fbf800000; '#9'// -1'
  197. #9'add.f32 '#9'%f6, %f4, %f5;'
  198. #9'ld.param.f32 '#9'%f7, [__cudaparm__Z6kernelP6float4jjf_time];'
  199. #9'mov.f32 '#9'%f8, 0f40800000; '#9'// 4'
  200. #9'mad.f32 '#9'%f9, %f8, %f6, %f7;'
  201. #9'abs.f32 '#9'%f10, %f9;'
  202. #9'mov.f32 '#9'%f11, 0f00000000; '#9'// 0'
  203. #9'set.eq.u32.f32 '#9'%r5, %f9, %f11;'
  204. #9'neg.s32 '#9'%r6, %r5;'
  205. #9'mov.f32 '#9'%f12, 0f7f800000; '#9'// 1.#INF'
  206. #9'set.eq.u32.f32 '#9'%r7, %f10, %f12;'
  207. #9'neg.s32 '#9'%r8, %r7;'
  208. #9'or.b32 '#9'%r9, %r6, %r8;'
  209. #9'mov.u32 '#9'%r10, 0;'
  210. #9'setp.eq.s32 '#9'%p1, %r9, %r10;'
  211. #9'@%p1 bra '#9'$Lt_0_23554;'
  212. #9'.loc'#9'18'#9'1639'#9'0'
  213. #9'mov.f32 '#9'%f13, 0f00000000; '#9'// 0'
  214. #9'mul.rn.f32 '#9'%f14, %f9, %f13;'
  215. #9'mov.u32 '#9'%r11, __cudart_i2opi_f;'
  216. #9'bra.uni '#9'$LDWendi___isinff_204_5;'
  217. '$Lt_0_23554:'
  218. #9'mov.f32 '#9'%f15, 0f473ba700; '#9'// 48039'
  219. #9'setp.gt.f32 '#9'%p2, %f10, %f15;'
  220. #9'@!%p2 bra '#9'$Lt_0_24066;'
  221. #9'.loc'#9'18'#9'1396'#9'0'
  222. #9'mov.u32 '#9'%r11, __cudart_i2opi_f;'
  223. #9'mov.b32 '#9'%r12, %f9;'
  224. #9'and.b32 '#9'%r13, %r12, -2147483648;'
  225. #9'mov.s32 '#9'%r14, %r13;'
  226. #9'.loc'#9'18'#9'1405'#9'0'
  227. #9'shl.b32 '#9'%r15, %r12, 1;'
  228. #9'shr.u32 '#9'%r16, %r15, 24;'
  229. #9'sub.u32 '#9'%r17, %r16, 128;'
  230. #9'shr.u32 '#9'%r18, %r17, 5;'
  231. #9'mov.s32 '#9'%r19, 4;'
  232. #9'sub.s32 '#9'%r20, %r19, %r18;'
  233. #9'.loc'#9'18'#9'24'#9'0'
  234. #9'mov.s32 '#9'%r21, %r11;'
  235. #9'add.u32 '#9'%r22, %r11, 24;'
  236. #9'mov.u32 '#9'%r23, __cuda_result_16;'
  237. #9'shl.b32 '#9'%r24, %r12, 8;'
  238. #9'or.b32 '#9'%r25, %r24, -2147483648;'
  239. #9'mov.u32 '#9'%r26, 0;'
  240. '$Lt_0_25090:'
  241. ' //<loop> Loop body line 24, nesting depth: 1, iterations: 6'
  242. #9'.loc'#9'18'#9'1411'#9'0'
  243. #9'ld.const.u32 '#9'%r27, [%r21+0];'
  244. #9'mul.lo.u32 '#9'%r28, %r27, %r25;'
  245. #9'add.u32 '#9'%r29, %r28, %r26;'
  246. #9'.loc'#9'18'#9'1412'#9'0'
  247. #9'set.gt.u32.u32 '#9'%r30, %r28, %r29;'
  248. #9'neg.s32 '#9'%r31, %r30;'
  249. #9'mul.hi.u32 '#9'%r32, %r27, %r25;'
  250. #9'add.u32 '#9'%r26, %r31, %r32;'
  251. #9'.loc'#9'18'#9'1413'#9'0'
  252. #9'st.local.u32 '#9'[%r23+0], %r29;'
  253. #9'add.u32 '#9'%r23, %r23, 4;'
  254. #9'add.u32 '#9'%r21, %r21, 4;'
  255. #9'setp.ne.u32 '#9'%p3, %r21, %r22;'
  256. #9'@%p3 bra '#9'$Lt_0_25090;'
  257. #9'.loc'#9'18'#9'1415'#9'0'
  258. #9'st.local.u32 '#9'[__cuda_result_16+24], %r26;'
  259. #9'.loc'#9'18'#9'1420'#9'0'
  260. #9'mul.lo.u32 '#9'%r33, %r20, 4;'
  261. #9'mov.u32 '#9'%r34, __cuda_result_16;'
  262. #9'add.u32 '#9'%r35, %r33, %r34;'
  263. #9'ld.local.u32 '#9'%r26, [%r35+8];'
  264. #9'.loc'#9'18'#9'1421'#9'0'
  265. #9'ld.local.u32 '#9'%r36, [%r35+4];'
  266. #9'and.b32 '#9'%r37, %r17, 31;'
  267. #9'mov.u32 '#9'%r38, 0;'
  268. #9'setp.eq.u32 '#9'%p4, %r37, %r38;'
  269. #9'@%p4 bra '#9'$Lt_0_25602;'
  270. #9'.loc'#9'18'#9'1423'#9'0'
  271. #9'mov.s32 '#9'%r39, 32;'
  272. #9'sub.s32 '#9'%r40, %r39, %r37;'
  273. #9'.loc'#9'18'#9'1424'#9'0'
  274. #9'shr.u32 '#9'%r41, %r36, %r40;'
  275. #9'shl.b32 '#9'%r42, %r26, %r37;'
  276. #9'add.u32 '#9'%r26, %r41, %r42;'
  277. #9'.loc'#9'18'#9'1425'#9'0'
  278. #9'ld.local.u32 '#9'%r43, [%r35+0];'
  279. #9'shr.u32 '#9'%r44, %r43, %r40;'
  280. #9'shl.b32 '#9'%r45, %r36, %r37;'
  281. #9'add.u32 '#9'%r36, %r44, %r45;'
  282. '$Lt_0_25602:'
  283. #9'.loc'#9'18'#9'1427'#9'0'
  284. #9'shr.u32 '#9'%r40, %r26, 30;'
  285. #9'.loc'#9'18'#9'1429'#9'0'
  286. #9'shr.u32 '#9'%r46, %r36, 30;'
  287. #9'shl.b32 '#9'%r47, %r26, 2;'
  288. #9'add.u32 '#9'%r26, %r46, %r47;'
  289. #9'.loc'#9'18'#9'1430'#9'0'
  290. #9'shl.b32 '#9'%r36, %r36, 2;'
  291. #9'mov.u32 '#9'%r48, 0;'
  292. #9'setp.eq.u32 '#9'%p5, %r36, %r48;'
  293. #9'@%p5 bra '#9'$Lt_0_26370;'
  294. #9'.loc'#9'18'#9'1431'#9'0'
  295. #9'add.u32 '#9'%r49, %r26, 1;'
  296. #9'mov.u32 '#9'%r50, -2147483648;'
  297. #9'set.gt.u32.u32 '#9'%r51, %r49, %r50;'
  298. #9'neg.s32 '#9'%r52, %r51;'
  299. #9'bra.uni '#9'$Lt_0_26114;'
  300. '$Lt_0_26370:'
  301. #9'mov.u32 '#9'%r53, -2147483648;'
  302. #9'set.gt.u32.u32 '#9'%r54, %r26, %r53;'
  303. #9'neg.s32 '#9'%r52, %r54;'
  304. '$Lt_0_26114:'
  305. #9'.loc'#9'18'#9'1432'#9'0'
  306. #9'add.u32 '#9'%r55, %r40, %r52;'
  307. #9'.loc'#9'18'#9'1431'#9'0'
  308. #9'neg.s32 '#9'%r56, %r55;'
  309. #9'mov.u32 '#9'%r57, 0;'
  310. #9'setp.ne.u32 '#9'%p6, %r13, %r57;'
  311. #9'selp.s32 '#9'%r40, %r56, %r55, %p6;'
  312. #9'mov.u32 '#9'%r58, 0;'
  313. #9'setp.eq.u32 '#9'%p7, %r52, %r58;'
  314. #9'@%p7 bra '#9'$Lt_0_26626;'
  315. #9'.loc'#9'18'#9'1437'#9'0'
  316. #9'neg.s32 '#9'%r36, %r36;'
  317. #9'.loc'#9'18'#9'1439'#9'0'
  318. #9'mov.u32 '#9'%r59, 0;'
  319. #9'set.eq.u32.u32 '#9'%r60, %r36, %r59;'
  320. #9'neg.s32 '#9'%r61, %r60;'
  321. #9'not.b32 '#9'%r62, %r26;'
  322. #9'add.u32 '#9'%r26, %r61, %r62;'
  323. #9'.loc'#9'18'#9'1440'#9'0'
  324. #9'xor.b32 '#9'%r14, %r13, -2147483648;'
  325. '$Lt_0_26626:'
  326. #9'.loc'#9'18'#9'1442'#9'0'
  327. #9'mov.s32 '#9'%r63, %r40;'
  328. #9'mov.u32 '#9'%r64, 0;'
  329. #9'setp.le.s32 '#9'%p8, %r26, %r64;'
  330. #9'@%p8 bra '#9'$Lt_0_34818;'
  331. #9'mov.u32 '#9'%r65, 0;'
  332. '$Lt_0_27650:'
  333. ' //<loop> Loop body line 1442, nesting depth: 1, estimated itera' +
  334. 'tions: unknown'
  335. #9'.loc'#9'18'#9'1446'#9'0'
  336. #9'shr.u32 '#9'%r66, %r36, 31;'
  337. #9'shl.b32 '#9'%r67, %r26, 1;'
  338. #9'add.u32 '#9'%r26, %r66, %r67;'
  339. #9'.loc'#9'18'#9'1447'#9'0'
  340. #9'shl.b32 '#9'%r36, %r36, 1;'
  341. #9'.loc'#9'18'#9'1448'#9'0'
  342. #9'sub.u32 '#9'%r65, %r65, 1;'
  343. #9'mov.u32 '#9'%r68, 0;'
  344. #9'setp.gt.s32 '#9'%p9, %r26, %r68;'
  345. #9'@%p9 bra '#9'$Lt_0_27650;'
  346. #9'bra.uni '#9'$Lt_0_27138;'
  347. '$Lt_0_34818:'
  348. #9'mov.u32 '#9'%r65, 0;'
  349. '$Lt_0_27138:'
  350. #9'.loc'#9'18'#9'1450'#9'0'
  351. #9'mul.lo.u32 '#9'%r36, %r26, -921707870;'
  352. #9'.loc'#9'18'#9'1451'#9'0'
  353. #9'mov.u32 '#9'%r69, -921707870;'
  354. #9'mul.hi.u32 '#9'%r26, %r26, %r69;'
  355. #9'mov.u32 '#9'%r70, 0;'
  356. #9'setp.le.s32 '#9'%p10, %r26, %r70;'
  357. #9'@%p10 bra '#9'$Lt_0_28162;'
  358. #9'.loc'#9'18'#9'1453'#9'0'
  359. #9'shr.u32 '#9'%r71, %r36, 31;'
  360. #9'shl.b32 '#9'%r72, %r26, 1;'
  361. #9'add.u32 '#9'%r26, %r71, %r72;'
  362. #9'.loc'#9'18'#9'1454'#9'0'
  363. #9'shl.b32 '#9'%r36, %r36, 1;'
  364. #9'.loc'#9'18'#9'1455'#9'0'
  365. #9'sub.u32 '#9'%r65, %r65, 1;'
  366. '$Lt_0_28162:'
  367. #9'.loc'#9'18'#9'1457'#9'0'
  368. #9'mov.u32 '#9'%r73, 0;'
  369. #9'set.ne.u32.u32 '#9'%r74, %r36, %r73;'
  370. #9'neg.s32 '#9'%r75, %r74;'
  371. #9'add.u32 '#9'%r26, %r75, %r26;'
  372. #9'.loc'#9'18'#9'1459'#9'0'
  373. #9'shl.b32 '#9'%r76, %r26, 24;'
  374. #9'mov.s32 '#9'%r77, 0;'
  375. #9'set.lt.u32.s32 '#9'%r78, %r76, %r77;'
  376. #9'neg.s32 '#9'%r79, %r78;'
  377. #9'shr.u32 '#9'%r80, %r26, 8;'
  378. #9'add.u32 '#9'%r81, %r65, 126;'
  379. #9'shl.b32 '#9'%r82, %r81, 23;'
  380. #9'add.u32 '#9'%r83, %r80, %r82;'
  381. #9'add.u32 '#9'%r84, %r79, %r83;'
  382. #9'or.b32 '#9'%r85, %r14, %r84;'
  383. #9'mov.b32 '#9'%f16, %r85;'
  384. #9'bra.uni '#9'$LDWendi___internal_fmad_204_6;'
  385. '$Lt_0_24066:'
  386. #9'.loc'#9'18'#9'1463'#9'0'
  387. #9'mov.f32 '#9'%f17, 0f3f22f983; '#9'// 0.63662'
  388. #9'mul.f32 '#9'%f18, %f9, %f17;'
  389. #9'cvt.rni.s32.f32 '#9'%r86, %f18;'
  390. #9'cvt.rn.f32.s32 '#9'%f19, %r86;'
  391. #9'neg.f32 '#9'%f20, %f19;'
  392. #9'.loc'#9'18'#9'1472'#9'0'
  393. #9'mov.s32 '#9'%r63, %r86;'
  394. #9'.loc'#9'18'#9'1473'#9'0'
  395. #9'mov.f32 '#9'%f21, 0f3fc90000; '#9'// 1.57031'
  396. #9'mad.f32 '#9'%f22, %f21, %f20, %f9;'
  397. #9'mov.f32 '#9'%f23, 0f39fd8000; '#9'// 0.000483513'
  398. #9'mad.f32 '#9'%f24, %f23, %f20, %f22;'
  399. #9'mov.f32 '#9'%f25, 0f34a88000; '#9'// 3.13856e-007'
  400. #9'mad.f32 '#9'%f26, %f25, %f20, %f24;'
  401. #9'mov.f32 '#9'%f27, 0f2e85a309; '#9'// 6.0771e-011'
  402. #9'mad.f32 '#9'%f16, %f27, %f20, %f26;'
  403. #9'mov.u32 '#9'%r11, __cudart_i2opi_f;'
  404. '$LDWendi___internal_fmad_204_6:'
  405. #9'.loc'#9'18'#9'1641'#9'0'
  406. #9'mul.f32 '#9'%f28, %f16, %f16;'
  407. #9'and.b32 '#9'%r87, %r63, 1;'
  408. #9'mov.u32 '#9'%r88, 0;'
  409. #9'setp.eq.s32 '#9'%p11, %r87, %r88;'
  410. #9'@%p11 bra '#9'$Lt_0_28930;'
  411. #9'.loc'#9'18'#9'1644'#9'0'
  412. #9'mov.f32 '#9'%f29, 0f3f800000; '#9'// 1'
  413. #9'mov.f32 '#9'%f30, 0fbf000000; '#9'// -0.5'
  414. #9'mov.f32 '#9'%f31, 0f3d2aaaa5; '#9'// 0.0416666'
  415. #9'mov.f32 '#9'%f32, 0fbab6061a; '#9'// -0.00138873'
  416. #9'mov.f32 '#9'%f33, 0f37ccf5ce; '#9'// 2.44332e-005'
  417. #9'mad.f32 '#9'%f34, %f33, %f28, %f32;'
  418. #9'mad.f32 '#9'%f35, %f28, %f34, %f31;'
  419. #9'mad.f32 '#9'%f36, %f28, %f35, %f30;'
  420. #9'mad.f32 '#9'%f37, %f28, %f36, %f29;'
  421. #9'bra.uni '#9'$Lt_0_28674;'
  422. '$Lt_0_28930:'
  423. #9'.loc'#9'18'#9'1646'#9'0'
  424. #9'mov.f32 '#9'%f38, 0fbe2aaaa3; '#9'// -0.166667'
  425. #9'mov.f32 '#9'%f39, 0f3c08839e; '#9'// 0.00833216'
  426. #9'mov.f32 '#9'%f40, 0fb94ca1f9; '#9'// -0.000195153'
  427. #9'mad.f32 '#9'%f41, %f40, %f28, %f39;'
  428. #9'mad.f32 '#9'%f42, %f28, %f41, %f38;'
  429. #9'mul.f32 '#9'%f43, %f28, %f42;'
  430. #9'mad.f32 '#9'%f37, %f43, %f16, %f16;'
  431. '$Lt_0_28674:'
  432. #9'.loc'#9'18'#9'1648'#9'0'
  433. #9'neg.f32 '#9'%f44, %f37;'
  434. #9'and.b32 '#9'%r89, %r63, 2;'
  435. #9'mov.s32 '#9'%r90, 0;'
  436. #9'setp.ne.s32 '#9'%p12, %r89, %r90;'
  437. #9'selp.f32 '#9'%f37, %f44, %f37, %p12;'
  438. #9'.loc'#9'18'#9'1651'#9'0'
  439. #9'mov.f32 '#9'%f14, %f37;'
  440. '$LDWendi___isinff_204_5:'
  441. #9'.loc'#9'18'#9'1702'#9'0'
  442. #9'mov.u16 '#9'%rh3, %ctaid.y;'
  443. #9'mov.u16 '#9'%rh4, %ntid.y;'
  444. #9'mul.wide.u16 '#9'%r91, %rh3, %rh4;'
  445. #9'ld.param.u32 '#9'%r92, [__cudaparm__Z6kernelP6float4jjf_height];'
  446. #9'cvt.rn.f32.u32 '#9'%f45, %r92;'
  447. #9'cvt.u32.u16 '#9'%r93, %tid.y;'
  448. #9'add.u32 '#9'%r94, %r93, %r91;'
  449. #9'cvt.rn.f32.u32 '#9'%f46, %r94;'
  450. #9'div.full.f32 '#9'%f47, %f46, %f45;'
  451. #9'add.f32 '#9'%f48, %f47, %f47;'
  452. #9'mov.f32 '#9'%f49, 0fbf800000; '#9'// -1'
  453. #9'add.f32 '#9'%f50, %f48, %f49;'
  454. #9'mov.f32 '#9'%f51, 0f40800000; '#9'// 4'
  455. #9'.loc'#9'18'#9'1638'#9'0'
  456. #9'ld.param.f32 '#9'%f7, [__cudaparm__Z6kernelP6float4jjf_time];'
  457. #9'.loc'#9'18'#9'1702'#9'0'
  458. #9'mad.f32 '#9'%f52, %f51, %f50, %f7;'
  459. #9'abs.f32 '#9'%f53, %f52;'
  460. #9'mov.f32 '#9'%f54, 0f7f800000; '#9'// 1.#INF'
  461. #9'setp.eq.f32 '#9'%p13, %f53, %f54;'
  462. #9'@!%p13 bra '#9'$Lt_0_29186;'
  463. #9'.loc'#9'18'#9'1703'#9'0'
  464. #9'neg.f32 '#9'%f55, %f52;'
  465. #9'add.rn.f32 '#9'%f56, %f52, %f55;'
  466. #9'bra.uni '#9'$LDWendi___isinff_204_1;'
  467. '$Lt_0_29186:'
  468. #9'mov.f32 '#9'%f57, 0f473ba700; '#9'// 48039'
  469. #9'setp.gt.f32 '#9'%p14, %f53, %f57;'
  470. #9'@!%p14 bra '#9'$Lt_0_29698;'
  471. #9'.loc'#9'18'#9'1396'#9'0'
  472. #9'mov.b32 '#9'%r95, %f52;'
  473. #9'and.b32 '#9'%r96, %r95, -2147483648;'
  474. #9'mov.s32 '#9'%r97, %r96;'
  475. #9'.loc'#9'18'#9'1405'#9'0'
  476. #9'shl.b32 '#9'%r98, %r95, 1;'
  477. #9'shr.u32 '#9'%r99, %r98, 24;'
  478. #9'sub.u32 '#9'%r100, %r99, 128;'
  479. #9'shr.u32 '#9'%r101, %r100, 5;'
  480. #9'mov.s32 '#9'%r102, 4;'
  481. #9'sub.s32 '#9'%r103, %r102, %r101;'
  482. #9'.loc'#9'18'#9'24'#9'0'
  483. #9'mov.s32 '#9'%r104, %r11;'
  484. #9'add.u32 '#9'%r22, %r11, 24;'
  485. #9'mov.u32 '#9'%r105, __cuda_result_44;'
  486. #9'shl.b32 '#9'%r106, %r95, 8;'
  487. #9'or.b32 '#9'%r107, %r106, -2147483648;'
  488. #9'mov.u32 '#9'%r108, 0;'
  489. '$Lt_0_30722:'
  490. ' //<loop> Loop body line 24, nesting depth: 1, iterations: 6'
  491. #9'.loc'#9'18'#9'1411'#9'0'
  492. #9'ld.const.u32 '#9'%r109, [%r104+0];'
  493. #9'mul.lo.u32 '#9'%r110, %r109, %r107;'
  494. #9'add.u32 '#9'%r111, %r110, %r108;'
  495. #9'.loc'#9'18'#9'1412'#9'0'
  496. #9'set.gt.u32.u32 '#9'%r112, %r110, %r111;'
  497. #9'neg.s32 '#9'%r113, %r112;'
  498. #9'mul.hi.u32 '#9'%r114, %r109, %r107;'
  499. #9'add.u32 '#9'%r108, %r113, %r114;'
  500. #9'.loc'#9'18'#9'1413'#9'0'
  501. #9'st.local.u32 '#9'[%r105+0], %r111;'
  502. #9'add.u32 '#9'%r105, %r105, 4;'
  503. #9'add.u32 '#9'%r104, %r104, 4;'
  504. #9'setp.ne.u32 '#9'%p15, %r104, %r22;'
  505. #9'@%p15 bra '#9'$Lt_0_30722;'
  506. #9'.loc'#9'18'#9'1415'#9'0'
  507. #9'st.local.u32 '#9'[__cuda_result_44+24], %r108;'
  508. #9'.loc'#9'18'#9'1420'#9'0'
  509. #9'mul.lo.u32 '#9'%r115, %r103, 4;'
  510. #9'mov.u32 '#9'%r116, __cuda_result_44;'
  511. #9'add.u32 '#9'%r117, %r115, %r116;'
  512. #9'ld.local.u32 '#9'%r108, [%r117+8];'
  513. #9'.loc'#9'18'#9'1421'#9'0'
  514. #9'ld.local.u32 '#9'%r118, [%r117+4];'
  515. #9'and.b32 '#9'%r119, %r100, 31;'
  516. #9'mov.u32 '#9'%r120, 0;'
  517. #9'setp.eq.u32 '#9'%p16, %r119, %r120;'
  518. #9'@%p16 bra '#9'$Lt_0_31234;'
  519. #9'.loc'#9'18'#9'1423'#9'0'
  520. #9'mov.s32 '#9'%r121, 32;'
  521. #9'sub.s32 '#9'%r122, %r121, %r119;'
  522. #9'.loc'#9'18'#9'1424'#9'0'
  523. #9'shr.u32 '#9'%r123, %r118, %r122;'
  524. #9'shl.b32 '#9'%r124, %r108, %r119;'
  525. #9'add.u32 '#9'%r108, %r123, %r124;'
  526. #9'.loc'#9'18'#9'1425'#9'0'
  527. #9'ld.local.u32 '#9'%r125, [%r117+0];'
  528. #9'shr.u32 '#9'%r126, %r125, %r122;'
  529. #9'shl.b32 '#9'%r127, %r118, %r119;'
  530. #9'add.u32 '#9'%r118, %r126, %r127;'
  531. '$Lt_0_31234:'
  532. #9'.loc'#9'18'#9'1427'#9'0'
  533. #9'shr.u32 '#9'%r122, %r108, 30;'
  534. #9'.loc'#9'18'#9'1429'#9'0'
  535. #9'shr.u32 '#9'%r128, %r118, 30;'
  536. #9'shl.b32 '#9'%r129, %r108, 2;'
  537. #9'add.u32 '#9'%r108, %r128, %r129;'
  538. #9'.loc'#9'18'#9'1430'#9'0'
  539. #9'shl.b32 '#9'%r118, %r118, 2;'
  540. #9'mov.u32 '#9'%r130, 0;'
  541. #9'setp.eq.u32 '#9'%p17, %r118, %r130;'
  542. #9'@%p17 bra '#9'$Lt_0_32002;'
  543. #9'.loc'#9'18'#9'1431'#9'0'
  544. #9'add.u32 '#9'%r131, %r108, 1;'
  545. #9'mov.u32 '#9'%r132, -2147483648;'
  546. #9'set.gt.u32.u32 '#9'%r133, %r131, %r132;'
  547. #9'neg.s32 '#9'%r134, %r133;'
  548. #9'bra.uni '#9'$Lt_0_31746;'
  549. '$Lt_0_32002:'
  550. #9'mov.u32 '#9'%r135, -2147483648;'
  551. #9'set.gt.u32.u32 '#9'%r136, %r108, %r135;'
  552. #9'neg.s32 '#9'%r134, %r136;'
  553. '$Lt_0_31746:'
  554. #9'.loc'#9'18'#9'1432'#9'0'
  555. #9'add.u32 '#9'%r137, %r122, %r134;'
  556. #9'.loc'#9'18'#9'1431'#9'0'
  557. #9'neg.s32 '#9'%r138, %r137;'
  558. #9'mov.u32 '#9'%r139, 0;'
  559. #9'setp.ne.u32 '#9'%p18, %r96, %r139;'
  560. #9'selp.s32 '#9'%r122, %r138, %r137, %p18;'
  561. #9'mov.u32 '#9'%r140, 0;'
  562. #9'setp.eq.u32 '#9'%p19, %r134, %r140;'
  563. #9'@%p19 bra '#9'$Lt_0_32258;'
  564. #9'.loc'#9'18'#9'1437'#9'0'
  565. #9'neg.s32 '#9'%r118, %r118;'
  566. #9'.loc'#9'18'#9'1439'#9'0'
  567. #9'mov.u32 '#9'%r141, 0;'
  568. #9'set.eq.u32.u32 '#9'%r142, %r118, %r141;'
  569. #9'neg.s32 '#9'%r143, %r142;'
  570. #9'not.b32 '#9'%r144, %r108;'
  571. #9'add.u32 '#9'%r108, %r143, %r144;'
  572. #9'.loc'#9'18'#9'1440'#9'0'
  573. #9'xor.b32 '#9'%r97, %r96, -2147483648;'
  574. '$Lt_0_32258:'
  575. #9'.loc'#9'18'#9'1442'#9'0'
  576. #9'mov.s32 '#9'%r145, %r122;'
  577. #9'mov.u32 '#9'%r146, 0;'
  578. #9'setp.le.s32 '#9'%p20, %r108, %r146;'
  579. #9'@%p20 bra '#9'$Lt_0_35074;'
  580. #9'mov.u32 '#9'%r147, 0;'
  581. '$Lt_0_33282:'
  582. ' //<loop> Loop body line 1442, nesting depth: 1, estimated itera' +
  583. 'tions: unknown'
  584. #9'.loc'#9'18'#9'1446'#9'0'
  585. #9'shr.u32 '#9'%r148, %r118, 31;'
  586. #9'shl.b32 '#9'%r149, %r108, 1;'
  587. #9'add.u32 '#9'%r108, %r148, %r149;'
  588. #9'.loc'#9'18'#9'1447'#9'0'
  589. #9'shl.b32 '#9'%r118, %r118, 1;'
  590. #9'.loc'#9'18'#9'1448'#9'0'
  591. #9'sub.u32 '#9'%r147, %r147, 1;'
  592. #9'mov.u32 '#9'%r150, 0;'
  593. #9'setp.gt.s32 '#9'%p21, %r108, %r150;'
  594. #9'@%p21 bra '#9'$Lt_0_33282;'
  595. #9'bra.uni '#9'$Lt_0_32770;'
  596. '$Lt_0_35074:'
  597. #9'mov.u32 '#9'%r147, 0;'
  598. '$Lt_0_32770:'
  599. #9'.loc'#9'18'#9'1450'#9'0'
  600. #9'mul.lo.u32 '#9'%r118, %r108, -921707870;'
  601. #9'.loc'#9'18'#9'1451'#9'0'
  602. #9'mov.u32 '#9'%r151, -921707870;'
  603. #9'mul.hi.u32 '#9'%r108, %r108, %r151;'
  604. #9'mov.u32 '#9'%r152, 0;'
  605. #9'setp.le.s32 '#9'%p22, %r108, %r152;'
  606. #9'@%p22 bra '#9'$Lt_0_33794;'
  607. #9'.loc'#9'18'#9'1453'#9'0'
  608. #9'shr.u32 '#9'%r153, %r118, 31;'
  609. #9'shl.b32 '#9'%r154, %r108, 1;'
  610. #9'add.u32 '#9'%r108, %r153, %r154;'
  611. #9'.loc'#9'18'#9'1454'#9'0'
  612. #9'shl.b32 '#9'%r118, %r118, 1;'
  613. #9'.loc'#9'18'#9'1455'#9'0'
  614. #9'sub.u32 '#9'%r147, %r147, 1;'
  615. '$Lt_0_33794:'
  616. #9'.loc'#9'18'#9'1457'#9'0'
  617. #9'mov.u32 '#9'%r155, 0;'
  618. #9'set.ne.u32.u32 '#9'%r156, %r118, %r155;'
  619. #9'neg.s32 '#9'%r157, %r156;'
  620. #9'add.u32 '#9'%r108, %r157, %r108;'
  621. #9'.loc'#9'18'#9'1459'#9'0'
  622. #9'shl.b32 '#9'%r158, %r108, 24;'
  623. #9'mov.s32 '#9'%r159, 0;'
  624. #9'set.lt.u32.s32 '#9'%r160, %r158, %r159;'
  625. #9'neg.s32 '#9'%r161, %r160;'
  626. #9'shr.u32 '#9'%r162, %r108, 8;'
  627. #9'add.u32 '#9'%r163, %r147, 126;'
  628. #9'shl.b32 '#9'%r164, %r163, 23;'
  629. #9'add.u32 '#9'%r165, %r162, %r164;'
  630. #9'add.u32 '#9'%r166, %r161, %r165;'
  631. #9'or.b32 '#9'%r167, %r97, %r166;'
  632. #9'mov.b32 '#9'%f58, %r167;'
  633. #9'bra.uni '#9'$LDWendi___internal_fmad_204_2;'
  634. '$Lt_0_29698:'
  635. #9'.loc'#9'18'#9'1463'#9'0'
  636. #9'mov.f32 '#9'%f59, 0f3f22f983; '#9'// 0.63662'
  637. #9'mul.f32 '#9'%f60, %f52, %f59;'
  638. #9'cvt.rni.s32.f32 '#9'%r168, %f60;'
  639. #9'cvt.rn.f32.s32 '#9'%f61, %r168;'
  640. #9'neg.f32 '#9'%f62, %f61;'
  641. #9'.loc'#9'18'#9'1472'#9'0'
  642. #9'mov.s32 '#9'%r145, %r168;'
  643. #9'.loc'#9'18'#9'1473'#9'0'
  644. #9'mov.f32 '#9'%f63, 0f3fc90000; '#9'// 1.57031'
  645. #9'mad.f32 '#9'%f64, %f63, %f62, %f52;'
  646. #9'mov.f32 '#9'%f65, 0f39fd8000; '#9'// 0.000483513'
  647. #9'mad.f32 '#9'%f66, %f65, %f62, %f64;'
  648. #9'mov.f32 '#9'%f67, 0f34a88000; '#9'// 3.13856e-007'
  649. #9'mad.f32 '#9'%f68, %f67, %f62, %f66;'
  650. #9'mov.f32 '#9'%f69, 0f2e85a309; '#9'// 6.0771e-011'
  651. #9'mad.f32 '#9'%f58, %f69, %f62, %f68;'
  652. '$LDWendi___internal_fmad_204_2:'
  653. #9'.loc'#9'18'#9'1705'#9'0'
  654. #9'add.s32 '#9'%r169, %r145, 1;'
  655. #9'mul.f32 '#9'%f70, %f58, %f58;'
  656. #9'and.b32 '#9'%r170, %r169, 1;'
  657. #9'mov.u32 '#9'%r171, 0;'
  658. #9'setp.eq.s32 '#9'%p23, %r170, %r171;'
  659. #9'@%p23 bra '#9'$Lt_0_34562;'
  660. #9'.loc'#9'18'#9'1709'#9'0'
  661. #9'mov.f32 '#9'%f71, 0f3f800000; '#9'// 1'
  662. #9'mov.f32 '#9'%f72, 0fbf000000; '#9'// -0.5'
  663. #9'mov.f32 '#9'%f73, 0f3d2aaaa5; '#9'// 0.0416666'
  664. #9'mov.f32 '#9'%f74, 0fbab6061a; '#9'// -0.00138873'
  665. #9'mov.f32 '#9'%f75, 0f37ccf5ce; '#9'// 2.44332e-005'
  666. #9'mad.f32 '#9'%f76, %f75, %f70, %f74;'
  667. #9'mad.f32 '#9'%f77, %f70, %f76, %f73;'
  668. #9'mad.f32 '#9'%f78, %f70, %f77, %f72;'
  669. #9'mad.f32 '#9'%f79, %f70, %f78, %f71;'
  670. #9'bra.uni '#9'$Lt_0_34306;'
  671. '$Lt_0_34562:'
  672. #9'.loc'#9'18'#9'1711'#9'0'
  673. #9'mov.f32 '#9'%f80, 0fbe2aaaa3; '#9'// -0.166667'
  674. #9'mov.f32 '#9'%f81, 0f3c08839e; '#9'// 0.00833216'
  675. #9'mov.f32 '#9'%f82, 0fb94ca1f9; '#9'// -0.000195153'
  676. #9'mad.f32 '#9'%f83, %f82, %f70, %f81;'
  677. #9'mad.f32 '#9'%f84, %f70, %f83, %f80;'
  678. #9'mul.f32 '#9'%f85, %f70, %f84;'
  679. #9'mad.f32 '#9'%f79, %f85, %f58, %f58;'
  680. '$Lt_0_34306:'
  681. #9'.loc'#9'18'#9'1713'#9'0'
  682. #9'neg.f32 '#9'%f86, %f79;'
  683. #9'and.b32 '#9'%r172, %r169, 2;'
  684. #9'mov.s32 '#9'%r173, 0;'
  685. #9'setp.ne.s32 '#9'%p24, %r172, %r173;'
  686. #9'selp.f32 '#9'%f79, %f86, %f79, %p24;'
  687. #9'.loc'#9'18'#9'1716'#9'0'
  688. #9'mov.f32 '#9'%f56, %f79;'
  689. '$LDWendi___isinff_204_1:'
  690. #9'.loc'#9'18'#9'1638'#9'0'
  691. #9'ld.param.u32 '#9'%r2, [__cudaparm__Z6kernelP6float4jjf_width];'
  692. #9'.loc'#9'28'#9'25'#9'0'
  693. #9'mul.lo.u32 '#9'%r174, %r94, %r2;'
  694. #9'add.u32 '#9'%r175, %r4, %r174;'
  695. #9'mul.lo.u32 '#9'%r176, %r175, 16;'
  696. #9'ld.param.u32 '#9'%r177, [__cudaparm__Z6kernelP6float4jjf_pos];'
  697. #9'add.u32 '#9'%r178, %r177, %r176;'
  698. #9'mul.f32 '#9'%f87, %f56, %f14;'
  699. #9'mov.f32 '#9'%f88, 0f3f000000; '#9'// 0.5'
  700. #9'mul.f32 '#9'%f89, %f87, %f88;'
  701. #9'mov.f32 '#9'%f90, 0f3f800000; '#9'// 1'
  702. #9'st.global.v4.f32 '#9'[%r178+0], {%f6,%f89,%f50,%f90};'
  703. #9'.loc'#9'28'#9'26'#9'0'
  704. #9'exit;'
  705. '$LDWend__Z6kernelP6float4jjf:'
  706. #9'} // _Z6kernelP6float4jjf'
  707. '')
  708. Compiler = GLCUDACompiler1
  709. object MakeDotField: TCUDAFunction
  710. KernelName = '_Z6kernelP6float4jjf'
  711. BlockShape.SizeX = 8
  712. BlockShape.SizeY = 8
  713. OnParameterSetup = MakeVertexBufferParameterSetup
  714. end
  715. end
  716. object DotFieldMapper: TCUDAGeometryResource
  717. FeedBackMesh = GLFeedBackMesh1
  718. Mapping = grmWriteDiscard
  719. Left = 248
  720. Top = 264
  721. end
  722. end
  723. object GLCUDACompiler1: TGLCUDACompiler
  724. NVCCPath = 'C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v3.2\bin\'
  725. CppCompilerPath = 'C:\Program Files (x86)\Microsoft Visual Studio 12.0\VC\bin\'
  726. ProjectModule = 'Simple kernel.cu'
  727. Left = 448
  728. Top = 128
  729. end
  730. object GLSLShader1: TGLSLShader
  731. FragmentProgram.Code.Strings = (
  732. '#version 330'
  733. 'out vec4 FragColor;'
  734. 'void main(void)'
  735. '{'
  736. ' FragColor = vec4(1.0, 0.0, 0.0, 1.0);'
  737. '};')
  738. FragmentProgram.Enabled = True
  739. VertexProgram.Code.Strings = (
  740. '#version 330'
  741. 'in vec4 Position;'
  742. 'uniform mat4 ModelViewProjectionMatrix;'
  743. 'void main()'
  744. '{'
  745. ' gl_Position = ModelViewProjectionMatrix * Position;'
  746. '};')
  747. VertexProgram.Enabled = True
  748. OnApply = GLSLShader1Apply
  749. FailedInitAction = fiaSilentDisable
  750. Left = 24
  751. Top = 128
  752. end
  753. end