{"text_format":"3","workbook":"xls","n_sheets":4,"sheets":["Calculator","Help","GPU Data","Copyright & License"],"not_rendered":{"charts":3,"macros":1,"external_links":1,"defined_names":45},"read":"whole","date_system":"1900","formulas":"cached_only"}

{"sheet":"Calculator","index":0,"n_rows":339,"n_cols":11,"rows_kept":339,"merged":[[0,4,0,11],[1,4,1,11],[3,0,3,1],[3,4,5,11],[44,4,44,10],[46,1,46,3],[47,1,47,3],[47,4,47,10],[51,0,51,1]],"flags":["formulas_cached_only"]}
["CUDA GPU Occupancy Calculator",{"$nulls":3},"Click Here for detailed instructions on how to use this occupancy calculator."]
[{"$nulls":4},"For more information on NVIDIA CUDA, visit http://developer.nvidia.com/cuda"]
[]
["Just follow steps 1, 2, and 3 below! (or click here for help)",{"$nulls":3},"Your chosen resource usage is indicated by the red triangle on the graphs.  The other data points represent the range of possible block sizes, register counts, and shared memory allocation."]
[]
["1.) Select Compute Capability (click):",3.5,"(Help)"]
["1.b) Select Shared Memory Size Config (bytes)",49152]
[]
["2.) Enter your resource usage:"]
["Threads Per Block",256,"(Help)"]
["Registers Per Thread",32]
["Shared Memory Per Block (bytes)",4096]
[]
["(Don't edit anything below this line)"]
[]
["3.) GPU Occupancy Data is displayed here and in the graphs:"]
["Active Threads per Multiprocessor",2048,"(Help)"]
["Active Warps per Multiprocessor",64]
["Active Thread Blocks per Multiprocessor",8]
["Occupancy of each Multiprocessor",{"v":1,"d":"100%"}]
[]
[]
["Physical Limits for GPU Compute Capability:",3.5]
["Threads per Warp",32]
["Warps per Multiprocessor",64]
["Threads per Multiprocessor",2048]
["Thread Blocks per Multiprocessor",16]
["Total # of 32-bit registers per Multiprocessor",65536]
["Register allocation unit size",256]
["Register allocation granularity","warp"]
["Registers per Thread",255]
["Shared Memory per Multiprocessor (bytes)",49152]
["Shared Memory Allocation unit size",256]
["Warp allocation granularity",4]
["Maximum Thread Block Size",1024]
[]
["Allocated Resources","Per Block","Limit Per SM","= Allocatable Blocks Per SM"]
["Warps              (Threads Per Block / Threads Per Warp)",8,64,8]
["Registers          (Warp limit per SM due to per-warp reg count)",8,64,8]
["Shared Memory (Bytes)",4096,49152,12]
["Note: SM is an abbreviation for (Streaming) Multiprocessor"]
[]
["Maximum Thread Blocks Per Multiprocessor","Blocks/SM","* Warps/Block","= Warps/SM"]
["Limited by Max Warps or Max Blocks per Multiprocessor",8,8,64]
["Limited by Registers per Multiprocessor",8,8,64]
["Limited by Shared Memory per Multiprocessor",12,8,0]
["Note: Occupancy limiter is shown in orange","Physical Max Warps/SM = 64"]
[null,"Occupancy = 64 / 64 = 100%"]
[]
["CUDA Occupancy Calculator"]
["Version:",5.1]
["Copyright and License"]
{"$rows_empty":30}
["Threads","Warps/Multiprocessor",{"$nulls":3},"Registers","Warps/Multiprocessor",null,null,"Shared Memory","Warps/Multiprocessor"]
[256,64,{"$nulls":3},32,64,null,null,4096,64]
[32,12,{"$nulls":3},1,64,null,null,0,64]
[64,24,{"$nulls":3},2,64,null,null,512,64]
[96,36,{"$nulls":3},3,64,null,null,1024,64]
[128,48,{"$nulls":3},4,64,null,null,1536,64]
[160,60,{"$nulls":3},5,64,null,null,2048,64]
[192,60,{"$nulls":3},6,64,null,null,2560,64]
[224,63,{"$nulls":3},7,64,null,null,3072,64]
[256,64,{"$nulls":3},8,64,null,null,3584,64]
[288,63,{"$nulls":3},9,64,null,null,4096,64]
[320,60,{"$nulls":3},10,64,null,null,4608,64]
[352,55,{"$nulls":3},11,64,null,null,5120,64]
[384,60,{"$nulls":3},12,64,null,null,5632,64]
[416,52,{"$nulls":3},13,64,null,null,6144,64]
[448,56,{"$nulls":3},14,64,null,null,6656,56]
[480,60,{"$nulls":3},15,64,null,null,7168,48]
[512,64,{"$nulls":3},16,64,null,null,7680,48]
[544,51,{"$nulls":3},17,64,null,null,8192,48]
[576,54,{"$nulls":3},18,64,null,null,8704,40]
[608,57,{"$nulls":3},19,64,null,null,9216,40]
[640,60,{"$nulls":3},20,64,null,null,9728,40]
[672,63,{"$nulls":3},21,64,null,null,10240,32]
[704,44,{"$nulls":3},22,64,null,null,10752,32]
[736,46,{"$nulls":3},23,64,null,null,11264,32]
[768,48,{"$nulls":3},24,64,null,null,11776,32]
[800,50,{"$nulls":3},25,64,null,null,12288,32]
[832,52,{"$nulls":3},26,64,null,null,12800,24]
[864,54,{"$nulls":3},27,64,null,null,13312,24]
[896,56,{"$nulls":3},28,64,null,null,13824,24]
[928,58,{"$nulls":3},29,64,null,null,14336,24]
[960,60,{"$nulls":3},30,64,null,null,14848,24]
[992,62,{"$nulls":3},31,64,null,null,15360,24]
[1024,64,{"$nulls":3},32,64,null,null,15872,24]
[1056,33,{"$nulls":3},33,48,null,null,16384,24]
[1088,34,{"$nulls":3},34,48,null,null,16896,16]
[1120,35,{"$nulls":3},35,48,null,null,17408,16]
[1152,36,{"$nulls":3},36,48,null,null,17920,16]
[1184,37,{"$nulls":3},37,48,null,null,18432,16]
[1216,38,{"$nulls":3},38,48,null,null,18944,16]
[1248,39,{"$nulls":3},39,48,null,null,19456,16]
[1280,40,{"$nulls":3},40,48,null,null,19968,16]
[1312,41,{"$nulls":3},41,40,null,null,20480,16]
[1344,42,{"$nulls":3},42,40,null,null,20992,16]
[1376,43,{"$nulls":3},43,40,null,null,21504,16]
[1408,44,{"$nulls":3},44,40,null,null,22016,16]
[1440,45,{"$nulls":3},45,40,null,null,22528,16]
[1472,46,{"$nulls":3},46,40,null,null,23040,16]
[1504,47,{"$nulls":3},47,40,null,null,23552,16]
[1536,48,{"$nulls":3},48,40,null,null,24064,16]
[1568,{"$nulls":4},49,32,null,null,24576,16]
[1600,{"$nulls":4},50,32,null,null,25088,8]
[1632,{"$nulls":4},51,32,null,null,25600,8]
[1664,{"$nulls":4},52,32,null,null,26112,8]
[1696,{"$nulls":4},53,32,null,null,26624,8]
[1728,{"$nulls":4},54,32,null,null,27136,8]
[1760,{"$nulls":4},55,32,null,null,27648,8]
[1792,{"$nulls":4},56,32,null,null,28160,8]
[1824,{"$nulls":4},57,32,null,null,28672,8]
[1856,{"$nulls":4},58,32,null,null,29184,8]
[1888,{"$nulls":4},59,32,null,null,29696,8]
[1920,{"$nulls":4},60,32,null,null,30208,8]
[1952,{"$nulls":4},61,32,null,null,30720,8]
[1984,{"$nulls":4},62,32,null,null,31232,8]
[2016,{"$nulls":4},63,32,null,null,31744,8]
[2048,{"$nulls":4},64,32,null,null,32256,8]
[2080,{"$nulls":4},65,24,null,null,32768,8]
[2112,{"$nulls":4},66,24,null,null,33280,8]
[2144,{"$nulls":4},67,24,null,null,33792,8]
[2176,{"$nulls":4},68,24,null,null,34304,8]
[2208,{"$nulls":4},69,24,null,null,34816,8]
[2240,{"$nulls":4},70,24,null,null,35328,8]
[2272,{"$nulls":4},71,24,null,null,35840,8]
[2304,{"$nulls":4},72,24,null,null,36352,8]
[2336,{"$nulls":4},73,24,null,null,36864,8]
[2368,{"$nulls":4},74,24,null,null,37376,8]
[2400,{"$nulls":4},75,24,null,null,37888,8]
[2432,{"$nulls":4},76,24,null,null,38400,8]
[2464,{"$nulls":4},77,24,null,null,38912,8]
[2496,{"$nulls":4},78,24,null,null,39424,8]
[2528,{"$nulls":4},79,24,null,null,39936,8]
[2560,{"$nulls":4},80,24,null,null,40448,8]
[2592,{"$nulls":4},81,16,null,null,40960,8]
[2624,{"$nulls":4},82,16,null,null,41472,8]
[2656,{"$nulls":4},83,16,null,null,41984,8]
[2688,{"$nulls":4},84,16,null,null,42496,8]
[2720,{"$nulls":4},85,16,null,null,43008,8]
[2752,{"$nulls":4},86,16,null,null,43520,8]
[2784,{"$nulls":4},87,16,null,null,44032,8]
[2816,{"$nulls":4},88,16,null,null,44544,8]
[2848,{"$nulls":4},89,16,null,null,45056,8]
[2880,{"$nulls":4},90,16,null,null,45568,8]
[2912,{"$nulls":4},91,16,null,null,46080,8]
[2944,{"$nulls":4},92,16,null,null,46592,8]
[2976,{"$nulls":4},93,16,null,null,47104,8]
[3008,{"$nulls":4},94,16,null,null,47616,8]
[3040,{"$nulls":4},95,16,null,null,48128,8]
[3072,{"$nulls":4},96,16,null,null,48640,8]
[{"$nulls":5},97,16,null,null,49152,8]
[{"$nulls":5},98,16]
[{"$nulls":5},99,16]
[{"$nulls":5},100,16]
[{"$nulls":5},101,16]
[{"$nulls":5},102,16]
[{"$nulls":5},103,16]
[{"$nulls":5},104,16]
[{"$nulls":5},105,16]
[{"$nulls":5},106,16]
[{"$nulls":5},107,16]
[{"$nulls":5},108,16]
[{"$nulls":5},109,16]
[{"$nulls":5},110,16]
[{"$nulls":5},111,16]
[{"$nulls":5},112,16]
[{"$nulls":5},113,16]
[{"$nulls":5},114,16]
[{"$nulls":5},115,16]
[{"$nulls":5},116,16]
[{"$nulls":5},117,16]
[{"$nulls":5},118,16]
[{"$nulls":5},119,16]
[{"$nulls":5},120,16]
[{"$nulls":5},121,16]
[{"$nulls":5},122,16]
[{"$nulls":5},123,16]
[{"$nulls":5},124,16]
[{"$nulls":5},125,16]
[{"$nulls":5},126,16]
[{"$nulls":5},127,16]
[{"$nulls":5},128,16]
[{"$nulls":5},129,8]
[{"$nulls":5},130,8]
[{"$nulls":5},131,8]
[{"$nulls":5},132,8]
[{"$nulls":5},133,8]
[{"$nulls":5},134,8]
[{"$nulls":5},135,8]
[{"$nulls":5},136,8]
[{"$nulls":5},137,8]
[{"$nulls":5},138,8]
[{"$nulls":5},139,8]
[{"$nulls":5},140,8]
[{"$nulls":5},141,8]
[{"$nulls":5},142,8]
[{"$nulls":5},143,8]
[{"$nulls":5},144,8]
[{"$nulls":5},145,8]
[{"$nulls":5},146,8]
[{"$nulls":5},147,8]
[{"$nulls":5},148,8]
[{"$nulls":5},149,8]
[{"$nulls":5},150,8]
[{"$nulls":5},151,8]
[{"$nulls":5},152,8]
[{"$nulls":5},153,8]
[{"$nulls":5},154,8]
[{"$nulls":5},155,8]
[{"$nulls":5},156,8]
[{"$nulls":5},157,8]
[{"$nulls":5},158,8]
[{"$nulls":5},159,8]
[{"$nulls":5},160,8]
[{"$nulls":5},161,8]
[{"$nulls":5},162,8]
[{"$nulls":5},163,8]
[{"$nulls":5},164,8]
[{"$nulls":5},165,8]
[{"$nulls":5},166,8]
[{"$nulls":5},167,8]
[{"$nulls":5},168,8]
[{"$nulls":5},169,8]
[{"$nulls":5},170,8]
[{"$nulls":5},171,8]
[{"$nulls":5},172,8]
[{"$nulls":5},173,8]
[{"$nulls":5},174,8]
[{"$nulls":5},175,8]
[{"$nulls":5},176,8]
[{"$nulls":5},177,8]
[{"$nulls":5},178,8]
[{"$nulls":5},179,8]
[{"$nulls":5},180,8]
[{"$nulls":5},181,8]
[{"$nulls":5},182,8]
[{"$nulls":5},183,8]
[{"$nulls":5},184,8]
[{"$nulls":5},185,8]
[{"$nulls":5},186,8]
[{"$nulls":5},187,8]
[{"$nulls":5},188,8]
[{"$nulls":5},189,8]
[{"$nulls":5},190,8]
[{"$nulls":5},191,8]
[{"$nulls":5},192,8]
[{"$nulls":5},193,8]
[{"$nulls":5},194,8]
[{"$nulls":5},195,8]
[{"$nulls":5},196,8]
[{"$nulls":5},197,8]
[{"$nulls":5},198,8]
[{"$nulls":5},199,8]
[{"$nulls":5},200,8]
[{"$nulls":5},201,8]
[{"$nulls":5},202,8]
[{"$nulls":5},203,8]
[{"$nulls":5},204,8]
[{"$nulls":5},205,8]
[{"$nulls":5},206,8]
[{"$nulls":5},207,8]
[{"$nulls":5},208,8]
[{"$nulls":5},209,8]
[{"$nulls":5},210,8]
[{"$nulls":5},211,8]
[{"$nulls":5},212,8]
[{"$nulls":5},213,8]
[{"$nulls":5},214,8]
[{"$nulls":5},215,8]
[{"$nulls":5},216,8]
[{"$nulls":5},217,8]
[{"$nulls":5},218,8]
[{"$nulls":5},219,8]
[{"$nulls":5},220,8]
[{"$nulls":5},221,8]
[{"$nulls":5},222,8]
[{"$nulls":5},223,8]
[{"$nulls":5},224,8]
[{"$nulls":5},225,8]
[{"$nulls":5},226,8]
[{"$nulls":5},227,8]
[{"$nulls":5},228,8]
[{"$nulls":5},229,8]
[{"$nulls":5},230,8]
[{"$nulls":5},231,8]
[{"$nulls":5},232,8]
[{"$nulls":5},233,8]
[{"$nulls":5},234,8]
[{"$nulls":5},235,8]
[{"$nulls":5},236,8]
[{"$nulls":5},237,8]
[{"$nulls":5},238,8]
[{"$nulls":5},239,8]
[{"$nulls":5},240,8]
[{"$nulls":5},241,8]
[{"$nulls":5},242,8]
[{"$nulls":5},243,8]
[{"$nulls":5},244,8]
[{"$nulls":5},245,8]
[{"$nulls":5},246,8]
[{"$nulls":5},247,8]
[{"$nulls":5},248,8]
[{"$nulls":5},249,8]
[{"$nulls":5},250,8]
[{"$nulls":5},251,8]
[{"$nulls":5},252,8]
[{"$nulls":5},253,8]
[{"$nulls":5},254,8]
[{"$nulls":5},255,8]

