CUDA 4.0中P2P与UVA的性特性使用方法

近日，CUDA 4.0已经对注册开发者开放，其中增加了不少的功能。其中P2P（Peer-to-Peer ）与UVA（Unified Virtual Address Space ）的引进最为大家关心。这里与大家一起分享下SDK中的simpleP2P这个例子，他展示了如何使用这两个功能。
代码如下：
[cpp] view plaincopyprint?
  1. /* 
  2.  * Copyright 1993-2011 NVIDIA Corporation.  All rights reserved. 
  3.  * 
  4.  * Please refer to the NVIDIA end user license agreement (EULA) associated 
  5.  * with this source code for terms and conditions that govern your use of 
  6.  * this software. Any use, reproduction, disclosure, or distribution of 
  7.  * this software and related documentation outside the terms of the EULA 
  8.  * is strictly prohibited. 
  9.  * 
  10.  */  
  11. /* 
  12.  * This sample demonstrates a combination of Peer-to-Peer (P2P) and Unified 
  13.  * Virtual Address Space (UVA) features new to SDK 4.0 
  14.  */  
  15. #include <stdlib.h>   
  16. #include <stdio.h>   
  17. #include <string.h>   
  18. #include <cutil_inline.h>   
  19. #include <cuda_runtime_api.h>   
  20. const char *sSDKsample = "simpleP2P";  
  21. __global__ void SimpleKernel(float *src, float *dst)  
  22. {  
  23.     // Just a dummy kernel, doing enough for us to verify that everything   
  24.     // worked   
  25.     const int idx = blockIdx.x * blockDim.x + threadIdx.x;  
  26.     dst[idx] = src[idx] * 2.0f;  
  27. }  
  28. int main(int argc, char **argv)  
  29. {  
  30.     printf("[%s] starting.../n", sSDKsample);  
  31.     // Number of GPUs   
  32.     printf("Checking for multiple GPUs.../n");  
  33.     int gpu_n;  
  34.     cutilSafeCall(cudaGetDeviceCount(&gpu_n));  
  35.     printf("CUDA-capable device count: %i/n", gpu_n);  
  36.     if (gpu_n < 2)  
  37.     {  
  38.         printf("Two or more Tesla(s) with (SM 2.0) class GPUs are required for %s./n", sSDKsample);  
  39.         printf("Waiving test./n");  
  40.         printf("PASSED/n");  
  41.         exit(EXIT_SUCCESS);  
  42.     }  
  43.     // Query device properties   
  44.     cudaDeviceProp prop_0, prop_1;  
  45.     cutilSafeCall(cudaGetDeviceProperties(∝_0, 0));  
  46.     cutilSafeCall(cudaGetDeviceProperties(∝_1, 1));  
  47.     // Check for TCC   
  48. #ifdef _WIN32   
  49.     if (prop_0.tccDriver == 0 || prop_1.tccDriver == 0)  
  50.     {  
  51.         printf("Need to have both GPUs running under TCC driver to use P2P / UVA functionality./n");  
  52.         printf("PASSED/n");  
  53.         exit(EXIT_SUCCESS);  
  54.     }  
  55. #endif // WIN32   
  56.     // Check possibility for peer access   
  57.     printf("Checking for peer access.../n");  
  58.     int can_access_peer_0_1, can_access_peer_1_0;  
  59.     cutilSafeCall(cudaDeviceCanAccessPeer(&can_access_peer_0_1, 0, 1));  
  60.     cutilSafeCall(cudaDeviceCanAccessPeer(&can_access_peer_1_0, 1, 0));  
  61.     if (can_access_peer_0_1 == 0 || can_access_peer_1_0 == 0)  
  62.     {  
  63.         printf("Two or more Tesla(s) with (SM 2.0) class GPUs are required for %s./n", sSDKsample);  
  64.         printf("Peer access is not available between GPU0 <-> GPU1, waiving test./n");  
  65.         printf("PASSED/n");  
  66.         exit(EXIT_SUCCESS);  
  67.     }  
  68.     // Enable peer access   
  69.     printf("Enabling peer access.../n");  
  70.     cutilSafeCall(cudaSetDevice(0));  
  71.     cutilSafeCall(cudaDeviceEnablePeerAccess(1, 0));  
  72.     cutilSafeCall(cudaSetDevice(1));  
  73.     cutilSafeCall(cudaDeviceEnablePeerAccess(0, 0));  
  74.     // Check that we got UVA on both devices   
  75.     printf("Checking for UVA.../n");  
  76.     const bool has_uva = prop_0.unifiedAddressing && prop_1.unifiedAddressing;  
  77.     if (has_uva == false)  
  78.     {  
  79.         printf("At least one of the two GPUs has no UVA support/n");  
  80.     }  
  81.     // Allocate buffers   
  82.     const size_t buf_size = 1024 * 1024 * 16 * sizeof(float);  
  83.     printf("Allocating buffers (%iMB on GPU0, GPU1 and Host).../n", int(buf_size / 1024 / 1024));  
  84.     cutilSafeCall(cudaSetDevice(0));  
  85.     float* g0;  
  86.     cutilSafeCall(cudaMalloc(&g0, buf_size));  
  87.     cutilSafeCall(cudaSetDevice(1));  
  88.     float* g1;  
  89.     cutilSafeCall(cudaMalloc(&g1, buf_size));  
  90.     float* h0;  
  91.     if (has_uva)  
  92.         cutilSafeCall(cudaMallocHost(&h0, buf_size)); // Automatically portable with UVA   
  93.     else  
  94.         cutilSafeCall(cudaHostAlloc(&h0, buf_size, cudaHostAllocPortable));  
  95.     float *g0_peer, *g1_peer;  
  96.     if (has_uva == false)  
  97.     {  
  98.         // Need explicit mapping without UVA   
  99.         cutilSafeCall(cudaSetDevice(0));  
  100.         cutilSafeCall(cudaPeerRegister(g1, 1, cudaPeerRegisterMapped));  
  101.         cutilSafeCall(cudaPeerGetDevicePointer((void **) &g1_peer, g1, 1, 0));  
  102.         cutilSafeCall(cudaSetDevice(1));  
  103.         cutilSafeCall(cudaPeerRegister(g0, 0, cudaPeerRegisterMapped));  
  104.         cutilSafeCall(cudaPeerGetDevicePointer((void **) &g0_peer, g0, 0, 0));  
  105.     }  
  106.     // Create CUDA event handles   
  107.     printf("Creating event handles.../n");  
  108.     cudaEvent_t start_event, stop_event;  
  109.     float time_memcpy;  
  110.     int eventflags = cudaEventBlockingSync;  
  111.     cutilSafeCall(cudaEventCreateWithFlags(&start_event, eventflags));  
  112.     cutilSafeCall(cudaEventCreateWithFlags(&stop_event, eventflags));  
  113.     // P2P memcopy() benchmark   
  114.     cutilSafeCall(cudaEventRecord(start_event, 0));  
  115.     for (int i=0; i<100; i++)  
  116.     {  
  117.         // With UVA we don't need to specify source and target devices, the   
  118.         // runtime figures this out by itself from the pointers   
  119.         if (has_uva)  
  120.         {  
  121.             // Ping-pong copy between GPUs   
  122.             if (i % 2 == 0)  
  123.                 cutilSafeCall(cudaMemcpy(g1, g0, buf_size, cudaMemcpyDefault));  
  124.             else  
  125.                 cutilSafeCall(cudaMemcpy(g0, g1, buf_size, cudaMemcpyDefault));  
  126.         }  
  127.         else  
  128.         {  
  129.             // Ping-pong copy between GPUs   
  130.             if (i % 2 == 0)  
  131.                 cutilSafeCall(cudaMemcpyPeer(g1, 1, g0, 0, buf_size));  
  132.             else  
  133.                 cutilSafeCall(cudaMemcpyPeer(g0, 0, g1, 1, buf_size));  
  134.         }  
  135.     }  
  136.     cutilSafeCall(cudaEventRecord(stop_event, 0));  
  137.     cutilSafeCall(cudaEventSynchronize(stop_event));  
  138.     cutilSafeCall(cudaEventElapsedTime(&time_memcpy, start_event, stop_event));  
  139.     printf("cudaMemcpyPeer / cudaMemcpy between GPU0 and GPU1: %.2fGB/s/n",  
  140.         (1.0f / (time_memcpy / 1000.0f)) * ((100.0f * buf_size)) / 1024.0f / 1024.0f / 1024.0f);  
  141.    
  142.     // Prepare host buffer and copy to GPU 0   
  143.     printf("Preparing host buffer and memcpy to GPU0.../n");  
  144.     for (int i=0; i<buf_size / sizeof(float); i++)  
  145.     {  
  146.         h0[i] = float(i % 4096);  
  147.     }  
  148.     cutilSafeCall(cudaSetDevice(0));  
  149.     if (has_uva)  
  150.         cutilSafeCall(cudaMemcpy(g0, h0, buf_size, cudaMemcpyDefault));  
  151.     else  
  152.         cutilSafeCall(cudaMemcpy(g0, h0, buf_size, cudaMemcpyHostToDevice));  
  153.     // Kernel launch configuration   
  154.     const dim3 threads(512, 1);  
  155.     const dim3 blocks((buf_size / sizeof(float)) / threads.x, 1);  
  156.    
  157.     // Run kernel on GPU 1, reading input from the GPU 0 buffer, writing   
  158.     // output to the GPU 1 buffer   
  159.     printf("Run kernel on GPU1, taking source data from GPU0 and writing to GPU1.../n");  
  160.     cutilSafeCall(cudaSetDevice(1));  
  161.     if (has_uva)  
  162.         SimpleKernel<<<blocks, threads>>> (g0, g1);  
  163.     else  
  164.         SimpleKernel<<<blocks, threads>>> (g0_peer, g1);  
  165.     // Run kernel on GPU 0, reading input from the GPU 1 buffer, writing   
  166.     // output to the GPU 0 buffer   
  167.     printf("Run kernel on GPU0, taking source data from GPU1 and writing to GPU0.../n");  
  168.     cutilSafeCall(cudaSetDevice(0));  
  169.     if (has_uva)  
  170.         SimpleKernel<<<blocks, threads>>> (g1, g0);  
  171.     else  
  172.         SimpleKernel<<<blocks, threads>>> (g1_peer, g0);  
  173.    
  174.     // Copy data back to host and verify   
  175.     printf("Copy data back to host from GPU0 and verify.../n");  
  176.     if (has_uva)  
  177.         cutilSafeCall(cudaMemcpy(h0, g0, buf_size, cudaMemcpyDefault));  
  178.     else  
  179.         cutilSafeCall(cudaMemcpy(h0, g0, buf_size, cudaMemcpyHostToDevice));  
  180.     int error_count = 0;  
  181.     for (int i=0; i<buf_size / sizeof(float); i++)  
  182.     {  
  183.         // Re-generate input data and apply 2x '* 2.0f' computation of both   
  184.         // kernel runs   
  185.         if (h0[i] != float(i % 4096) * 2.0f * 2.0f)  
  186.         {  
  187.             printf("Verification error, element %i/n", i);  
  188.             if (error_count++ > 10)  
  189.                 break;  
  190.         }  
  191.     }  
  192.     printf((error_count == 0) ? "PASSED/n" : "FAILED/n");  
  193.     // Disable peer access (also unregisters memory for non-UVA cases)   
  194.     printf("Enabling peer access.../n");  
  195.     cutilSafeCall(cudaSetDevice(0));  
  196.     cutilSafeCall(cudaDeviceDisablePeerAccess(1));  
  197.     cutilSafeCall(cudaSetDevice(1));  
  198.     cutilSafeCall(cudaDeviceDisablePeerAccess(0));  
  199.     // Cleanup and shutdown   
  200.     printf("Shutting down.../n");  
  201.     cutilSafeCall(cudaEventDestroy(start_event));  
  202.     cutilSafeCall(cudaEventDestroy(stop_event));  
  203.     cutilSafeCall(cudaSetDevice(0));  
  204.     cutilSafeCall(cudaFree(g0));  
  205.     cutilSafeCall(cudaSetDevice(1));  
  206.     cutilSafeCall(cudaFree(g1));  
  207.     cutilSafeCall(cudaFreeHost(h0));  
  208.     cudaDeviceReset();  
  209.     cutilExit(argc, argv);  
  210. }  
 
从这段代码可以看出，目前仅有Fermi架构的tesla卡才能支持到P2P功能。由于UVA的需要，想成功编译运行程序，需要编译成64程序。  
 
而且如果支持UVA（其实现在如果是支持P2P的卡理论上应该都是支持UVA的），可以使用cudaMemcpyDefault代替原有的cudaMemcpyHostToDevice等方式，  
 
而且在内核函数等的调用上，不在需要分别获取各设备内存单独地址，大大的缩减了代码的编写量。
posted on 2012-03-21 11:21  carekee  阅读(1874)  评论(0)  编辑  收藏  举报

Powered by: 博客园 Copyright © 2024 carekee
Powered by .NET 8.0 on Kubernetes