насколько я помню, так код не пишется. При распараллеливании цикла нет. Потоки объединяются в блоки, оптимальный размер блока определяется видекартой (количеством потоковых процессоров, объединяющихся в мультипроцессорный блок), а блоки объединяются в сетку-грид. Там есть следующие понятия:
dim3 gridDim - размер сетки-грида в блоках (под одно ядро вроде бы выделяется одна сетка, ядром называется кусок кода-микропрограмма-функция-которую надо выполнить над массивом данных)
dim3 blockDim - размер блока в нитях-тредах(все нити одного блока используют одну кеш-память)
dim3 blockIdx - индекс блока в сетке
dim3 threadIdx - индекс нити в блоке

код ниже всё поясняет

__global__ void increment_gpu(float *a, float b, int N)
{
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N)
a[idx] = a[idx] + b;
}
void main()
{
dim3 dimBlock (blocksize);
dim3 dimGrid( ceil( N / (float)blocksize) );
increment_gpu<<>>(a, b, N);
}

видимо эти переменные (blockIdx.x, blockDim.x, threadIdx.x) принимают своё значение для каждой копии функции increment_gpu своё, тем самым при заходе в тело функции можно узнать абсолютный номер вычислительного элемента, т.е. нужно написать следующее:

__device__ BYTE K[16];
__device__ WORD K1, K2, K3, K4 ;

__device__ void SetKey( unsigned char * key )
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
K[i]=key[i];
}
__global__ void chiper( BYTE *plain, DWORD newkey, DWORD *retdata )
{
SetKey(newkey);
.....
DoCrypt(plain);
....
}

и вызов делать например так chiper<<1,16>> или так chiper<<16,1>> .
Удачи.

P.S. это алгоритм AES?