Envoyé par
Alexko
Quel pourrait être l'intérêt de migrer des tâches au vol ?
Par exemple, tu as du code de ce genre (tiré du CUDA prog guide)
Code:
#include <stdlib.h>
__global__ void mallocTest()
{
__shared__ int* data;
// The first thread in the block does the allocation
// and then shares the pointer with all other threads
// through shared memory, so that access can easily be
// coalesced. 64 bytes per thread are allocated.
if (threadIdx.x == 0)
data = (int*)malloc(blockDim.x * 64);
__syncthreads();
// Check for failure
if (data == NULL)
return;
// Threads index into the memory, ensuring coalescence
int* ptr = data;
for (int i = 0; i < 64; ++i)
ptr[i * blockDim.x + threadIdx.x] = threadIdx.x;
// Ensure all threads complete before freeing
__syncthreads();
// Only one thread may free the memory!
if (threadIdx.x == 0)
free(data);
}
int main()
{
cudaDeviceSetLimit(cudaLimitMallocHeapSize, 128*1024*1024);
mallocTest<<<10, 128>>>();
cudaDeviceSynchronize();
return 0;
}
Le malloc et le free sont du code scalaire sur le chemin critique. Bon, dans ce cas, ça ne vaut probablement pas le coup de migrer. C'est plutôt un datapath scalaire optimisé latence dans chaque SM qu'il faudrait. (On y revient toujours...)