wgan example failed to run on H100 with cuda backend

[npuichigo]

0.17.0 version
H100 with CUDA 12.1

    Finished `release` profile [optimized] target(s) in 3.77s
     Running `/home/yuchao/workspace/burn/target/release/examples/wgan-mnist`

thread '<unnamed>' panicked at /home/yuchao/.cargo/registry/src/index.crates.io-1949cf8c6b5b557f/cubecl-cuda-0.5.0/src/compute/server.rs:682:17:
[Compilation Error] 
    default_program(36): warning #550-D: variable "l_arr_0" was set but never used
      uint32 l_arr_0[4];
             ^
    Remark: The warnings can be suppressed with "-diag-suppress <warning-number>"
    default_program(27): Error: Formal parameter space overflowed (36800 bytes required, max 32764 bytes allowed) in function elemwise_fuse
[Source]  
#include <cuda_runtime.h>
typedef unsigned int uint;
typedef unsigned char uint8;
typedef unsigned short uint16;
typedef unsigned int uint32;
typedef unsigned long long int uint64;
typedef signed char int8;
typedef signed short int16;
typedef signed int int32;
typedef signed long long int int64;

struct __align__(16) float_4 {
    float i_0;
    float i_1;
    float i_2;
    float i_3;
};

...

buffer_1018[idxGlobal] = reinterpret_cast< float_4 const&>(l_18917);
const float_4 l_18918 = l_18337;
buffer_1019[idxGlobal] = reinterpret_cast< float_4 const&>(l_18918);
const float_4 l_18919 = l_18373;
buffer_1020[idxGlobal] = reinterpret_cast< float_4 const&>(l_18919);
const float_4 l_18920 = l_18409;
buffer_1021[idxGlobal] = reinterpret_cast< float_4 const&>(l_18920);
}

}

thread 'main' panicked at /home/yuchao/workspace/burn/crates/burn-core/src/data/dataloader/multithread.rs:231:29:
called `Result::unwrap()` on an `Err` value: RecvError

[amanPradhan05]

hey, in which branch are you getting this error?

[npuichigo]

Thanks. I just use the main branch.

[npuichigo]

Same error on 0.17.0

[laggui]

Thanks for the bug report! Will take a look.

/edit: this is clearly a fusion bug

default_program(27): Error: Formal parameter space overflowed (36800 bytes required, max 32764 bytes allowed) in function elemwise_fuse

If you disable the default features (which includes fusion) it should work.

[dependencies]
-burn = { path = "../../crates/burn", features=["train", "vision"] }
+burn = { path = "../../crates/burn", features=["train", "vision", "std"], default-features = false }
image = { workspace = true }

Still a bug though 😅

[laggui]

Looks like we maxed out the parameter space with element-wise fusion.

elemwise_fuse def

