커널에서 for 루프에서 멀티 스레딩으로 변경

Question

현재 멀티 스레딩과 관련하여 그리드 보간과 일부 문제가 있습니다. 이 코드는 2x2 행렬로 표시된 맵을 읽은 다음이를 보간하여 점의 수를 100 배로 늘리는 것으로 가정합니다. 커널에서 for 루프를 사용하면 큰 효과가 있습니다. 보간 전

: http://bildr.no/view/OWV1UDRO

보간 후에는 :

내가 스레드 루프의 변경을 시도 http://bildr.no/view/eTlmNmpo, 그것은 몇 가지 이상한 결과를 생산했다. 숫자 대신, 그것은 -1 결과 행렬을 가득 채웠다. #의 QNAN

여기 내 작업 코드가 커널 여기

#include <stdlib.h> 
#include <stdio.h> 
#include <math.h> 
#include <fstream> 
#include "cuda.h" 

using namespace std; 

float Z[41][41]; 

// Macro to catch CUDA errors in CUDA runtime calls 
#define CUDA_SAFE_CALL(call)           \ 
do {                 \ 
    cudaError_t err = call;           \ 
    if (cudaSuccess != err) {           \ 
     fprintf (stderr, "Cuda error in file '%s' in line %i : %s.\n",\ 
       __FILE__, __LINE__, cudaGetErrorString(err));  \ 
     exit(EXIT_FAILURE);           \ 
    }                 \ 
} while (0) 

// Macro to catch CUDA errors in kernel launches 
#define CHECK_LAUNCH_ERROR()           \ 
do {                 \ 
    /* Check synchronous errors, i.e. pre-launch */     \ 
    cudaError_t err = cudaGetLastError();        \ 
    if (cudaSuccess != err) {           \ 
     fprintf (stderr, "Cuda error in file '%s' in line %i : %s.\n",\ 
       __FILE__, __LINE__, cudaGetErrorString(err));  \ 
     exit(EXIT_FAILURE);           \ 
    }                 \ 
    /* Check asynchronous errors, i.e. kernel failed (ULF) */   \ 
    err = cudaThreadSynchronize();         \ 
    if (cudaSuccess != err) {           \ 
     fprintf (stderr, "Cuda error in file '%s' in line %i : %s.\n",\ 
       __FILE__, __LINE__, cudaGetErrorString(err));  \ 
     exit(EXIT_FAILURE);           \ 
    }                 \ 
} while (0) 

texture<float, 2, cudaReadModeElementType> tex; 


__global__ void kernel (int m, int n, float *f, float numberOfInterpolationsPerSquare) 
{ 
    int k = sqrt(numberOfInterpolationsPerSquare); 


    for (float i=0; i<n*k; i++) 
    { 
     for (float j=0; j<m*k; j++) 
     { 
     f[(int)(j+(m*k*i))] = tex2D (tex, j/k+0.5f, i/k+0.5f); 
     } 
    } 

} 

