128 bit inteiro no cuda?

Acabei de instalar meu cuda SDK no Linux Ubuntu 10.04. Minha placa gráfica é uma NVIDIA GeForce GT 425M, e eu gostaria de usá-la para algum problema computacional pesado. O que eu me pergunto é: existe alguma maneira de usar alguns int de 128 bits não assinados? Ao usar o gcc para executar meu programa na CPU, eu estava usando o tipo __uint128_t, mas usá-lo com o cuda parece não funcionar. Existe alguma coisa que eu possa fazer para ter 128 bits inteiros no cuda?

Muito obrigado Matteo Monti Msoft Programação

Para um melhor desempenho, seria desejável mapear o tipo de 128 bits em cima de um tipo de vetor CUDA adequado, como uint4, e implementar a funcionalidade usando a assembly em linha PTX. A adição seria algo como isto:

typedef uint4 my_uint128_t; __device__ my_uint128_t add_uint128 (my_uint128_t addend, my_uint128_t augend) { my_uint128_t res; asm ("add.cc.u32 %0, %4, %8;\n\t" "addc.cc.u32 %1, %5, %9;\n\t" "addc.cc.u32 %2, %6, %10;\n\t" "addc.u32 %3, %7, %11;\n\t" : "=r"(res.x), "=r"(res.y), "=r"(res.z), "=r"(res.w) : "r"(addend.x), "r"(addend.y), "r"(addend.z), "r"(addend.w), "r"(augend.x), "r"(augend.y), "r"(augend.z), "r"(augend.w)); return res; } 

A multiplicação também pode ser construída usando a assembly em linha PTX, dividindo os números de 128 bits em blocos de 32 bits, calculando os produtos parciais de 64 bits e adicionando-os adequadamente. Obviamente, isso requer um pouco de trabalho. Pode-se obter um desempenho razoável no nível C, dividindo o número em blocos de 64 bits e usando __umul64hi () em conjunto com a multiplicação regular de 64 bits e algumas adições. Isso resultaria no seguinte:

 __device__ my_uint128_t mul_uint128 (my_uint128_t multiplicand, my_uint128_t multiplier) { my_uint128_t res; unsigned long long ahi, alo, bhi, blo, phi, plo; alo = ((unsigned long long)multiplicand.y << 32) | multiplicand.x; ahi = ((unsigned long long)multiplicand.w << 32) | multiplicand.z; blo = ((unsigned long long)multiplier.y << 32) | multiplier.x; bhi = ((unsigned long long)multiplier.w << 32) | multiplier.z; plo = alo * blo; phi = __umul64hi (alo, blo) + alo * bhi + ahi * blo; res.x = (unsigned int)(plo & 0xffffffff); res.y = (unsigned int)(plo >> 32); res.z = (unsigned int)(phi & 0xffffffff); res.w = (unsigned int)(phi >> 32); return res; } 

Abaixo está uma versão da multiplicação de 128 bits que usa a assembly em linha PTX. Ele requer o PTX 3.0, que é fornecido com o CUDA 4.2, e o código requer uma GPU com pelo menos capacidade de computação 2.0, ou seja, um dispositivo de class Fermi ou Kepler. O código usa o número mínimo de instruções, pois dezesseis multiplicações de 32 bits são necessárias para implementar uma multiplicação de 128 bits. Por comparação, a variante acima usando intrínsecos CUDA compila 23 instruções para um alvo sm_20.

 __device__ my_uint128_t mul_uint128 (my_uint128_t a, my_uint128_t b) { my_uint128_t res; asm ("{\n\t" "mul.lo.u32 %0, %4, %8; \n\t" "mul.hi.u32 %1, %4, %8; \n\t" "mad.lo.cc.u32 %1, %4, %9, %1;\n\t" "madc.hi.u32 %2, %4, %9, 0;\n\t" "mad.lo.cc.u32 %1, %5, %8, %1;\n\t" "madc.hi.cc.u32 %2, %5, %8, %2;\n\t" "madc.hi.u32 %3, %4,%10, 0;\n\t" "mad.lo.cc.u32 %2, %4,%10, %2;\n\t" "madc.hi.u32 %3, %5, %9, %3;\n\t" "mad.lo.cc.u32 %2, %5, %9, %2;\n\t" "madc.hi.u32 %3, %6, %8, %3;\n\t" "mad.lo.cc.u32 %2, %6, %8, %2;\n\t" "madc.lo.u32 %3, %4,%11, %3;\n\t" "mad.lo.u32 %3, %5,%10, %3;\n\t" "mad.lo.u32 %3, %6, %9, %3;\n\t" "mad.lo.u32 %3, %7, %8, %3;\n\t" "}" : "=r"(res.x), "=r"(res.y), "=r"(res.z), "=r"(res.w) : "r"(ax), "r"(ay), "r"(az), "r"(aw), "r"(bx), "r"(by), "r"(bz), "r"(bw)); return res; } 

CUDA não suporta números inteiros de 128 bits nativamente. Você pode falsificar as operações usando dois inteiros de 64 bits.

Olhe para este post :

 typedef struct { unsigned long long int lo; unsigned long long int hi; } my_uint128; my_uint128 add_uint128 (my_uint128 a, my_uint128 b) { my_uint128 res; res.lo = a.lo + b.lo; res.hi = a.hi + b.hi + (res.lo < a.lo); return res; } 

Uma resposta muito tardia, mas você poderia considerar o uso desta biblioteca:

https://github.com/curtisseizert/CUDA-uint128

que define uma estrutura de tamanho de 128 bits, com methods e funções utilitárias independentes para fazer com que funcione conforme o esperado, o que permite que ele seja usado como um inteiro regular. Na maioria das vezes.

Intereting Posts