A100 显卡关键参数

全局视图

首先看top图，
GA100是无损卡，就是说上面所有的unit都是好的，如下，

A100包含有108个SM，每个SM最大可以容纳1024个threads.
说白了就是GA100的有缺陷的卡就为A100,撇去那个不能用的GPC

SM视图

下面是一个SM的视图：

算力

算力：

加工工艺

工艺采用的是7nm工艺：

关键参数

实际测量参数

下面是关键参数：

device properties : 
	name : NVIDIA A100-PCIE-40GB
	totalGlobalMem : 42298834944
	sharedMemPerBlock : 49152
	regsPerBlock : 65536
	warpSize : 32
	memPitch : 2147483647
	maxThreadsPerBlock : 1024
	maxThreadsDim[0] : 1024
	maxThreadsDim[1] : 1024
	maxThreadsDim[2] : 64
	maxGridSize[0] : 2147483647
	maxGridSize[1] : 65535
	maxGridSize[2] : 65535
	clockRate : 1410000
	totalConstMem : 65536
	major : 8
	minor : 0
	textureAlignment : 512
	texturePitchAlignment : 32
	deviceOverlap : 1
	multiProcessorCount : 108
	kernelExecTimeoutEnabled : 0
	integrated : 0
	canMapHostMemory : 1
	computeMode : 0
	concurrentKernels : 1
	ECCEnabled : 1
	pciBusID : 64
	pciDeviceID : 0
	pciDomainID : 0
	tccDriver : 0
	asyncEngineCount : 3
	unifiedAddressing : 1
	memoryClockRate : 1215000
	memoryBusWidth : 5120
	l2CacheSize : 41943040
	persistingL2CacheMaxSize : 31457280
	maxThreadsPerMultiProcessor : 2048
	streamPrioritiesSupported : 1
	globalL1CacheSupported : 1
	localL1CacheSupported : 1
	sharedMemPerMultiprocessor : 167936
	regsPerMultiprocessor : 65536
	managedMemory : 1
	isMultiGpuBoard : 0
	multiGpuBoardGroupID : 0
	singleToDoublePrecisionPerfRatio : 2
	pageableMemoryAccess : 0
	concurrentManagedAccess : 1
	computePreemptionSupported : 1
	canUseHostPointerForRegisteredMem : 1
	cooperativeLaunch : 1
	cooperativeMultiDeviceLaunch : 1
	pageableMemoryAccessUsesHostPageTables : 0
	directManagedMemAccessFromHost : 0
	accessPolicyMaxWindowSize : 134213632

device limit : 
	deviceLimitStackSize : 1024
	deviceLimitPrintfFifoSize : 7077888
	deviceLimitMallocHeapSize : 8388608
	deviceLimitDevRuntimeSyncDepth : 2
	deviceLimitDevRuntimePendingLaunchCount : 2048
	deviceLimitMaxL2FetchGranularity : 64
	deviceLimitPersistingL2CacheSize : 7864320

summary : 
	register total size : 6.75 MiB
	shared memory size per sm : 164.00 KiB
	shared memory total size : 17.30 MiB
	constant memory total size : 64.00 KiB
	level 2 cache total size : 40.00 MiB
	device memory total size : 39.39 GiB
	device memory bandwidth : 1.56 TB/s
	stack memory total size : 216.00 MiB

block 在SM上的分布

• sm上是以block为单位进行分配的。
• 先分配偶数标号的sm，接着再分配奇数标号的sm
• <<<108,1024>>全部sm占满。
  