{"sheet":"Help","index":1,"n_rows":37,"n_cols":2,"rows_kept":37,"flags":["formulas_cached_only"]}
[]
[null,"IMPORTANT"]
[null,"This spreadsheet requires Excel macros for full functionality.  When you load this file, make sure you enable macros "]
[null,"because they are often disabled by default by Excel."]
[]
[null,"Overview"]
[null,"The CUDA Occupancy Calculator allows you to compute the multiprocessor occupancy of a GPU by a given CUDA kernel.  The multiprocessor occupancy is the ratio of active warps to the maximum number of warps supported on a multiprocessor of the GPU.  Each multiprocessor on the device has a set of N registers available for use by CUDA program threads.  These registers are a shared resource that are allocated among the thread blocks executing on a multiprocessor.  The CUDA compiler attempts to minimize register usage to maximize the number of thread blocks that can be active in the machine simultaneously.   If a program tries to launch a kernel for which the registers used per thread times the thread block size is greater than N, the launch will fail.  "]
[]
[null,"The size of N on GPUs with compute capability 1.0-1.1 is 8192 32-bit registers per multiprocessor.  On GPUs with compute "]
[null,"capability 1.2-1.3, N = 16384.  On GPUs with compute capability 2.0-2.1, N = 32768.  On GPUs with compute capability 3.0,"]
[null,"N=65536."]
[null,"Maximizing the occupancy can help to cover latency during global memory loads that are followed by a __syncthreads().  The occupancy is determined by the amount of shared memory and registers used by each thread block.  Because of this, programmers need to choose the size of thread blocks with care in order to maximize occupancy.  This GPU Occupancy Calculator can assist in choosing thread block size based on shared memory and register requirements."]
[]
[null,"Instructions"]
[null,"Using the CUDA Occupancy Calculator is as easy as 1-2-3.  Change to the calculator sheet and follow these three steps."]
[null,"1.) First select your device's compute capability in the green box."]
[null,"Click to go there…"]
[null,"1.b) If your compute capability supports it, you will be shown a second green box in which you can select the size in bytes of the shared memory (configurable at run time in CUDA)."]
[null,"Click to go there…"]
[null,"2.) For the kernel you are profiling, enter the number of threads per thread block, the registers used per thread, and the total shared memory used per thread block in bytes in the orange block.  See below for how to find the registers used per thread."]
[null,"Click to go there…"]
[null,"3.) Examine the blue box and the graph to the right.  This will tell you the occupancy, as well as the number of active threads, warps, and thread blocks per multiprocessor, and the maximum number of active blocks on the GPU.  The graph will show you the occupancy for your chosen block size as a red triangle, and for all other possible block sizes as a line graph."]
[null,"Click to go there…"]
[null,"You can now experiment with how different thread block sizes, register counts, and shared memory usages can affect your GPU occupancy."]
[]
[null,"Determining Registers Per Thread and Shared Memory Per Thread Block"]
[null,"To determine the number of registers used per thread in your kernel, simply compile the kernel code using the option               --ptxas-options=-v to nvcc.  This will output information about register, local memory, shared memory, and constant memory usage for each kernel in the .cu file.  However, if your kernel declares any external shared memory that is allocated dynamically, you will need to add the (statically allocated) shared memory reported by ptxas to the amount you dynamically allocate at run time to get the correct shared memory usage.  An example of the verbose ptxas output is as follows:"]
[]
[null,"ptxas info    : Compiling entry function '_Z8my_kernelPf' for 'sm_10'"]
[null,"ptxas info    : Used 5 registers, 8+16 bytes smem"]
[]
[null,"Let's say \"my_kernel\" contains an external shared memory array which is allocated to be 2048 bytes at run time.  Then our total shared memory usage per block is 2048+8+16 = 2072 bytes.  We enter this into the box labeled \"shared memory per block (bytes)\" in this occupancy calculator, and we also enter the number of registers used by my_kernel, 5, in the box labeled registers per thread.  We then enter our thread block size and the calculator will display the occupancy."]
[]
[null,"Notes about Occupancy"]
[null,"Higher occupancy does not necessarily mean higher performance. If a kernel is not bandwidth-limited or latency-limited, then increasing occupancy will not necessarily increase performance.  If a kernel grid is already running at least one thread block per multiprocessor in the GPU, and it is bottlenecked by computation and not by global memory accesses, then increasing occupancy may have no effect. In fact, making changes just to increase occupancy can have other effects, such as additional instructions, more register spills to local memory (which is off-chip), more divergent branches, etc. As with any optimization, you should experiment to see how changes affect the *wall clock time* of the kernel execution.  For bandwidth-bound applications, on the other hand, increasing occupancy can help better hide the latency of memory accesses, and therefore improve performance."]
[]
[null,"For more information on NVIDIA CUDA, visit http://developer.nvidia.com/cuda"]