int main (void) 
{ 
    // Start timer 
    clock_t tStart = clock(); 

    // Size of map 
    int n=41; 
    int m=41; 

    int g = 0; 

    float numberOfInterpolationsPerSquare = 100; 
    float numberOfElements = pow(sqrt(numberOfInterpolationsPerSquare)*n,2); 

    size_t pitch, tex_ofs; 
    float *f; 
    float *r; 
    float *map_d = 0; 

    // Build read-Streams 
    ifstream map; 

    //Create and open a txt file for MATLAB 
    ofstream file; 

    // Open data 
    map.open("Map.txt", ios_base::in); 
    file.open("Bilinear.txt"); 

    // Store the map in a 2D array 
    for (int i=0; i<n; i++) 
    { 
     for (int j=0; j<m; j++) 
     { 
      map >> Z[i][j]; 
     } 
    } 

    // Allocate memory on host and device 
    CUDA_SAFE_CALL(cudaMallocPitch((void**)&map_d,&pitch,n*sizeof(*map_d),m)); 
    CUDA_SAFE_CALL(cudaMalloc((void**)&f, numberOfElements*sizeof(float))); 
    r = (float*)malloc(numberOfElements*sizeof(float)); 

    // Copy map from host to device 
    CUDA_SAFE_CALL(cudaMemcpy2D(map_d, pitch, Z, n*sizeof(Z[0][0]), n*sizeof(Z[0][0]),m,cudaMemcpyHostToDevice)); 

    // Set texture mode to bilinear interpolation 
    tex.normalized = false; 
    tex.filterMode = cudaFilterModeLinear; 

    // Bind the map to texture 
    CUDA_SAFE_CALL (cudaBindTexture2D (&tex_ofs, &tex, map_d, &tex.channelDesc, n, m, pitch)); 

    // Checking for offset 
    if (tex_ofs !=0) { 
     printf ("tex_ofs = %zu\n", tex_ofs); 
     return EXIT_FAILURE; 
    } 

    // Launch Kernel 
    kernel <<< 1,1 >>> (m, n, f, numberOfInterpolationsPerSquare); 
    CHECK_LAUNCH_ERROR();  
    CUDA_SAFE_CALL (cudaDeviceSynchronize()); 

    // Copy result from device to host 
    cudaMemcpy(r, f, numberOfElements*sizeof(float), cudaMemcpyDeviceToHost); 

    // Write results to file 
    for(int h=0;h<numberOfElements;h++) 
    { 
     if(g==sqrt(numberOfElements)) 
     { 
      file << endl; 
      g=0; 
     } 
     file << r[h] << " "; 
     g++; 
    } 

    // Free memory 
    CUDA_SAFE_CALL (cudaUnbindTexture (tex)); 
    CUDA_SAFE_CALL (cudaFree (map_d)); 
    CUDA_SAFE_CALL (cudaFree (f)); 
    free(r); 

    // Print out execution time 
    printf("Time taken: %.3fs\n", (double)(clock() - tStart)/CLOCKS_PER_SEC); 

    return EXIT_SUCCESS; 
} 

에서 루프로의이

작동하지 않는 멀티 스레딩과 커널의

__global__ void kernel (int m, int n, float *f, float numberOfInterpolationsPerSquare) 
{ 
    int k = sqrt(numberOfInterpolationsPerSquare); 

    int i= blockIdx.x * blockDim.x + threadIdx.x; 
    int j= blockIdx.y * blockDim.y + threadIdx.y; 


    if(i>=n*k || j>=m*k) 
     return; 

    f[(int)(j+(m*k*i))] = tex2D (tex, j/k+0.5f, i/k+0.5f); 

} 

멀티 스레드 버전이 작동하지 않는 이유를 알고 있습니까? 제 커널

안부

Sondre는

Answer 1

, i 및 j은 int 대신 float이다. 따라서 j/k과 i/k은 tex2D에 정수로 나뉩니다. 정수 나누기를 피하기 위해 k를 float로 선언하십시오.

//Find number of blocks 
int nthreads = 1024; 
int blocksize = 512; 
int nblocks = ceil((n*m*numberOfInterpolationsPerSquare)/nthreads); 

// Launch Kernel 
kernel <<< nblocks,blocksize >>> (m, n, f, numberOfInterpolationsPerSquare); 

위의 코드의 문제는 그것이 1D 블록의 1 차원 격자를 시작할 것이라고하지만, 커널 내부, 2D 인덱싱이 사용됩니다

처음 커널은 다음과 같은 구성으로 시작되었다. 커널을 올바르게 작동 시키려면 2D 그리드/블록 구성이 필요합니다. 그리드 다음 커널 코드의 외모에서/블록 구성이 작동합니다 :

float k = sqrt(numberOfInterpolationsPerSquare); 

const int threads_x = (int)ceil(n * k); 
const int threads_y = (int)ceil(m * k); 

const dim3 dimBlock(16,16); 

dim3 dimGrid; 
dimGrid.x = (threads_x + dimBlock.x - 1)/dimBlock.x; 
dimGrid.y = (threads_y + dimBlock.y - 1)/dimBlock.y; 

kernel<<<dimGrid,dimBlock>>>(m, n, f, numberOfInterpolationsPerSquare);