extern "C" __global__ void elemwise_fuse (
    float_4 buffer_0[], float_4 buffer_1[], float_4 buffer_2[], float_4 buffer_3[], 
float_4 buffer_4[], float_4 buffer_5[], float_4 buffer_6[], float_4 buffer_7[], float_4 
buffer_8[], float_4 buffer_9[], float_4 buffer_10[], float_4 buffer_11[], float_4 
buffer_12[], float_4 buffer_13[], float_4 buffer_14[], float_4 buffer_15[], float_4 
buffer_16[], float_4 buffer_17[], float_4 buffer_18[], float_4 buffer_19[], float_4 
buffer_20[], float_4 buffer_21[], float_4 buffer_22[], float_4 buffer_23[], float_4 
buffer_24[], float_4 buffer_25[], float_4 buffer_26[], float_4 buffer_27[], float_4 
buffer_28[], float_4 buffer_29[], float_4 buffer_30[], float_4 buffer_31[], float_4 
buffer_32[], float_4 buffer_33[], float_4 buffer_34[], float_4 buffer_35[], float_4 
buffer_36[], float_4 buffer_37[], float_4 buffer_38[], float_4 buffer_39[], float_4 
buffer_40[], float_4 buffer_41[], float_4 buffer_42[], float_4 buffer_43[], float_4 
buffer_44[], float_4 buffer_45[], float_4 buffer_46[], float_4 buffer_47[], float_4 
buffer_48[], float_4 buffer_49[], float_4 buffer_50[], float_4 buffer_51[], float_4 
buffer_52[], float_4 buffer_53[], float_4 buffer_54[], float_4 buffer_55[], float_4 
buffer_56[], float_4 buffer_57[], float_4 buffer_58[], float_4 buffer_59[], float_4 
buffer_60[], float_4 buffer_61[], float_4 buffer_62[], float_4 buffer_63[], float_4 
buffer_64[], float_4 buffer_65[], float_4 buffer_66[], float_4 buffer_67[], float_4 
buffer_68[], float_4 buffer_69[], float_4 buffer_70[], float_4 buffer_71[], float_4 
buffer_72[], float_4 buffer_73[], float_4 buffer_74[], float_4 buffer_75[], float_4 
buffer_76[], float_4 buffer_77[], float_4 buffer_78[], float_4 buffer_79[], float_4 
buffer_80[], float_4 buffer_81[], float_4 buffer_82[], float_4 buffer_83[], float_4 
buffer_84[], float_4 buffer_85[], float_4 buffer_86[], float_4 buffer_87[], float_4 
buffer_88[], float_4 buffer_89[], float_4 buffer_90[], float_4 buffer_91[], float_4 
buffer_92[], float_4 buffer_93[], float_4 buffer_94[], float_4 buffer_95[], float_4 
buffer_96[], float_4 buffer_97[], float_4 buffer_98[], float_4 buffer_99[], float_4 
buffer_100[], float_4 buffer_101[], float_4 buffer_102[], float_4 buffer_103[], float_4 
buffer_104[], float_4 buffer_105[], float_4 buffer_106[], float_4 buffer_107[], float_4 
buffer_108[], float_4 buffer_109[], float_4 buffer_110[], float_4 buffer_111[], float_4 
buffer_112[], float_4 buffer_113[], float_4 buffer_114[], float_4 buffer_115[], float_4 
buffer_116[], float_4 buffer_117[], float_4 buffer_118[], float_4 buffer_119[], float_4 
buffer_120[], float_4 buffer_121[], float_4 buffer_122[], float_4 buffer_123[], float_4 
buffer_124[], float_4 buffer_125[], float_4 buffer_126[], float_4 buffer_127[], float_4 
buffer_128[], float_4 buffer_129[], float_4 buffer_130[], float_4 buffer_131[], float_4 
buffer_132[], float_4 buffer_133[], float_4 buffer_134[], float_4 buffer_135[], float_4 
buffer_136[], float_4 buffer_137[], float_4 buffer_138[], float_4 buffer_139[], float_4 
buffer_140[], float_4 buffer_141[], float_4 buffer_142[], float_4 buffer_143[], float_4 
buffer_144[], float_4 buffer_145[], float_4 buffer_146[], float_4 buffer_147[], float_4 
buffer_148[], float_4 buffer_149[], float_4 buffer_150[], float_4 buffer_151[], float_4 
buffer_152[], float_4 buffer_153[], float_4 buffer_154[], float_4 buffer_155[], float_4 
buffer_156[], float_4 buffer_157[], float_4 buffer_158[], float_4 buffer_159[], float_4 
buffer_160[], float_4 buffer_161[], float_4 buffer_162[], float_4 buffer_163[], float_4 
buffer_164[], float_4 buffer_165[], float_4 buffer_166[], float_4 buffer_167[], float_4 
buffer_168[], float_4 buffer_169[], float_4 buffer_170[], float_4 buffer_171[], float_4 
buffer_172[], float_4 buffer_173[], float_4 buffer_174[], float_4 buffer_175[], float_4 
buffer_176[], float_4 buffer_177[], float_4 buffer_178[], float_4 buffer_179[], float_4 
buffer_180[], float_4 buffer_181[], float_4 buffer_182[], float_4 buffer_183[], float_4 
buffer_184[], float_4 buffer_185[], float_4 buffer_186[], float_4 buffer_187[], float_4 
buffer_188[], float_4 buffer_189[], float_4 buffer_190[], float_4 buffer_191[], float_4 
buffer_192[], float_4 buffer_193[], float_4 buffer_194[], float_4 buffer_195[], float_4 
buffer_196[], float_4 buffer_197[], float_4 buffer_198[], float_4 buffer_199[], float_4 
buffer_200[], float_4 buffer_201[], float_4 buffer_202[], float_4 buffer_203[], float_4 
buffer_204[], float_4 buffer_205[], float_4 buffer_206[], float_4 buffer_207[], float_4 
buffer_208[], float_4 buffer_209[], float_4 buffer_210[], float_4 buffer_211[], float_4 
buffer_212[], float_4 buffer_213[], float_4 buffer_214[], float_4 buffer_215[], float_4 
buffer_216[], float_4 buffer_217[], float_4 buffer_218[], float_4 buffer_219[], float_4 
buffer_220[], float_4 buffer_221[], float_4 buffer_222[], float_4 buffer_223[], float_4 
buffer_224[], float_4 buffer_225[], float_4 buffer_226[], float_4 buffer_227[], float_4 
buffer_228[], float_4 buffer_229[], float_4 buffer_230[], float_4 buffer_231[], float_4 
buffer_232[], float_4 buffer_233[], float_4 buffer_234[], float_4 buffer_235[], float_4 
buffer_236[], float_4 buffer_237[], float_4 buffer_238[], float_4 buffer_239[], float_4 
buffer_240[], float_4 buffer_241[], float_4 buffer_242[], float_4 buffer_243[], float_4 
buffer_244[], float_4 buffer_245[], float_4 buffer_246[], float_4 buffer_247[], float_4 
buffer_248[], float_4 buffer_249[], float_4 buffer_250[], float_4 buffer_251[], float_4 
buffer_252[], float_4 buffer_253[], float_4 buffer_254[], float_4 buffer_255[], float_4 
buffer_256[], float_4 buffer_257[], float_4 buffer_258[], float_4 buffer_259[], float_4 
buffer_260[], float_4 buffer_261[], float_4 buffer_262[], float_4 buffer_263[], float_4 
buffer_264[], float_4 buffer_265[], float_4 buffer_266[], float_4 buffer_267[], float_4 
buffer_268[], float_4 buffer_269[], float_4 buffer_270[], float_4 buffer_271[], float_4 
buffer_272[], float_4 buffer_273[], float_4 buffer_274[], float_4 buffer_275[], float_4 
buffer_276[], float_4 buffer_277[], float_4 buffer_278[], float_4 buffer_279[], float_4 
buffer_280[], float_4 buffer_281[], float_4 buffer_282[], float_4 buffer_283[], float_4 
buffer_284[], float_4 buffer_285[], float_4 buffer_286[], float_4 buffer_287[], float_4 
buffer_288[], float_4 buffer_289[], float_4 buffer_290[], float_4 buffer_291[], float_4 
buffer_292[], float_4 buffer_293[], float_4 buffer_294[], float_4 buffer_295[], float_4 
buffer_296[], float_4 buffer_297[], float_4 buffer_298[], float_4 buffer_299[], float_4 
buffer_300[], float_4 buffer_301[], float_4 buffer_302[], float_4 buffer_303[], float_4 
buffer_304[], float_4 buffer_305[], float_4 buffer_306[], float_4 buffer_307[], float_4 
buffer_308[], float_4 buffer_309[], float_4 buffer_310[], float_4 buffer_311[], float_4 
buffer_312[], float_4 buffer_313[], float_4 buffer_314[], float_4 buffer_315[], float_4 
buffer_316[], float_4 buffer_317[], float_4 buffer_318[], float_4 buffer_319[], float_4 
buffer_320[], float_4 buffer_321[], float_4 buffer_322[], float_4 buffer_323[], float_4 
buffer_324[], float_4 buffer_325[], float_4 buffer_326[], float_4 buffer_327[], float_4 
buffer_328[], float_4 buffer_329[], float_4 buffer_330[], float_4 buffer_331[], float_4 
buffer_332[], float_4 buffer_333[], float_4 buffer_334[], float_4 buffer_335[], float_4 
buffer_336[], float_4 buffer_337[], float_4 buffer_338[], float_4 buffer_339[], float_4 
buffer_340[], float_4 buffer_341[], float_4 buffer_342[], float_4 buffer_343[], float_4 
buffer_344[], float_4 buffer_345[], float_4 buffer_346[], float_4 buffer_347[], float_4 
buffer_348[], float_4 buffer_349[], float_4 buffer_350[], float_4 buffer_351[], float_4 
buffer_352[], float_4 buffer_353[], float_4 buffer_354[], float_4 buffer_355[], float_4 
buffer_356[], float_4 buffer_357[], float_4 buffer_358[], float_4 buffer_359[], float_4 
buffer_360[], float_4 buffer_361[], float_4 buffer_362[], float_4 buffer_363[], float_4 
buffer_364[], float_4 buffer_365[], float_4 buffer_366[], float_4 buffer_367[], float_4 
buffer_368[], float_4 buffer_369[], float_4 buffer_370[], float_4 buffer_371[], float_4 
buffer_372[], float_4 buffer_373[], float_4 buffer_374[], float_4 buffer_375[], float_4 
buffer_376[], float_4 buffer_377[], float_4 buffer_378[], float_4 buffer_379[], float_4 
buffer_380[], float_4 buffer_381[], float_4 buffer_382[], float_4 buffer_383[], float_4 
buffer_384[], float_4 buffer_385[], float_4 buffer_386[], float_4 buffer_387[], float_4 
buffer_388[], float_4 buffer_389[], float_4 buffer_390[], float_4 buffer_391[], float_4 
buffer_392[], float_4 buffer_393[], float_4 buffer_394[], float_4 buffer_395[], float_4 
buffer_396[], float_4 buffer_397[], float_4 buffer_398[], float_4 buffer_399[], float_4 
buffer_400[], float_4 buffer_401[], float_4 buffer_402[], float_4 buffer_403[], float_4 
buffer_404[], float_4 buffer_405[], float_4 buffer_406[], float_4 buffer_407[], float_4 
buffer_408[], float_4 buffer_409[], float_4 buffer_410[], float_4 buffer_411[], float_4 
buffer_412[], float_4 buffer_413[], float_4 buffer_414[], float_4 buffer_415[], float_4 
buffer_416[], float_4 buffer_417[], float_4 buffer_418[], float_4 buffer_419[], float_4 
buffer_420[], float_4 buffer_421[], float_4 buffer_422[], float_4 buffer_423[], float_4 
buffer_424[], float_4 buffer_425[], float_4 buffer_426[], float_4 buffer_427[], float_4 
buffer_428[], float_4 buffer_429[], float_4 buffer_430[], float_4 buffer_431[], float_4 
buffer_432[], float_4 buffer_433[], float_4 buffer_434[], float_4 buffer_435[], float_4 
buffer_436[], float_4 buffer_437[], float_4 buffer_438[], float_4 buffer_439[], float_4 
buffer_440[], float_4 buffer_441[], float_4 buffer_442[], float_4 buffer_443[], float_4 
buffer_444[], float_4 buffer_445[], float_4 buffer_446[], float_4 buffer_447[], float_4 
buffer_448[], float_4 buffer_449[], float_4 buffer_450[], float_4 buffer_451[], float_4 
buffer_452[], float_4 buffer_453[], float_4 buffer_454[], float_4 buffer_455[], float_4 
buffer_456[], float_4 buffer_457[], float_4 buffer_458[], float_4 buffer_459[], float_4 
buffer_460[], float_4 buffer_461[], float_4 buffer_462[], float_4 buffer_463[], float_4 
buffer_464[], float_4 buffer_465[], float_4 buffer_466[], float_4 buffer_467[], float_4 
buffer_468[], float_4 buffer_469[], float_4 buffer_470[], float_4 buffer_471[], float_4 
buffer_472[], float_4 buffer_473[], float_4 buffer_474[], float_4 buffer_475[], float_4 
buffer_476[], float_4 buffer_477[], float_4 buffer_478[], float_4 buffer_479[], float_4 
buffer_480[], float_4 buffer_481[], float_4 buffer_482[], float_4 buffer_483[], float_4 
buffer_484[], float_4 buffer_485[], float_4 buffer_486[], float_4 buffer_487[], float_4 
buffer_488[], float_4 buffer_489[], float_4 buffer_490[], float_4 buffer_491[], float_4 
buffer_492[], float_4 buffer_493[], float_4 buffer_494[], float_4 buffer_495[], float_4 
buffer_496[], float_4 buffer_497[], float_4 buffer_498[], float_4 buffer_499[], float_4 
buffer_500[], float_4 buffer_501[], float_4 buffer_502[], float_4 buffer_503[], float_4 
buffer_504[], float_4 buffer_505[], float_4 buffer_506[], float_4 buffer_507[], float_4 
buffer_508[], float_4 buffer_509[], float_4 buffer_510[], float_4 buffer_511[], float_4 
buffer_512[], float_4 buffer_513[], float_4 buffer_514[], float_4 buffer_515[], float_4 
buffer_516[], float_4 buffer_517[], float_4 buffer_518[], float_4 buffer_519[], float_4 
buffer_520[], float_4 buffer_521[], float_4 buffer_522[], float_4 buffer_523[], float_4 
buffer_524[], float_4 buffer_525[], float_4 buffer_526[], float_4 buffer_527[], float_4 
buffer_528[], float_4 buffer_529[], float_4 buffer_530[], float_4 buffer_531[], float_4 
buffer_532[], float_4 buffer_533[], float_4 buffer_534[], float_4 buffer_535[], float_4 
buffer_536[], float_4 buffer_537[], float_4 buffer_538[], float_4 buffer_539[], float_4 
buffer_540[], float_4 buffer_541[], float_4 buffer_542[], float_4 buffer_543[], float_4 
buffer_544[], float_4 buffer_545[], float_4 buffer_546[], float_4 buffer_547[], float_4 
buffer_548[], float_4 buffer_549[], float_4 buffer_550[], float_4 buffer_551[], float_4 
buffer_552[], float_4 buffer_553[], float_4 buffer_554[], float_4 buffer_555[], float_4 
buffer_556[], float_4 buffer_557[], float_4 buffer_558[], float_4 buffer_559[], float_4 
buffer_560[], float_4 buffer_561[], float_4 buffer_562[], float_4 buffer_563[], float_4 
buffer_564[], float_4 buffer_565[], float_4 buffer_566[], float_4 buffer_567[], float_4 
buffer_568[], float_4 buffer_569[], float_4 buffer_570[], float_4 buffer_571[], float_4 
buffer_572[], float_4 buffer_573[], float_4 buffer_574[], float_4 buffer_575[], float_4 
buffer_576[], float_4 buffer_577[], float_4 buffer_578[], float_4 buffer_579[], float_4 
buffer_580[], float_4 buffer_581[], float_4 buffer_582[], float_4 buffer_583[], float_4 
buffer_584[], float_4 buffer_585[], float_4 buffer_586[], float_4 buffer_587[], float_4 
buffer_588[], float_4 buffer_589[], float_4 buffer_590[], float_4 buffer_591[], float_4 
buffer_592[], float_4 buffer_593[], float_4 buffer_594[], float_4 buffer_595[], float_4 
buffer_596[], float_4 buffer_597[], float_4 buffer_598[], float_4 buffer_599[], float_4 
buffer_600[], float_4 buffer_601[], float_4 buffer_602[], float_4 buffer_603[], float_4 
buffer_604[], float_4 buffer_605[], float_4 buffer_606[], float_4 buffer_607[], float_4 
buffer_608[], float_4 buffer_609[], float_4 buffer_610[], float_4 buffer_611[], float_4 
buffer_612[], float_4 buffer_613[], float_4 buffer_614[], float_4 buffer_615[], float_4 
buffer_616[], float_4 buffer_617[], float_4 buffer_618[], float_4 buffer_619[], float_4 
buffer_620[], float_4 buffer_621[], float_4 buffer_622[], float_4 buffer_623[], float_4 
buffer_624[], float_4 buffer_625[], float_4 buffer_626[], float_4 buffer_627[], float_4 
buffer_628[], float_4 buffer_629[], float_4 buffer_630[], float_4 buffer_631[], float_4 
buffer_632[], float_4 buffer_633[], float_4 buffer_634[], float_4 buffer_635[], float_4 
buffer_636[], float_4 buffer_637[], float_4 buffer_638[], float_4 buffer_639[], float_4 
buffer_640[], float_4 buffer_641[], float_4 buffer_642[], float_4 buffer_643[], float_4 
buffer_644[], float_4 buffer_645[], float_4 buffer_646[], float_4 buffer_647[], float_4 
buffer_648[], float_4 buffer_649[], float_4 buffer_650[], float_4 buffer_651[], float_4 
buffer_652[], float_4 buffer_653[], float_4 buffer_654[], float_4 buffer_655[], float_4 
buffer_656[], float_4 buffer_657[], float_4 buffer_658[], float_4 buffer_659[], float_4 
buffer_660[], float_4 buffer_661[], float_4 buffer_662[], float_4 buffer_663[], float_4 
buffer_664[], float_4 buffer_665[], float_4 buffer_666[], float_4 buffer_667[], float_4 
buffer_668[], float_4 buffer_669[], float_4 buffer_670[], float_4 buffer_671[], float_4 
buffer_672[], float_4 buffer_673[], float_4 buffer_674[], float_4 buffer_675[], float_4 
buffer_676[], float_4 buffer_677[], float_4 buffer_678[], float_4 buffer_679[], float_4 
buffer_680[], float_4 buffer_681[], float_4 buffer_682[], float_4 buffer_683[], float_4 
buffer_684[], float_4 buffer_685[], float_4 buffer_686[], float_4 buffer_687[], float_4 
buffer_688[], float_4 buffer_689[], float_4 buffer_690[], float_4 buffer_691[], float_4 
buffer_692[], float_4 buffer_693[], float_4 buffer_694[], float_4 buffer_695[], float_4 
buffer_696[], float_4 buffer_697[], float_4 buffer_698[], float_4 buffer_699[], float_4 
buffer_700[], float_4 buffer_701[], float_4 buffer_702[], float_4 buffer_703[], float_4 
buffer_704[], float_4 buffer_705[], float_4 buffer_706[], float_4 buffer_707[], float_4 
buffer_708[], float_4 buffer_709[], float_4 buffer_710[], float_4 buffer_711[], float_4 
buffer_712[], float_4 buffer_713[], float_4 buffer_714[], float_4 buffer_715[], float_4 
buffer_716[], float_4 buffer_717[], float_4 buffer_718[], float_4 buffer_719[], float_4 
buffer_720[], float_4 buffer_721[], float_4 buffer_722[], float_4 buffer_723[], float_4 
buffer_724[], float_4 buffer_725[], float_4 buffer_726[], float_4 buffer_727[], float_4 
buffer_728[], float_4 buffer_729[], float_4 buffer_730[], float_4 buffer_731[], float_4 
buffer_732[], float_4 buffer_733[], float_4 buffer_734[], float_4 buffer_735[], float_4 
buffer_736[], float_4 buffer_737[], float_4 buffer_738[], float_4 buffer_739[], float_4 
buffer_740[], float_4 buffer_741[], float_4 buffer_742[], float_4 buffer_743[], float_4 
buffer_744[], float_4 buffer_745[], float_4 buffer_746[], float_4 buffer_747[], float_4 
buffer_748[], float_4 buffer_749[], float_4 buffer_750[], float_4 buffer_751[], float_4 
buffer_752[], float_4 buffer_753[], float_4 buffer_754[], float_4 buffer_755[], float_4 
buffer_756[], float_4 buffer_757[], float_4 buffer_758[], float_4 buffer_759[], float_4 
buffer_760[], float_4 buffer_761[], float_4 buffer_762[], float_4 buffer_763[], float_4 
buffer_764[], float_4 buffer_765[], float_4 buffer_766[], float_4 buffer_767[], float_4 
buffer_768[], float_4 buffer_769[], float_4 buffer_770[], float_4 buffer_771[], float_4 
buffer_772[], float_4 buffer_773[], float_4 buffer_774[], float_4 buffer_775[], float_4 
buffer_776[], float_4 buffer_777[], float_4 buffer_778[], float_4 buffer_779[], float_4 
buffer_780[], float_4 buffer_781[], float_4 buffer_782[], float_4 buffer_783[], float_4 
buffer_784[], float_4 buffer_785[], float_4 buffer_786[], float_4 buffer_787[], float_4 
buffer_788[], float_4 buffer_789[], float_4 buffer_790[], float_4 buffer_791[], float_4 
buffer_792[], float_4 buffer_793[], float_4 buffer_794[], float_4 buffer_795[], float_4 
buffer_796[], float_4 buffer_797[], float_4 buffer_798[], float_4 buffer_799[], float_4 
buffer_800[], float_4 buffer_801[], float_4 buffer_802[], float_4 buffer_803[], float_4 
buffer_804[], float_4 buffer_805[], float_4 buffer_806[], float_4 buffer_807[], float_4 
buffer_808[], float_4 buffer_809[], float_4 buffer_810[], float_4 buffer_811[], float_4 
buffer_812[], float_4 buffer_813[], float_4 buffer_814[], float_4 buffer_815[], float_4 
buffer_816[], float_4 buffer_817[], float_4 buffer_818[], float_4 buffer_819[], float_4 
buffer_820[], float_4 buffer_821[], float_4 buffer_822[], float_4 buffer_823[], float_4 
buffer_824[], float_4 buffer_825[], float_4 buffer_826[], float_4 buffer_827[], float_4 
buffer_828[], float_4 buffer_829[], float_4 buffer_830[], float_4 buffer_831[], float_4 
buffer_832[], float_4 buffer_833[], float_4 buffer_834[], float_4 buffer_835[], float_4 
buffer_836[], float_4 buffer_837[], float_4 buffer_838[], float_4 buffer_839[], float_4 
buffer_840[], float_4 buffer_841[], float_4 buffer_842[], float_4 buffer_843[], float_4 
buffer_844[], float_4 buffer_845[], float_4 buffer_846[], float_4 buffer_847[], float_4 
buffer_848[], float_4 buffer_849[], float_4 buffer_850[], float_4 buffer_851[], float_4 
buffer_852[], float_4 buffer_853[], float_4 buffer_854[], float_4 buffer_855[], float_4 
buffer_856[], float_4 buffer_857[], float_4 buffer_858[], float_4 buffer_859[], float_4 
buffer_860[], float_4 buffer_861[], float_4 buffer_862[], float_4 buffer_863[], float_4 
buffer_864[], float_4 buffer_865[], float_4 buffer_866[], float_4 buffer_867[], float_4 
buffer_868[], float_4 buffer_869[], float_4 buffer_870[], float_4 buffer_871[], float_4 
buffer_872[], float_4 buffer_873[], float_4 buffer_874[], float_4 buffer_875[], float_4 
buffer_876[], float_4 buffer_877[], float_4 buffer_878[], float_4 buffer_879[], float_4 
buffer_880[], float_4 buffer_881[], float_4 buffer_882[], float_4 buffer_883[], float_4 
buffer_884[], float_4 buffer_885[], float_4 buffer_886[], float_4 buffer_887[], float_4 
buffer_888[], float_4 buffer_889[], float_4 buffer_890[], float_4 buffer_891[], float_4 
buffer_892[], float_4 buffer_893[], float_4 buffer_894[], float_4 buffer_895[], float_4 
buffer_896[], float_4 buffer_897[], float_4 buffer_898[], float_4 buffer_899[], float_4 
buffer_900[], float_4 buffer_901[], float_4 buffer_902[], float_4 buffer_903[], float_4 
buffer_904[], float_4 buffer_905[], float_4 buffer_906[], float_4 buffer_907[], float_4 
buffer_908[], float_4 buffer_909[], float_4 buffer_910[], float_4 buffer_911[], float_4 
buffer_912[], float_4 buffer_913[], float_4 buffer_914[], float_4 buffer_915[], float_4 
buffer_916[], float_4 buffer_917[], float_4 buffer_918[], float_4 buffer_919[], float_4 
buffer_920[], float_4 buffer_921[], float_4 buffer_922[], float_4 buffer_923[], float_4 
buffer_924[], float_4 buffer_925[], float_4 buffer_926[], float_4 buffer_927[], float_4 
buffer_928[], float_4 buffer_929[], float_4 buffer_930[], float_4 buffer_931[], float_4 
buffer_932[], float_4 buffer_933[], float_4 buffer_934[], float_4 buffer_935[], float_4 
buffer_936[], float_4 buffer_937[], float_4 buffer_938[], float_4 buffer_939[], float_4 
buffer_940[], float_4 buffer_941[], float_4 buffer_942[], float_4 buffer_943[], float_4 
buffer_944[], float_4 buffer_945[], float_4 buffer_946[], float_4 buffer_947[], float_4 
buffer_948[], float_4 buffer_949[], float_4 buffer_950[], float_4 buffer_951[], float_4 
buffer_952[], float_4 buffer_953[], float_4 buffer_954[], float_4 buffer_955[], float_4 
buffer_956[], float_4 buffer_957[], float_4 buffer_958[], float_4 buffer_959[], float_4 
buffer_960[], float_4 buffer_961[], float_4 buffer_962[], float_4 buffer_963[], float_4 
buffer_964[], float_4 buffer_965[], float_4 buffer_966[], float_4 buffer_967[], float_4 
buffer_968[], float_4 buffer_969[], float_4 buffer_970[], float_4 buffer_971[], float_4 
buffer_972[], float_4 buffer_973[], float_4 buffer_974[], float_4 buffer_975[], float_4 
buffer_976[], float_4 buffer_977[], float_4 buffer_978[], float_4 buffer_979[], float_4 
buffer_980[], float_4 buffer_981[], float_4 buffer_982[], float_4 buffer_983[], float_4 
buffer_984[], float_4 buffer_985[], float_4 buffer_986[], float_4 buffer_987[], float_4 
buffer_988[], float_4 buffer_989[], float_4 buffer_990[], float_4 buffer_991[], float_4 
buffer_992[], float_4 buffer_993[], float_4 buffer_994[], float_4 buffer_995[], float_4 
buffer_996[], float_4 buffer_997[], float_4 buffer_998[], float_4 buffer_999[], float_4 
buffer_1000[], float_4 buffer_1001[], float_4 buffer_1002[], float_4 buffer_1003[], 
float_4 buffer_1004[], float_4 buffer_1005[], float_4 buffer_1006[], float_4 
buffer_1007[], float_4 buffer_1008[], float_4 buffer_1009[], float_4 buffer_1010[], 
float_4 buffer_1011[], float_4 buffer_1012[], float_4 buffer_1013[], float_4 
buffer_1014[], float_4 buffer_1015[], float_4 buffer_1016[], float_4 buffer_1017[], 
float_4 buffer_1018[], float_4 buffer_1019[], float_4 buffer_1020[], float_4 
buffer_1021[], const uint32* __restrict__ info, const __grid_constant__ metadata_st 
static_info, const __grid_constant__ scalars_uint32_st scalars_uint32
)

CC'ing @nathanielsimard