{"sheet":"GPU Data","index":2,"n_rows":22,"n_cols":10,"rows_kept":22,"flags":["formulas_cached_only"]}
[]
["Compute Capability",null,{"v":1,"d":"1.0"},1.1,1.2,1.3,{"v":2,"d":"2.0"},2.1,{"v":3,"d":"3.0"},3.5]
["SM Version",null,"sm_10","sm_11","sm_12","sm_13","sm_20","sm_21","sm_30","sm_35"]
["Threads / Warp",null,32,32,32,32,32,32,32,32]
["Warps / Multiprocessor",null,24,24,32,32,48,48,64,64]
["Threads / Multiprocessor",null,768,768,1024,1024,1536,1536,2048,2048]
["Thread Blocks / Multiprocessor",null,8,8,8,8,8,8,16,16]
["Max Shared Memory / Multiprocessor (bytes)",null,16384,16384,16384,16384,49152,49152,49152,49152]
["Register File Size",null,8192,8192,16384,16384,32768,32768,65536,65536]
["Register Allocation Unit Size",null,256,256,512,512,64,64,256,256]
["Allocation Granularity",null,"block","block","block","block","warp","warp","warp","warp"]
["Max Registers / Thread",null,124,124,124,124,63,63,63,255]
["Shared Memory Allocation Unit Size",null,512,512,512,512,128,128,256,256]
["Warp allocation granularity",null,2,2,2,2,2,2,4,4]
["Max Thread Block Size",null,512,512,512,512,1024,1024,1024,1024]
[]
["Shared Memory Size Configurations (bytes)",null,16384,16384,16384,16384,49152,49152,49152,49152]
["[note: default at top of list]",{"$nulls":5},16384,16384,16384,16384]
[{"$nulls":8},32768,32768]
[]
["Warp register allocation granularities",{"$nulls":5},64,64,256,256]
["[note: default at top of list]","21,22,29,30,37,38,45,46,",{"$nulls":4},128,128]

