Ядро CUDA не вызывается всеми блоками - PullRequest
Новая
       7

Ядро CUDA не вызывается всеми блоками

0 голосов
/ 11 июня 2018

Я получил странное поведение, пытаясь запустить простое векторное сложение.Если я запускаю приведенный ниже код с помощью функции printf, все работает нормально, и я получаю ожидаемый результат, 5050.

Теперь, если я удаляю функцию printf, выполняется только первый блок, и я получаю 2080, который являетсяожидаемый результат на сумму до 64.

Кто-нибудь знает, что здесь происходит?

Заранее спасибо за вашу помощь.

vecSum.cu: 

#include <cuda.h>
#include <cuda_runtime.h>
#include <iostream>
#include <math.h>

#define BLOCK_SIZE 64

__global__
void vecSumKernel(int N, float *d_v, float *d_out)
{
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
    unsigned int t = threadIdx.x;

    printf("Processing block #: %i\n", blockIdx.x);

    __shared__ float partialSum[BLOCK_SIZE];
    if(idx < N)
        partialSum[t] = d_v[idx];
    else
        partialSum[t] = 0;


    for(unsigned int stride=1; stride < BLOCK_SIZE; stride *= 2)
    {
        __syncthreads();
        if(t % (2*stride) == 0)
            partialSum[t] += partialSum[t + stride];
    }

    __syncthreads();
    *d_out += partialSum[0];
}

void vecSum_wrapper(int N, float *v, float &out, cudaDeviceProp devProp)
{
    float *d_v;
    float *d_out;
    size_t size = N*sizeof(float);

    cudaMalloc(&d_v, size);
    cudaMalloc(&d_out, sizeof(float));

    cudaMemcpy(d_v, v, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_out, &out, sizeof(float), cudaMemcpyHostToDevice);

    int nbrBlocks = ceil((float)N / (float)BLOCK_SIZE);
    vecSumKernel<<<nbrBlocks, BLOCK_SIZE>>>(N, d_v, d_out);

    cudaDeviceSynchronize();

    cudaMemcpy(&out, d_out, sizeof(float), cudaMemcpyDeviceToHost);

    cudaFree(d_v);
}

main.cpp: 

int main()
{
    ...

    int N = 100;

    float *vec = new float[N];

    for(int i=0; i < N; ++i)
        vec[i] = i + 1;

    std::chrono::time_point<timer> start = timer::now();

    float result = 0;
    vecSum_wrapper(N, vec, result, devProp);

    std::cout << "Operation executed in " << std::chrono::duration_cast<chrono>(timer::now() - start).count() << " ms \n";

    std::cout << "Result: " << result << '\n';

    delete[] vec;

    return 0;
}

1 Ответ

0 голосов
/ 11 июня 2018

Кажется, что последняя строка вашего ядра *d_out += partialSum[0] может представлять некоторые проблемы с параллелизмом, поскольку вы наверняка знаете, __syncthreads не синхронизирует блоки.atomicAdd может решить эту проблему параллелизма.

Что касается причины, по которой он работает лучше с printf, я бы предположил, что printf требует некоторой синхронизации, поэтому блоки не будут входить в негопоследняя инструкция в то же время, но у меня нет ничего, чтобы доказать это.

...