grid_dim	block_dim	sm0	sm1	sm2	sm3	sm4	sm5	sm6	sm7	sm8	sm9	sm10	sm11	sm12	sm13	sm14	sm15	sm16	sm17	sm18	sm19	sm20	sm21	sm22	sm23	sm24	sm25	sm26	sm27	sm28	sm29	sm30	sm31	sm32	sm33	sm34	sm35	sm36	sm37	sm38	sm39	sm40	sm41	sm42	sm43	sm44	sm45	sm46	sm47	sm48	sm49	sm50	sm51	sm52	sm53	sm54	sm55	sm56	sm57	sm58	sm59	sm60	sm61	sm62	sm63	sm64	sm65	sm66	sm67	sm68	sm69	sm70	sm71	sm72	sm73	sm74	sm75	sm76	sm77	sm78	sm79	sm80	sm81	sm82	sm83	sm84	sm85	sm86	sm87	sm88	sm89	sm90	sm91	sm92	sm93	sm94	sm95	sm96	sm97	sm98	sm99	sm100	sm101	sm102	sm103	sm104	sm105	sm106	sm107
1	1	1	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0
1	32	32	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0
1	64	64	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0
1	128	128	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0
1	256	256	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0
1	512	512	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0
1	1024	1024	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0
16	1	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0
16	32	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0
16	64	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0
16	128	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0
16	256	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0
16	512	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0
16	1024	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0
32	1	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0
32	32	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0
32	64	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0
32	128	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0
32	256	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0
32	512	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0
32	1024	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0	0
64	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0	1	0
64	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0	32	0
64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0	64	0
64	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0	128	0
64	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0	256	0
64	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0	512	0
64	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0	1024	0
108	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1
108	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32	32
108	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64	64
108	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128	128
108	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256	256
108	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512	512
108	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024	1024

A100上SM/TPC/GPC分组关系

注意，SM都是逻辑idx

具体做法：

//__cooperative__
__global__ void KERNEL_NAME(__TEST_CASE_NAME__)(int *id, uint64_t *clocks) {
  int threadInBlock = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.x * blockDim.y;
  int blockInGrid = blockIdx.x + blockIdx.y * gridDim.x + blockIdx.z * gridDim.x * gridDim.y;
  int oneBlockSize = blockDim.x * blockDim.y * blockDim.z;
  int tidx = threadInBlock + oneBlockSize * blockInGrid;

#pragma unroll
  for (int i = 0; i < SYNC_LOOP; i++) {
      __syncthreads();
  }

  uint64_t start = rt::Clock();

  id[tidx] = __mysmid();
  clocks[__mysmid()] = start;  
}

static void gpc_test_kernel(int grid_dim, int block_dim, uint32_t *h_id, uint64_t *h_clocks) {
  rt::Error_t err;
  Stream_t stream;
  uint32_t *d_id;
  CHECK_ERROR(rt::Malloc((void **)&d_id, sizeof(uint32_t) * grid_dim * block_dim));
  CHECK_ERROR(rt::Memset(d_id, 0, sizeof(uint32_t) * grid_dim * block_dim));


  uint64_t *d_clocks;
  CHECK_ERROR(rt::Malloc((void **)&d_clocks, sizeof(uint64_t) * grid_dim * block_dim));
  CHECK_ERROR(rt::Memset(d_clocks, 0, sizeof(uint64_t) * grid_dim * block_dim));

  CHECK_ERROR(rt::StreamCreate(&stream));

  // kernel function
  void *args[] = {(void *)&d_id, (void *)&d_clocks};
  err = rt::LaunchCooperativeKernel((const void *)(KERNEL_NAME(__TEST_CASE_NAME__)), grid_dim, block_dim, args, 0,
                                    stream);
  CHECK_ERROR(err);
  CHECK_ERROR(rt::GetLastError());

  CHECK_ERROR(rt::StreamSynchronize(stream));
  CHECK_ERROR(rt::DeviceSynchronize());
  CHECK_ERROR(rt::StreamSynchronize(stream));

  CHECK_ERROR(rt::Memcpy(h_id, d_id, sizeof(uint32_t) * 1 * grid_dim * block_dim, rt::MemcpyDeviceToHost));
  CHECK_ERROR(rt::Memcpy(h_clocks, d_clocks, sizeof(uint64_t) * 1 * grid_dim * block_dim, rt::MemcpyDeviceToHost));

  CHECK_ERROR(rt::StreamDestroy(stream));
  CHECK_ERROR(rt::Free(d_id));
  CHECK_ERROR(rt::Free(d_clocks));
}