{"sheet":"Copyright & License","index":3,"n_rows":6,"n_cols":2,"rows_kept":6,"flags":["formulas_cached_only"]}
[]
[null,"\nCopyright 1993-2012 NVIDIA Corporation.  All rights reserved. \n"]
[null,"NOTICE TO USER:\n\nThis spreadsheet and data is subject to NVIDIA ownership rights under U.S. and international Copyright laws.  Users and possessors of this spreadsheet and data are hereby granted a nonexclusive, royalty-free license to use it in individual and commercial software.\n\nNVIDIA MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THIS SPREADSHEET AND DATA FOR ANY PURPOSE.  IT IS PROVIDED \"AS IS\" WITHOUT EXPRESS OR IMPLIED WARRANTY OF ANY KIND.  NVIDIA DISCLAIMS ALL WARRANTIES WITH REGARD TO THIS SPREADSHEET AND DATA, INCLUDING ALL IMPLIED WARRANTIES OF MERCHANTABILITY, NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE.  IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL, OR CONSEQUENTIAL DAMAGES, OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS OF USE, DATA OR PROFITS,  WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE OR OTHER TORTIOUS ACTION,  ARISING OUT OF OR IN CONNECTION WITH THE USE OR PERFORMANCE OF THIS SPREADSHEET AND DATA.\n\n"]
[null,"U.S. Government End Users.   This spreadsheet and data are a \"commercial item\" as that term is defined at  48 C.F.R. 2.101 (OCT 1995), consisting  of \"commercial computer software\"  and \"commercial computer software documentation\" as such terms are  used in 48 C.F.R. 12.212 (SEPT 1995) and is provided to the U.S. Government only as a commercial end item. Consistent with 48 C.F.R.12.212 and 48 C.F.R. 227.7202-1 through 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the spreadsheet and data with only those rights set forth herein.  Any use of this spreadsheet and data in individual and commercial software must include, in the user documentation and internal comments to the code, the above Disclaimer and U.S. Government End Users Notice."]
[]
[null,"For more information on NVIDIA CUDA, visit http://www.nvidia.com/cuda"]