int mainc(){
......
  rt::Error_t err;
  std::ofstream file2(std::string(test_name) + std::string("_gpc_sm_layout.csv"));


  err = rt::SetDevice(0);
  CHECK_ERROR(err);
  rt::DeviceProp device_prop;
  err = rt::GetDeviceProperties(&device_prop, 0);
  CHECK_ERROR(err);

  grid_dim=device_prop.multiProcessorCount;
  block_dim = 1;

  gpc_test_kernel(grid_dim, block_dim, id, h_clocks);
...
}

sm0	sm1	sm2	sm3	sm4	sm5	sm6	sm7	sm8	sm9	sm10	sm11	sm12	sm13	sm14	sm15	sm16	sm17	sm18	sm19	sm20	sm21	sm22	sm23	sm24	sm25	sm26	sm27	sm28	sm29	sm30	sm31	sm32	sm33	sm34	sm35	sm36	sm37	sm38	sm39	sm40	sm41	sm42	sm43	sm44	sm45	sm46	sm47	sm48	sm49	sm50	sm51	sm52	sm53	sm54	sm55	sm56	sm57	sm58	sm59	sm60	sm61	sm62	sm63	sm64	sm65	sm66	sm67	sm68	sm69	sm70	sm71	sm72	sm73	sm74	sm75	sm76	sm77	sm78	sm79	sm80	sm81	sm82	sm83	sm84	sm85	sm86	sm87	sm88	sm89	sm90	sm91	sm92	sm93	sm94	sm95	sm96	sm97	sm98	sm99	sm100	sm101	sm102	sm103	sm104	sm105	sm106	sm107
1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1	1
3.25863E+14	3.25863E+14	3.25447E+14	3.25447E+14	3.25654E+14	3.25654E+14	3.26144E+14	3.26144E+14	3.2569E+14	3.2569E+14	3.25881E+14	3.25881E+14	3.25386E+14	3.25386E+14	3.25863E+14	3.25863E+14	3.25447E+14	3.25447E+14	3.25654E+14	3.25654E+14	3.26144E+14	3.26144E+14	3.2569E+14	3.2569E+14	3.25881E+14	3.25881E+14	3.25386E+14	3.25386E+14	3.25863E+14	3.25863E+14	3.25447E+14	3.25447E+14	3.25654E+14	3.25654E+14	3.26144E+14	3.26144E+14	3.2569E+14	3.2569E+14	3.25881E+14	3.25881E+14	3.25386E+14	3.25386E+14	3.25863E+14	3.25863E+14	3.25447E+14	3.25447E+14	3.25654E+14	3.25654E+14	3.26144E+14	3.26144E+14	3.2569E+14	3.2569E+14	3.25881E+14	3.25881E+14	3.25386E+14	3.25386E+14	3.25863E+14	3.25863E+14	3.25447E+14	3.25447E+14	3.25654E+14	3.25654E+14	3.26144E+14	3.26144E+14	3.2569E+14	3.2569E+14	3.25881E+14	3.25881E+14	3.25386E+14	3.25386E+14	3.25863E+14	3.25863E+14	3.25447E+14	3.25447E+14	3.25654E+14	3.25654E+14	3.26144E+14	3.26144E+14	3.2569E+14	3.2569E+14	3.25881E+14	3.25881E+14	3.25386E+14	3.25386E+14	3.25863E+14	3.25863E+14	3.25447E+14	3.25447E+14	3.25654E+14	3.25654E+14	3.26144E+14	3.26144E+14	3.2569E+14	3.2569E+14	3.25881E+14	3.25881E+14	3.25386E+14	3.25386E+14	3.25863E+14	3.25863E+14	3.25447E+14	3.25447E+14	3.25654E+14	3.25654E+14	3.26144E+14	3.26144E+14	3.2569E+14	3.2569E+14

按照这个图就可以直到SM和GPC的划分了

计算能力

MIG

一个GPUA100,可以划分7个独立的Instance.

参数列表

A100 一个block中最大可以容纳1024个threads
一个SM最大可以荣达2048个threads
所以最大可以容纳<216,1024>个线程